Implementing nested loop in OpenCL?

Question

I'm new to OpenCL, been trying to implement a 3 level nested loop in Kernel function. Guess my understanding is not enough. Below is the C code of the logic

void scale(float *output, float *scales, int batch, int n, int size)
{
    int i,j,b;
    for(b = 0; b < batch; ++b){
        for(i = 0; i < n; ++i){
            for(j = 0; j < size; ++j){
                output[(b*n+i)*size+j] *= scales[i];
            }
        }
    }
}

Where output and scales are 1D arrays. Ex:

float output[18] = {1,2,3,4,5,6,7,8,9,1,2,3,4,5,6,7,8,9};
float scales[9] = {1,0,1,0,1,0,1,0,1};

int n = 9;
int size = 2;
int batch = 1;

The expected output is Output:

1.000000  2.000000  0.000000  0.000000  5.000000  6.000000  
0.000000  0.000000  9.000000  1.000000  0.000000  0.000000 
4.000000  5.000000  0.000000  0.000000  8.000000  9.000000

Below is my OpenCL kernel

__kernel void scale_kernel(__global float *output, __global float *biases, int n, int size)
{
    int j = get_global_id(0);
    int i = get_group_id(1);
    int b = get_group_id(2);

    if(j < size) output[(b*n+i)*size + j] *= biases[i];
}

I hope this implementation is correct and the way I'm launching the NDkernel is wrong. My BLOCK size is 16 (Think this is where my understanding is wrong).

size_t global_work_size[3] = {size-1)/BLOCK + 1, n, batch};
size_t local_work_size[3] = {BLOCK, 1, 1};
cl.error = clEnqueueNDRangeKernel(queue, kernel, 3, 0, global_work_size, local_work_size, 0, 0, NULL);

EDIT 1:

Changing the global_work_size as below produces the expected output, I've set local_work_size as NULL in this case. This might not provide the best performance.

size_t global_work_size[3] = {size, n, batch};
    cl.error = clEnqueueNDRangeKernel(queue, kernel, 3, 0, global_work_size, NULL, 0, 0, NULL);

Please let me know how to choose global_work_size , local_work_size.

Implementing nested loop in OpenCL?

0 Answers