CLH_ERROR_NO_BINARY_FOR_GPU for image2d_t

This is on MacOS 10.6.1. Here’s my kernel:

__kernel void Convolve(
	__read_only __global image2d_t img, 
	__write_only __global image2d_t newImg,
	__constant uint *filter, const int filterWidth, const int filterHeight,
	const int w, const int h)
{
	__private const sampler_t sampler = 
		CLK_NORMALIZED_COORDS_TRUE |
		CLK_ADDRESS_CLAMP |
		CLK_FILTER_NEAREST;

	uint4 val = read_imageui(img, sampler, (int2)(0, 0));
}

The problem is that the build log shows this:

Build Log:
Error while compiling the ptx module: CLH_ERROR_NO_BINARY_FOR_GPU
PTX Info log: 
PTX Error log: 

And that’s no fun! Commenting out the read_imageui line causes a successful compilation, but is hardly a solution. The specs say specifically (6.11.8) “Image memory objects that are being read by a kernel should be declared with the __read_only qualifier.” Removing the __read_only qualifier from img also doesn’t help. Removing the __global qualifier does help, but why should that be? Global is the only space that could possibly hold the image because of its size.

Should I file a bug with Apple? Or is it a case of ur doin it wrong?

Thanks,

–Rob

Rob,
I would agree that you should not get this error because of the global option. Either you should get a compiler error or it should just work with global. (You don’t need the “__” before private/global etc.) The no binary for GPU error suggests that the GPU’s compiler failed to compile for some reason, which would be consistent with why you didn’t get a compiler error string. (Those tend to come from the front-end and not the GPU back-end.) If removing the global makes it work then I would file a bug against Apple saying that putting in global fails to compile and doesn’t generate an compiler error. I suspect this works fine on the CPU, though.

Thanks, I have filed bug ID 7372367. Let’s hope we see some action from Apple on this!

Here is what Apple had to say about this issue. Apparently they are looking to Khronos to update the spec. They effectively say that image2d_t and image3d_t (as well as sampler) are structures, and you don’t get to specify where the data in the structure gets stored. In addition, regardless of where the data in the structure is stored, the image content is always stored in global space.

The image data types are opaque data structures. The contents of the image are allocated in __global address space but the variable defined to be of type image2d_t or image3d_t is not. Since these are not pointers, it does not make sense to use the __global qualifier for the image data type declaration. The spec needs to be clarified so that image type behave the same as sampler types i.e. the address space qualifier should not be specified.