CUDAのプログラムをOpenCLに移植(VecAdd編)

CUDA by Example 汎用GPUプログラミング入門

CUDA by Exampleのサンプルプログラムを移植して、OpenCLで動作させてみようと思う。

改訂新版 OpenCL入門 1.2対応 マルチコアCPU・GPUのための並列プログラミング

デバイスIDとデバイスタイプを調査するためのOpenCLプログラムを記述する

OpenCLは、CPUもGPUも同じAPIを使って動作させることができる。つまり、設定を間違えると、GPUではなくCPUで動作させてしまう可能性もある。 まず、OpenCLでデバイスIDを調査するプログラムを書いてみた。

clGetPlatformIDs

platformIDを利用して、OpenCLが認識可能なデバイスを特定する。

 cl_platform_id platform_id[MAX_PLATFORM_ID] = { NULL };
...
    /* Obtain information of Platform Device */
    ret = clGetPlatformIDs(MAX_PLATFORM_ID, platform_id, &ret_num_platform);

    for (cl_uint pid = 0; pid < ret_num_platform; pid++) {
        ret = clGetDeviceIDs(platform_id[pid], CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
        cl_device_type device_type;
        size_t type_size;
        ret = clGetDeviceInfo(device_id, CL_DEVICE_TYPE, sizeof(device_type), &device_type, &type_size);
        if (type_size == sizeof(device_type)){
            fprintf(stderr, "Success to get device type. Device_Type=%d\n", device_type);
            if ((device_type & CL_DEVICE_TYPE_CPU) != 0) {
                fprintf(stderr, "Device_Type : CPU\n");
            }
            else if ((device_type & CL_DEVICE_TYPE_GPU) != 0) {
                fprintf(stderr, "Device Type : GPU\n");
            }
            else if ((device_type & CL_DEVICE_TYPE_ACCELERATOR) != 0) {
                fprintf(stderr, "Device Type : Accelerator\n");
            }
            else if ((device_type & CL_DEVICE_TYPE_DEFAULT) != 0) {
                fprintf(stderr, "Device Type : Default\n");
            }
        }
        else {
            fprintf(stderr, "Failed to get device type");
        }
    }

これを実行してみると、以下のようになった。

Success to get device type. Device_Type=2
Device_Type : CPU
Success to get device type. Device_Type=4
Device Type : GPU

やはり、CPUとGPUの2つが認識されている。

VecAdd を移植する (OpenCLとCUDAを比較)

では、以下のプログラムをOpenCLで記述してみよう。

for (int i = 0; i < MEM_SIZE; i++) {
    c[i] = a[i] + b[i];
}

基本的に、インデックスiで並列化して動作させる。

• OpenCL

 /* Craete Memory Buffer */
    memobj_a = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(int), NULL, &ret);
    memobj_b = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(int), NULL, &ret);
    memobj_c = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(int), NULL, &ret);
    /* Translate Memory Buffer */
    ret = clEnqueueWriteBuffer(command_queue, memobj_a, CL_TRUE, 0, MEM_SIZE * sizeof(int), a, 0, NULL, NULL);
    ret = clEnqueueWriteBuffer(command_queue, memobj_b, CL_TRUE, 0, MEM_SIZE * sizeof(int), b, 0, NULL, NULL);

まずは、メモリオブジェクトを作成した。これは、CUDAの場合は、CudaMemAllocとCudaMemcpyに相当する記述だ。

• CUDA

    // Allocate GPU's memory
    HANDLE_ERROR (cudaMalloc ((void **)&dev_a, N * sizeof (int)));
    HANDLE_ERROR (cudaMalloc ((void **)&dev_b, N * sizeof (int)));
    HANDLE_ERROR (cudaMalloc ((void **)&dev_c, N * sizeof (int)));

    // Set Array a and b by CPU
    for (int i = 0; i < N; i++) {
        a[i] = -i;
        b[i] = i * i;
    }

    // Copy Array a and b into GPU
    HANDLE_ERROR (cudaMemcpy (dev_a, a, N * sizeof (int),
                              cudaMemcpyHostToDevice));
    HANDLE_ERROR (cudaMemcpy (dev_b, b, N * sizeof (int),
                              cudaMemcpyHostToDevice));

次に、カーネルプログラムを記述してみよう。

• OpenCL

__kernel void vecAdd(__global int *a, __global int *b, __global int *c)
{
    int gid = get_global_id(0);

    c[gid] = a[gid] + b[gid];
}

• CUDA

__global__ void add (int *a, int *b, int *c)
{
    int tid = blockIdx.x;

    if (tid < N) {
        c[tid] = a[tid] + b[tid];
    }
}

カーネルが完成したので、カーネルを起動するプログラムを記述する。

• OpenCL

 /* Make kernel program from read source */
    program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);

    /* Build kernel program */
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

    /* Build OpenCL kernel */
    kernel = clCreateKernel(program, "vecAdd", &ret);

    /* Setting OpenCL kernel arguments */
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobj_a);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobj_b);
    ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memobj_c);

    size_t global_work_size[3] = { MEM_SIZE, 0, 0 };
    size_t local_work_size[3] = { MEM_SIZE, 0, 0 };

    /* Execute OpenCL kernel */
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);

• CUDA

    add<<<N, 1>>> (dev_a, dev_b, dev_c);

CUDAの方が、記述が簡単なイメージがある。OpenCLの方が汎用性を求めている分、記述が冗長だな。

結果を回収して表示するプログラムを記述する。

• OpenCL

 clEnqueueReadBuffer(command_queue, memobj_c, CL_TRUE, 0, MEM_SIZE * sizeof(int), c, 0, NULL, NULL);

    /* Show Results */
    for (i = 0; i < MEM_SIZE; i++) {
        printf("mem[%d] : %d %d %d\n", i, a[i], b[i], c[i]);
    }

• CUDA

    // Copy Array c into CPU
    HANDLE_ERROR (cudaMemcpy (c, dev_c, N * sizeof(int),
                              cudaMemcpyDeviceToHost));
    // Display result
    for (int i = 0; i < N; i++) {
        printf ("%d + %d = %d\n", a[i], b[i], c[i]);
    }

実行結果

• OpenCL

add_loop_gpu.exe
Success to get device type. Device_Type=2
Device_Type : CPU
Success to get device type. Device_Type=4
Device Type : GPU
mem[0] : 0 0 0
mem[1] : -1 1 0
mem[2] : -2 4 2
mem[3] : -3 9 6
mem[4] : -4 16 12
mem[5] : -5 25 20
mem[6] : -6 36 30
mem[7] : -7 49 42
mem[8] : -8 64 56
mem[9] : -9 81 72

ちゃんと動いた！

OpenCLとCUDAのプログラミングモデルの違い

基本的には同じ、という印象だ。ただし、CUDAの方がGPUに特化している分、記述が簡単に見える。まあ、今回のプログラムは非常に簡単だから、そう見えるだけなのかな。 CUDAもOpenCLも、メモリの転送というところでプログラムの記述量が増えている。 どちらもGPUのメモリを確保し、転送、計算して、回収という部分が必要になっている。ここをプログラミングモデルとして隠蔽できれば、もうちょっと簡単に書けそうなのにな...