OpenCL kernel question

Question

I have taken the Kernel from the great OpenCL SpMV article for AMD by Bryan Catanzaro. I have given it a toy problem where the input is A= [0 0 0 6 1 3 5 7 2 4 0 0] offsets= [-3 0 2] x= [1 2 3 4] and the output y should be [7 22 15 34]

Here is the kernel:

__kernel
void dia_spmv(__global float *A, __const int rows,
              __const int diags, __global int *offsets,
              __global float *x, __global float *y) {        
    int row = get_global_id(0);
    float accumulator = 0;
    for(int diag = 0; diag < diags; diag++) {
        int col = row + offsets[diag];
        if ((col >= 0) && (col < rows)) {
            float m = A[diag*rows + row]; 
            float v = x[col];
            accumulator += m * v;
        }
    }
    y[row] = accumulator;
}

After loading and writing the input arguments I execute the kernel like this:

size_t global_work_size;
global_work_size = 4; 

err = clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL, &global_work_size,NULL, 0, NULL, NULL);
    err = clFinish(cmd_queue);

And I get the correct result when I read y back from gpu memory. I.e. I get y = [7 22 15 34]

I am new to OpenCL (and GPGPU in general) so I want to try and understand how to extend the problem correctly for much larger matrices of arbitrary dimension. So lets say I have 1000 000 rows. What should I set global_work_size to be? And should I set local_work_size or should I leave it as NULL?

Rick-Rainer Ludwig Rick-Rainer Ludwig · Accepted Answer · 2011-06-27T08:03:28

To use the kernel for arbitrary matrix sizes you should think about the problem and rewrite the kernel. The issue is the limited memory size of the GPU and limited size for a single buffer. You can get the maximum size for a buffer with clGetDeviceInfo and CL_DEVICE_MAX_MEM_ALLOC_SIZE.

You need to split your problem into smaller pieces. Calculate them separately and merge the results afterwards.

I do not know the problem above and can not give you any hint which helps you to implement this. I can only give you the general direction.

OpenCL kernel question

1 Answers