clBuildProgram fails when write_imageui is used

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi!

I’m trying to build a simple kernel, but it fails when I use write_imageui. The kernel should just copy an image to another. I know that there isn’t probably any sense in making such kernel, but I’m just trying to learn the language. Anyway, here is the kernel:

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_NEAREST;

__kernel void copy(image2d_t src, image2d_t dst)
{
int2 pos;
uint4 input;
int x = get_global_id(0);
int y = get_global_id(1);

pos.x = x;
pos.y = y;

input = read_imageui(src, sampler, pos);
write_imageui(dst, pos, input);

}

If I put it like above, clBuildProgram gives me error code -42 (btw how can I determine what error that is?). But if I comment the last line out it compiles just fine. Any idea what is going on?

-hnyk

2

You could check error numbers in your system CL/cl.h header file; specifically, -42 is CL_INVALID_BINARY (and clBuildProgram() documentation is stating that this one is returned “… if program is created with clCreateWithProgramBinary and
devices listed in device_list do not have a valid program binary loaded”). I ran your kernel quickly through AMD CL compiler (easier to check for syntax errors with an existing external program), and some errors get reported, so I’d suggest you use clGetProgramBuildInfo() to check what NVIDIA compiler is going to report; also I’d check for potential errors of OpenCL calls before this one (it doesn’t seem logical for above error to get reported, as you provided the kernel source, so I guess you used clCreateProgramWithSource() when creating the program object).

3

Thanks for the answer. Every command before clBuildProgram() executes without an error and the build log gives this:

Build Log:

ptxas ptx input, line 114; error : Unknown symbol ‘copy_param_0’

: Retrieving binary for ‘anonymous_jit_identity’, for gpu=‘sm_11’, usage mode=’

So the problem might be in input parameters? Any ideas what should I try next?

-hnyk

4

Okay, I fixed it. I just added __write_only and __read_only specifiers to input parameters.