OpenCLで実際に計算してみる

以前にOpenCLについて少しだけ触れたが、その時はOpenCL環境を確認したのみだった。今回はOpenCLを使って実際に簡単な計算をしてみようと思う。

OpenCLで実際に計算を行うまでの流れは次のようになる。

1. プラットフォーム情報を取得
2. コンテキストを生成
3. コンテキスト内のデバイスを取得
4. 使用するデバイスにコマンドキューを生成
5. プログラムオブジェクトの生成
6. プログラムのビルド
7. カーネルの生成
8. メモリオブジェクトの生成
9. カーネルをコマンドキューに入れる
10. 結果の取得
11. リソースの解放

今回のコードは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はこちら