以前にOpenCLについて少しだけ触れたが、その時はOpenCL環境を確認したのみだった。今回はOpenCLを使って実際に簡単な計算をしてみようと思う。
OpenCLで実際に計算を行うまでの流れは次のようになる。
- プラットフォーム情報を取得
- コンテキストを生成
- コンテキスト内のデバイスを取得
- 使用するデバイスにコマンドキューを生成
- プログラムオブジェクトの生成
- プログラムのビルド
- カーネルの生成
- メモリオブジェクトの生成
- カーネルをコマンドキューに入れる
- 結果の取得
- リソースの解放
今回のコードはOpenCL入門―GPU&マルチコアCPU並列プログラミング for MacOS Windows LinuxのChapter 3を参考に、エラー処理を省き、実行タスクを処理が重いものに変更し、実行時間計測を入れている。実行時間の計測はBoost.Chronoを使用した。
ソースは次のようになる。
#include <iostream> #include <cmath> #include <boost/chrono.hpp> #include <boost/chrono/duration.hpp> using namespace boost; #ifdef __APPLE__ #include #else #include #endif //__APPLE__ const int nElements = 9000000; const int maxDevices = 10; float input1[nElements]; float input2[nElements]; float output[nElements]; void addVector(float* input1, float* input2, float* output){ for(int i = 0; i < nElements; i++){ output[i] = sin(input1[i]) * sin(input2[i]); output[i] = cos(output[i]); output[i] = pow(output[i], output[i]); } } int main(int argc, char* argv[]) { cl_int status; cl_platform_id platforms[10]; cl_uint num_platforms; status = clGetPlatformIDs(sizeof(platforms) / sizeof(platforms[0]), platforms, &num_platforms); // 1. プラットフォーム情報を取得 cl_context_properties properties[] = {CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[0], 0}; cl_context context = clCreateContextFromType(properties, CL_DEVICE_TYPE_GPU, NULL, NULL, &status); // 2. コンテキストを生成 cl_device_id devices[maxDevices]; size_t size_return; status = clGetContextInfo(context, CL_CONTEXT_DEVICES, sizeof(devices), devices, &size_return); // 3. コンテキスト内のデバイスを取得 cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, &status); // 4. 使用するデバイスにコマンドキューを生成 const char *sources[] = { "__kernel void\n\ addVector(__global const float *input1,\n\ __global const float *input2,\n\ __global float *output)\n\ {\n\ int index = get_global_id(0);\n\ output[index] = sin(input1[index]) * sin(input2[index]);\n\ output[index] = cos(output[index]);\n\ output[index] = pow(output[index], output[index]);\n\ }\n"}; // カーネルのソースファイル cl_program program = clCreateProgramWithSource(context, 1, (const char**)&sources, NULL, &status); // 5. プログラムオブジェクトの生成 status = clBuildProgram(program, 1, devices, NULL, NULL, NULL); clUnloadCompiler(); // 6. プログラムのビルド cl_kernel kernel = clCreateKernel(program, "addVector", &status); // 7. カーネルの生成 for (int i = 0; i < nElements; i++) { input1[i] = (float)i * 10.0f; input2[i] = (float)i / 20.0f; output[i] = 0.0f; } cl_mem memInput1 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * nElements, input1, &status); // 8. メモリオブジェクトの生成 cl_mem memInput2 = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * nElements, input2, &status); // 8. メモリオブジェクトの生成
cl_mem memOutput = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * nElements, NULL, &status); // 8. メモリオブジェクトの生成 status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memInput1); // カーネル引数にセット status = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memInput2); // カーネル引数にセット status = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memOutput); // カーネル引数にセット boost::chrono::steady_clock::time_point start = boost::chrono::steady_clock::now(); // 時間計測開始 size_t globalSize[] = {nElements}; status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, globalSize, 0, 0, NULL, NULL); // 9. カーネルをコマンドキューに入れる status = clEnqueueReadBuffer(queue, memOutput, CL_TRUE, 0, sizeof(cl_float) * nElements, output, 0, NULL, NULL); // 10. 結果の取得 boost::chrono::duration sec = boost::chrono::steady_clock::now() - start; // 時間計測終了 std::cout << "input1, input2, output" << std::endl; for (int i = 0; i < 100; i++) { std::cout << input1[i] << ", " << input2[i] << ", " << output[i] << std::endl; } std::cout << "OpenCL addVector " << sec.count() << " seconds\n"; clReleaseMemObject(memOutput); // 11. リソース解放 clReleaseMemObject(memInput2); // 11. リソース解放 clReleaseMemObject(memInput1); // 11. リソース解放 clReleaseKernel(kernel); // 11. リソース解放 clReleaseProgram(program); // 11. リソース解放 clReleaseCommandQueue(queue); // 11. リソース解放 clReleaseContext(context); // 11. リソース解放 boost::chrono::steady_clock::time_point start2 = boost::chrono::steady_clock::now(); // 時間計測開始 addVector(input1, input2, output); // 通常の関数 boost::chrono::duration sec2 = boost::chrono::steady_clock::now() - start2; // 時間計測終了 std::cout << "Normal addVector " << sec2.count() << " seconds\n"; return 0; }
これをビルドし実行したところ次のような結果となった。
OpenCL addVector 0.126067 seconds
Normal addVector 1.13217 seconds
ちなみに実行環境は次のようになっている。
- CPU : Core i7-2670QM (2.20GHz 4Core/8Thread VT:enable)
- GPU : AMD RADEON 6770M
- MEM : 8GB
- HDD : 160GB SSD (Intel 320)
- OS : Windows 7 Ultimate (x64)
コメント