以前に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)
このブログの開発用PCはこちら
