Wierd behavior of pointers to structs in constant memory.

Is there something special I have to do in order to send pointers to structs in constant memory?

I have a memory object created with CL_MEM_READ_ONLY and filled using a standard clEnqueueWriteBuffer.
I also have a .cl-file that contains two functions and both of these have “__constant Parameters *params” in the argument list.

The problem is that I can only run the kernel defined lastly in the file. Calling the first one causes a crash in clEnqueueNDRangeKernel, trying to access address 0. Swapping the locations in the code makes the other kernel runnable, and the crash happens if the one that previously worked is enqueued.

Changing __constant to __global in the second kernel causes the first to work again.

Has anyone had similar problems? Any ideas on what I may be doing wrong?

I’m using NVIDIAS OpenCL and would have asked in the NVIDIA forums since I believe this to be a bug in their implementation, but their forums are down for the next few days.

I have a related problem to __constant memory pointer using Apple’s implementation on a MacBook Pro (which I suspects uses NVIDIA’s compiler for the GPU).

In the kernel, the __constant memory variable (denote it buf[] here) cannot index properly.
basically the following problem:

for(loopidx = 0; loopidx < MAX; loopidx++) {
x = … * buf[loopidx] * …;
// do other stuff
}

works but:
x = …* buf[0] *…;
for(loopidx = 0; loopidx < MAX; loopidx++) {
// do other stuff
x = … * buf[loopidx + 1] * …;
}
only returns zeros for the buf content (except for the first assignment using buf[0])

I belive that NVIDIA has some problems with their constant memory pointer handling.
Even if this does not help you in your problem, you (and NVIDIA if they read this) knows that there seems to be several problems with __constant memory in their implementation.

BR
/Patrik

It is my understanding that the pointer access to constant memory under OS X is a known bug that is being looked into. (I have no idea when a fix might appear.) It seems like the problem is that constant values are stored in a different address space on the hardware which means access to them has to be handled carefully by the backend. I have no idea if this is what is causing the problem with the Nvidia drivers, but since the kernels work by themselves I doubt that is the case. It sounds more like an Nvidia CL compiler bug.