CUDAのプログラムをOpenCLに移植(ドット積)

CUDA by Exampleには、ドット積をCUDAで実装する方法が記載されている。

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

これを、OpenCLに移植してみよう。

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

基本的なドット積のアルゴリズム

        float mem_Ans[MEM_SIZE];
        for (int i = 0; i < MEM_SIZE; i++) {
                mem_Ans[i] = mem_A[i] * mem_B[i];
        }
        for (int i = 1; i < MEM_SIZE; i++) {
                mem_Ans[0] += mem_Ans[i];
        }
        printf("Ans = %f\n", mem_Ans[0]);

簡単！

github.com

要点: CUDAとOpenCLのメモリアクセス方法

このプログラムの要点は、CUDAコア間の通信をどのように実現するかという点と、CUDA間の通信をどのように同期するかという点だ。 CUDAとOpenCLでは、その点は殆ど同じだった(今のところは) グローバルRAMとして配置している場合は、配列の値は、インデックスを用いてアクセスできるし、隣りのCUDAコアで実行されたデータも、インデックスを変化させることでアクセス可能だ。

...
    while (tid < N) {
        temp += a[tid] * b[tid];
        tid += blockDim.x * gridDim.x;
    }
...
    // This code, threadsPerBlock should be power of 2
    int i = blockDim.x / 2;
    while (i != 0) {
        if (cacheIndex < i) {
            cache[cacheIndex] += cache[cacheIndex+i];
        }
        __syncthreads();
        i /= 2;
    }
...

OpenCLでも基本的に同じ実装にした。

        int gid = get_global_id(0);
        array_C[gid] = array_A[gid] * array_B[gid];

要点: CUDAとOpenCLのメモリアクセス同期

また、同期についても同様に実装できる。CUDAでは __syncthreads(); を使ったが、OpenCLではスレッドを使うほど細かく制御はしていないので、barrier()を使っている。

        int i = 128 / 2;
        while (i != 0) {
            if (gid < i) {
                        array_C[gid] += array_C[gid + i];
                }
                barrier(CLK_GLOBAL_MEM_FENCE);
                i /= 2;
        }

barrier()を外してしまうと、答えが間違ってしまうことまで確認できた。すばらしい！ でも、今はスレッドを細かく使いこなすことができていない。もうちょっと頑張ってみよう。