OpenCLで正方行列の転置をとってみた
いままで、ホストコードでのOpeCL APIについて説明してきたが、諸般の都合*1により残りは一旦省略し、実際のコードを見ていく。
前まではホストコードをC言語で記述するとして説明してきたのに、いきなりC++バインディングで書いてあるとか気にしない。C++つええよC++
概要
今回はとりあえず定番な行列の転置をとってみることにした。アルゴリズムについてが本題ではないので、単純化のため正方行列のみを対象とすることにした。
お前の書いた汚ないコードなんぞ読めん!という人は(そうでなくともいいけれど)AMD Stream SDKなりNVIDIA OpenCL SDKなりのサンプルコードにきちんと含まれているのでそちらも見るとよい。
C++バインディング自体はこちらを参考にした。http://developer.amd.com/gpu/ATIStreamSDK/pages/TutorialOpenCL.aspx
あとはこちらにdoxygenで生成された公式のドキュメントがある。http://www.khronos.org/registry/cl/
ちなみにこのコードはWinXP x64+Visual Studio 2008+AMD Stream SDK 2.0 bate4で動作確認をしている。ただし32bitコードとして。
ソースコード
host.cpp
#include <utility> #define __NO_STD_VECTOR // Use cl::vector and cl::string and #define __NO_STD_STRING // not STL versions, more on this later #include <CL/cl.hpp> #include <cstdio> #include <cstdlib> #include <fstream> #include <iostream> #include <string> #include <iterator> #define MATRIX_SIZE 8 #define BLOCK_SIZE 4 inline void checkErr(cl_int err, const char * name) { if (err != CL_SUCCESS) { std::cerr << "ERROR: " << name << " (" << err << ")" << std::endl; exit(EXIT_FAILURE); } } int _tmain(int argc, _TCHAR* argv[]) { cl_int err; int ary[MATRIX_SIZE][MATRIX_SIZE]; int outb[MATRIX_SIZE][MATRIX_SIZE]; int i,j; cl::Context context(CL_DEVICE_TYPE_CPU, NULL, NULL, NULL, &err); checkErr(err, "cl::Context()"); // 転置を取る行列を生成 for (i = 0;i < MATRIX_SIZE; i++) { for (j = 0; j < MATRIX_SIZE; j++) { ary[i][j] = rand()%8; std::cout << ary[i][j] << " "; } std::cout << std::endl; } std::cout << std::endl; // カーネルとの情報交換用のバッファを確保 cl::Buffer in(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, ary, &err); checkErr(err, "alloc input buffer"); cl::Buffer out(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, outb, &err); checkErr(err, "alloc output buffer"); cl::vector<cl::Device> devices; devices = context.getInfo<CL_CONTEXT_DEVICES>(); std::ifstream file("kernel.cl"); checkErr(file.is_open() ? CL_SUCCESS:-1, "kernel.cl"); std::string prog( std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>())); //std::cout << prog; cl::Program::Sources src(1, std::make_pair(prog.c_str(), prog.length()+1)); cl::Program p(context, src); err = p.build(devices, "", NULL, NULL); if (err != CL_SUCCESS) { cl::string log = p.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]); std::cout << log.c_str() << std::endl; exit(EXIT_FAILURE); } cl::Kernel kernel(p, "trans", &err); checkErr(err, "alloc kernel"); checkErr(kernel.setArg(0, in), "Set Arg[0]"); checkErr(kernel.setArg(1, out), "Set Arg[1]"); checkErr(kernel.setArg(2,sizeof(cl_int)*BLOCK_SIZE*BLOCK_SIZE, NULL), "Set Arg[2]"); checkErr(kernel.setArg(3, MATRIX_SIZE), "Set Arg[3]"); checkErr(kernel.setArg(4, BLOCK_SIZE), "Set Arg[4]"); cl::CommandQueue queue(context, devices[0], NULL, &err); cl::Event eve; queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(MATRIX_SIZE,MATRIX_SIZE),cl::NDRange(BLOCK_SIZE,BLOCK_SIZE), NULL, &eve); eve.wait(); queue.enqueueReadBuffer(out, CL_TRUE,0, sizeof(int)*MATRIX_SIZE*MATRIX_SIZE,outb, NULL,&eve); for (i = 0; i < MATRIX_SIZE; ++i) { for (j = 0; j < MATRIX_SIZE; j++) printf("%d ", outb[i][j]); std::cout << std::endl; } return 0; }
kernel.cl
__kernel void trans(__global int in[], __global int out[], __local int block[], int matrix_size, int block_size) { size_t gx = get_global_id(0); size_t gy = get_global_id(1); size_t lx = get_local_id(0); size_t ly = get_local_id(1); block[lx + ly*block_size] = in[gx + gy*matrix_size]; barrier(CLK_LOCAL_MEM_FENCE); size_t groupx = get_group_id(0); size_t groupy = get_group_id(1); uint tx = groupy * block_size + ly; uint ty = groupx * block_size + lx; out[tx + ty*matrix_size] = block[lx + ly*block_size]; }
ホストコードの解説
#include <CL/cl.hpp>
Cならば
#include <CL/cl.h>
となる。
はじめの
inline void checkErr(cl_int err, const char *name)
は、OpenCL APIがエラーを返したときに、エラー箇所を表示して終了する手続きになっている。
はじめの
cl::Context context(CL_DEVICE_TYPE_CPU, NULL, NULL, NULL, &err);
は
clCraeteContextFromType(NULL, CL_DEVICE_CPU, NULL, NULL, &err);
に相当し、CPUなデバイスを含むコンテキストを作成している。
その後に乱数を発生させて適当な行列を生成し、表示する。
cl::Buffer in(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, ary, &err); cl::Buffer out(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, outb, &err);
は
in = clCraeteBuffer(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, ary, &err); out = clCraeteBuffer(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_int)*MATRIX_SIZE*MATRIX_SIZE, outb, &err);
にあたり、行列分の入出力用のバッファを作成する。
次に
context.getInfo<CL_CONTEXT_DEVICES>();
(getContextInfo()に相当)とデバイスIDのベクトルを取得する。
この後の部分では、kernel.clをメモリ上にロードし、それを文字列と文字列長のタプルにして、cl_programを作成している。
ここが、clCreateProgramWithSource()に相当する。
そして、カーネルコードをビルド(p.build())している。
if (err != CL_SUCCESS) { cl::string log = p.getBuildInfo<CL_PROGRAM_BUILD_LOG>(devices[0]); std::cerr << log.c_str() << std::endl; exit(EXIT_FAILURE); }
は、ビルドが正常に終了しなかったときに、ビルドログを出力し終了している。
cl::Kernel kernel(p, "trans", &err); checkErr(err, "alloc kernel"); checkErr(kernel.setArg(0, in), "Set Arg[0]"); checkErr(kernel.setArg(1, out), "Set Arg[1]"); checkErr(kernel.setArg(2,sizeof(cl_int)*BLOCK_SIZE*BLOCK_SIZE, NULL), "Set Arg[2]"); checkErr(kernel.setArg(3, MATRIX_SIZE), "Set Arg[3]"); checkErr(kernel.setArg(4, BLOCK_SIZE), "Set Arg[4]");
では、実行するカーネルを作成し引数を設定している。ローカルメモリ上の変数を渡す場合には、データの大きさだけ渡し、データ部分にはNULLを渡せばよい。
その後、命令キューを作成し、カーネルを投入する。このとき、global_work_sizeには実行してほしいwork-itemの総数を渡す。実行して欲しいworkgroupの数ではないことに注意。そしてlocal_work_sizeには、ひとつのworkgroupあたりのwork-item数を渡す。この数値からworkgroup数はOpenCL Runtimeが自分で計算してくれる。なお、local_work_sizeにNULLを渡すとworkgroup内のwork-item数も自動で設定してくれるらしい。
のこりで、カーネルの実行終了を待ち(eve.wait)結果を取得して(out.enqueueReadBuffer())出力している。
カーネルコードの解説
といってもあまりすることはないが、まずグローバルIDとworkgroup内のローカルIDを取得して、高速なローカルメモリにフェッチする。そして書き込み先のインデックスを計算し書き込んでいる。
ここで新出なのがbarrier()である。これはworkgroup内での同期のための組み込み関数(BUild-in function)で、次のどれかを引数に取る。
- CLK_MEM_LOCAL_FENCE
- CLK_MEM_GLOBAL_FENCE
意味はそれぞれ、それ以前に現れたローカルメモリないしグローバルメモリへの書き込みを全て終了するまで待つことである。
おわりに
今回は以上。不審な点、疑問や訂正等あればご気楽にどうぞ。
