How to pass array and int to opencl kernel in same time-CodePudding

I've an openCL code to do some calculate. Can OpenCL pass array and int to kernel in the same time? I want to pass 'myint' to kernel not via buffer.

    int myint = 100;
    cl_mem memObjects[3] = {0, 0, 0};

    memObjects[0] = clCreateBuffer(opencl_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   sizeof(int) * ARR_SIZE, data, NULL);
    memObjects[1] = clCreateBuffer(opencl_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                                   sizeof(int), new int[]{myint}, NULL);
    memObjects[2] = clCreateBuffer(opencl_context, CL_MEM_READ_WRITE,
                                   sizeof(float) * ARR_SIZE, NULL, NULL);

    opencl_errNum = clSetKernelArg(opencl_kernel, 0, sizeof(cl_mem), &memObjects[0]);
    opencl_errNum |= clSetKernelArg(opencl_kernel, 1, sizeof(cl_mem), &memObjects[1]);
    opencl_errNum |= clSetKernelArg(opencl_kernel, 2, sizeof(cl_mem), &memObjects[2]);
    
    
    //Kernel code
    const char *opencl_code = "__kernel void opencl_code(__global const int *src,\n"
                            "                                __global const int *value,\n"
                            "                                __global float *result) {\n"
                            "    int gid = get_global_id(0);\n"
                            "    result[gid] = src[gid] - value[0];\n"
                            "\n"
                            "}";

Want to change be like this.

    const char *opencl_code = "__kernel void opencl_code(__global const int *src,\n"
                            "                                __global const int value,\n"
                            "                                __global float *result) {\n"
                            "    int gid = get_global_id(0);\n"
                            "    result[gid] = src[gid] - value;\n"
                            "\n"
                            "}";

CodePudding user response：

Yes you can pass either 1-dimensional arrays or constants as kernel arguments. The syntax is as follows:

int myint = 100;
cl_mem memObjects[2] = {0, 0};
memObjects[0] = clCreateBuffer(opencl_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int) * ARR_SIZE, data, NULL);
memObjects[1] = clCreateBuffer(opencl_context, CL_MEM_READ_WRITE, sizeof(float) * ARR_SIZE, NULL, NULL);

opencl_errNum  = clSetKernelArg(opencl_kernel, 0, sizeof(cl_mem), &memObjects[0]); // pass an array as kernel parameter
opencl_errNum |= clSetKernelArg(opencl_kernel, 1, sizeof(int), (void*)&myint); // pass a constant as kernel parameter
opencl_errNum |= clSetKernelArg(opencl_kernel, 2, sizeof(cl_mem), &memObjects[1]); // pass an array as kernel parameter
    
// note: the second parameter is just "const int value", without the "__global" keyword or "*" pointer
const char *opencl_code = "__kernel void opencl_code(__global const int *src,\n"
                          "                          const int value,\n"
                          "                          __global float *result) {\n"
                          "    int gid = get_global_id(0);\n"
                          "    result[gid] = src[gid] - value;\n" // note: result is type float, and you write an int to it here
                          "}\n";

To make this all a lot easier, I created this OpenCL-Wrapper. With that, you can directly pass arrays (Memory objects) or numbers as kernel parameters, and the C code would look like this:

int main() {
    Device device(select_device_with_most_flops()); // compile OpenCL C code for the fastest available device
    int myint = 100;
    Memory<int> src(device, ARR_SIZE); // allocate memory on both host and device
    Memory<float> result(device, ARR_SIZE);
    Kernel opencl_kernel(device, ARR_SIZE, "opencl_code", src, myint, result); // kernel that runs on the device
    // initialize src, for example with "src[5] = 3;"
    src.write_to_device(); // copy data from host memory to device memory
    opencl_kernel.run(); // run kernel
    result.read_from_device(); // copy data from device memory to host memory
    // read result with for example "float test = result[5];"
    // ...
    return 0;
}