Computing on variable length arrays in OpenCL

Question

I am using OpenCL (Xcode, Intel GPU), and I am trying to implement a kernel that calculates moving averages and deviations. I want to pass several double arrays of varying lengths to the kernel. Is this possible to implement, or do I need to pad smaller arrays with zeroes so all the arrays are the same size?

I am new to OpenCL and GPGPU, so please forgive my ignorance of any nomenclature.

Answer 1

You can pass to the kernel any buffer, the kernel does not need to use it all. For example, if your kernel reduces a buffer you can query at run time the amount of work items (items to reduce) using get_global_size(0). And then call the kernel with the proper parameters.

An example (unoptimized):

__kernel reduce_step(__global float* data)
{
    int id = get_global_id(0);
    int size = get_global_size(0);
    int size2 = size/2;
    int size2p = (size+1)/2;
    if(id<size2) //Only reduce up to size2, the odd element will remain in place
       data[id] += data[id+size2p];
}

Then you can call it like this.

void reduce_me(std::vector<cl_float>& data){
    size_t size = data.size();

    //Copy to a buffer already created, equal or bigger size than data size
    // ... TODO, check sizes of buffer or change the buffer set to the kernel args.
    queue.enqueueWriteBuffer(buffer,CL_FALSE,0,sizeof(cl_float)*size,data.data());

    //Reduce until 1024
    while(size > 1024){
        queue.enqueueNDRangeKernel(reduce_kernel,cl::NullRange,cl::NDRange(size),cl::NullRange);
        size /= 2; 
    }

    //Read out and trim
    queue.enqueueReadBuffer(buffer,CL_TRUE,0,sizeof(cl_float)*size,data.data());
    data.resize(size);
}