Failed to compile kernel with read_imageui() - "unknown symbol"

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

Attached is the simple program with even simpler kernel, which tries to read data from given image. For those who do not want to download the program itself, the kernel looks like:

__kernel void test(read_only image3d_t img)

{

	uint4 image_data = read_imageui(img,

		(sampler_t)(CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP | CLK_NORMALIZED_COORDS_FALSE),

		(int4)(0, 0, 0, 0));

}

The whole program was created from one of SDK examples and has to be built analogously.

The program creates context, compiles and builds this kernel. Under OSX 10.6.2 it works fine. Under Ubuntu 9.10 x64 + nv drivers 195.17 clBuildProgram() fails and output is the following:

-----------------------------------------------------------

Build Log:

ptxas application ptx input, line 30; error   : Unknown symbol 'test_param_32741'

: Retrieving binary for 'anonymous_jit_identity', for gpu='sm_11', usage mode=''

-----------------------------------------------------------

Can anyone please tell me whether there is a way to make this kernel work under Linux too? (of course, it is not the real kernel I use, just a simplified bug-reproduction case)
oclTest.tar.gz (3.2 KB)

2

Try the following (untested, but I did have a similar issue where this worked):

__kernel void test(read_only image3d_t img)

{

	sampler_t sampler = CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP | CLK_NORMALIZED_COORDS_FALSE;

	uint4 image_data = read_imageui(img, sampler, (int4)(0, 0, 0, 0));

}

It seems that NVidia’s (Linux) tools cannot deal with function arguments of type sampler_t that are not a sampler_t variable.

3

Thank you very much, it works!

4

Just a small related issue… the opencl spec says read_only is default for image??_t kernel args, but /w linux drivers it will not compile if you omit the qualifier and still read from the image via e.g. read_imagef . it will give a similar cryptic ptxas log error about a undefined symbol. i found this thread in the search for a solution to that problem and i thought it might be nice to know for others