実現したいこと

OpenCLで+1するプログラムを作成してみました。
https://peta.okechan.net/blog/archives/2538
を参考にしています。

int index = get_global_id(0); 
int index2 = get_global_id(1);

を用いて data[index*M+index2] += 1.0; のようにアクセスしたいと思っています。

発生している問題

int index = get_global_id(0); 
int index2 = get_global_id(1);

を行うと常にindexとindex2に同じ値が格納されてしまいます。
ほとんどの部分が+1されず終わっています。

該当のソースコード

//
//  main3.cpp
//
#include <iostream>
#include <vector>
#include <OpenCL/opencl.h>
#include <numeric>

#define PLATFORM_MAX 4
#define DEVICE_MAX 4


void EC(cl_int result, const char *title)
{
    if (result != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << result << ")\n";
    }
}


cl_int err = CL_SUCCESS;
void EC2(const char *title)
{
    if (err != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << err << ")\n";
    }
    err = CL_SUCCESS;
}


int main(int argc, const char * argv[])
{
    // プラットフォーム一覧を取得
    cl_platform_id platforms[PLATFORM_MAX];
    cl_uint platformCount;
    EC(clGetPlatformIDs(PLATFORM_MAX, platforms, &platformCount), "clGetPlatformIDs");
    if (platformCount == 0) {
        std::cerr << "No platform.\n";
        return EXIT_FAILURE;
    }

    // 見つかったプラットフォームの情報を印字
    for (int i = 0; i < platformCount; i++) {
        char vendor[100] = {0};
        char version[100] = {0};
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, nullptr), "clGetPlatformInfo");
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(version), version, nullptr), "clGetPlatformInfo");
        std::cout << "Platform id: " << platforms[i] << ", Vendor: " << vendor << ", Version: " << version << "\n";
    }

    // デバイス一覧を取得
    cl_device_id devices[DEVICE_MAX];
    cl_uint deviceCount;
    EC(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, DEVICE_MAX, devices, &deviceCount), "clGetDeviceIDs");
    if (deviceCount == 0) {
        std::cerr << "No device.\n";
        return EXIT_FAILURE;
    }

    // 見つかったデバイスの情報を印字
    std::cout << deviceCount << " device(s) found.\n";
    for (int i = 0; i < deviceCount; i++) {
        char name[100] = {0};
        size_t len;
        EC(clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, &len), "clGetDeviceInfo");
        std::cout << "Device id: " << i << ", Name: " << name << "\n";
    }

    // コンテキストの作成
    cl_context ctx = clCreateContext(nullptr, 1, devices, nullptr, nullptr, &err);
    EC2("clCreateContext");

    // コンパイル済みclプログラムの読み込み
    const char* bitcode_path = "kernel2.cl.gpu_32.bc";
    size_t len = strlen(bitcode_path);
    cl_program program = clCreateProgramWithBinary(ctx, 1, devices, &len, (const unsigned char**)&bitcode_path, nullptr, &err);
    EC2("clCreateProgramWithBinary");

    // プログラムのビルド
    EC(clBuildProgram(program, 1, devices, nullptr, nullptr, nullptr), "clBuildProgram");

    // カーネルの作成
    cl_kernel kernel = clCreateKernel(program, "addone", &err);
    EC2("clCreateKernel");

    // データを用意
    int n = 10;
    std::vector<float> data(n*n,0.0f);
    // デバイスメモリを確保しつつデータをコピー
    cl_mem device_mem = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float) * n*n, data.data(), &err);
    EC2("clCreateBuffer");

    // カーネルの引数をセット
    EC(clSetKernelArg(kernel, 0, sizeof(cl_mem), &device_mem), "clSetKernelArg");
    EC(clSetKernelArg(kernel, 1, sizeof(int), &n), "clSetKernelArg");

    // コマンドキューの作成
    cl_command_queue q = clCreateCommandQueue(ctx, devices[0], 0, &err);
    EC2("clCreateCommandQueue");

    // カーネルの実行
    size_t global[2],local[2],offset[2];
    offset[0] = 0;
    offset[0] = 0;
    global[0] = n;
    global[1] = n;
    local[0] = 1;
    local[1] = 1;
    EC(clEnqueueNDRangeKernel(q, kernel, 2, offset, global, NULL, 0, nullptr, nullptr), "clEnqueueNDRangeKernel");

    // 結果を読み込み
    EC(clEnqueueReadBuffer(q, device_mem, CL_TRUE, 0, sizeof(float) * n*n, data.data(), 0, nullptr, nullptr), "clEnqueueReadBuffer");

    // 結果の印字
    for (int i = 0; i < n*n; i++) {
        std::cout << data[i] << ", ";
    }
    std::cout << "\n";
    float total = std::accumulate(data.begin(),data.end(),0.0);
    std::cout << total << std::endl;
    // コマンドキューの解放
    EC(clReleaseCommandQueue(q), "clReleaseCommandQueue");

    // デバイスメモリを解放
    EC(clReleaseMemObject(device_mem), "clReleaseMemObject");

    // カーネルの解放
    EC(clReleaseKernel(kernel), "clReleaseKernel");

    // プログラムの解放
    EC(clReleaseProgram(program), "clReleaseProgram");

    // コンテキストの解放
    EC(clReleaseContext(ctx), "clReleaseContext");

    std::cout << "Done.\n";
    return EXIT_SUCCESS;
}



//ここからカーネル部分
//kernel2.cl
__kernel
void addone(__global float* data,const int n)
{
    int index = get_global_id(0);
    int index2 = get_global_id(1);
    int dim = get_work_dim();
    printf("get_work_dim = %d\n",dim);
    printf("index = %d , index2 = %d \n",index,index2);
    data[index*n+index2] += 1.0f;
}

試したこと

EC(clEnqueueNDRangeKernel(q, kernel, 2, offset, global, local, 0, nullptr, nullptr), "clEnqueueNDRangeKernel");

のところのglobal, localなどの数値を変えるなどして動かしてみました。

実行環境

MacBook Pro 13インチ、OSはel capitanです。

Terminal上で 

/System/Library/Frameworks/OpenCL.framework/Libraries/openclc -c -o kernel2.cl.gpu_32.bc -arch gpu_32 -emit-llvm kernel2.cl 
g++ -O3 -std=c++11 -framework opencl main3.cpp -o test 
./test

で実行しました。