Constant memory execution error

Hello everybody,

I’ve been facing a strange problem with constant memory. Look at the header of my kernel:


__kernel void Lvo 
						( 
						 __global int* Rt,
 						 __constant int* St,
 						 __constant int* Me,
 						 __constant int* Ie,
 						 __global int* sl,
 						 __global char* sqs,
 						 __constant int* Q,
 						 __constant int* amount,
 						 __global int* A,
 						 __global int* B,
 						 __global int* C,
 						 __global int* D,
 						 __global int* E,
 						 __global int* F,
  						 __global int* J,
 						 __global int* try_me,
 						 __global int* s,
 						 __global int* ss,
 						 __global int* i,
 						 __global int* k
 						)
{
...
}

This works just fine, but when I try to use the Rt argument in constant memory (by putting __constant instead of __global) I sometimes get a segmentation fault, sometimes a malloc corruption and sometimes a double-linked list error.

My Rt argument has a 2KB size and is the biggest constant structure, so the size limit isn’t a problem. Besides, it is used 9 times inside the kernel (always being read).

Can anybody help me or at least suggest something? I am tired of trying on my own.

Thanks
Samuel

I’m not sure what we can say other than “it looks like a bug in the compiler; please contact your vendor”.

The issue could be related to the number of __constant arguments rather than their size. Each literal value in the kernel source counts as one additional __constant argument and the limit is 8 per kernel IIRC.

Unfortunately this problem happens even if I put Rt as the only constant argument. You may be right, it looks like a bug in the compiler, but I’ll keep trying to find a way to solve this. Thanks.

Any other suggestions are welcome.

I still haven’t found the solution for my issue and searching on OpenCL reference guide I’ve found the following:

"A pointer to address space A can only be assigned to a pointer to the same address space A. Casting a pointer to address space A to a pointer to address space B is illegal. "

Does it mean that a variable allocated on private address space must not receive a value from a pointer to the constant address space?

Thank you
Samuel

Does it mean that a variable allocated on private address space must not receive a value from a pointer to the constant address space?

No, that’s not a problem. I’ll explain it with the following example:


__kernel void foo(__global int *a, __local int *b, __constant int *c)
{
    // Illegal
    a = b;

    // Illegal
    a = c;

    // Illegal
    b = c;

    // Legal
    a[0] = b[0];

    // Legal
    a[0] = c[0];

    // Legal
    b[0] = c[0];
}