OpenCL: 2D Array

Untuk memproses array 2D di OpenCL, kita perlu nge-flatten array 2D itu jadi 1D. Bahkan mau sampai 100D tetap perlu di-flatten ke 1D. Jadi teknik untuk memproses data yang lebih dari 1D adalah dengan cara Indexing

Contoh grid 3×4 (H=3, W=4):

Maka kita perlu flatten ke 1D. Kira-kira gini:

Lalu untuk indexing 2D Array, formula sakti nya adalah

idx = y * width + x

Misal kita mau ambil y=1, x=2, dan kita punya 4 kolom (width=4). elemen array yang nilainya 6.

idx = y * width + x
idx = 1 * 4 + 2
idx = 6

Code

Contoh buat data. Vector a, b, c di sini semuanya tetap 1D. Tapi untuk menentukan dimensinya ditentukan ketika kita mengkonfigurasi kernel.

const int W = 4;
const int H = 3;

std::vector<float> a(W * H);
    std::vector<float> b(W * H);
    std::vector<float> c(W * H);
    
    // isi data
    for (int y = 0; y < H; y++) {
        for (int x = 0; x < W; x++) {
            a[y * W + x] = y * 10 + x;
            b[y * W + x] = 1.0f;
        }
    }

Lalu buat buffer, yang ini sama aja. Perhatikan ketika kita menentukan sizenya

cl_mem bufA = clCreateBuffer(
    context, 
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    sizeof(float) * W * H, 
    a.data(), 
    nullptr
);
cl_mem bufB = clCreateBuffer(
    context,
    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
    sizeof(float) * W * H, 
    b.data(), 
    nullptr
);
cl_mem bufC = clCreateBuffer(
    context,
    CL_MEM_WRITE_ONLY,
    sizeof(float) * W * H, 
    nullptr, 
    nullptr
);

code kernel:

__kernel void mat_add(__global const float* A,
                      __global const float* B,
                      __global float* C,
                      int width) {

    int x = get_global_id(0);
    int y = get_global_id(1);

    int idx = y * width + x;
    C[idx] = A[idx] + B[idx];

    printf("x=%d, y=%d, idx=%d, a=%f, b=%f, c=%f\n",
            x,    y,    idx,    A[idx], B[idx], C[idx]);
}

Lalu untuk execute

size_t global[2] = { (size_t)W, (size_t)H };
clEnqueueNDRangeKernel(
  queue,  
  kernel, 
  2,   // DIMENSI 
  nullptr, 
  global, 
  nullptr, 
  0, 
  nullptr,   
  nullptr
);

Lalu read result

clEnqueueReadBuffer(queue, bufC, CL_TRUE, 0, sizeof(float) * W * H, c.data(), 0, nullptr, nullptr);