MATLAB: Cuda PTX kernel function boolean argument not supported

Question

Hello,

i have a cuda kernel function:

__global__ void backprop_float(float* holo, const float* inputs, int spotn, const int holoWidth, const int holoHeight, bool addcorr)

when i try to create the kernel object in matlab:

backpropk=parallel.gpu.CUDAKernel('backprop_kernels.ptx', 'backprop_kernels.cu','backprop_float');

I get the following error:

__Error using iCheckPTXEntryAgainstCProto (line 420)
Found invalid mapping type < logical, .u8 >__

if i change the type of the last function parameter bool addcorr to int, matlab creates the kernel object succesfully.

(i use matlab r2012a, vstudio2008, cuda toolkit v4.2 card is a gtx 560ti)

Answer 1

Thanks for providing the compile line. It seems that recent versions of NVCC switch from mapping bools as "s8" to "u8" according to the target architecture ("s8" for <sm_20, "u8" for >=sm_20). parallel.gpu.CUDAKernel can't cope with the latter. That's a bug.

Whilst we work to get a permanent fix for this, I can think of three work-arounds. None of these are great, but should let you move on:

1. Edit the resulting PTX to accept "s8" instead of "u8". You would have to be careful that this doesn't cause problems elsewhere in the PTX entry.
2. Change the CUDA kernel to take the boolean arguments as unisgned chars (or some other unsigned 8-bit type). Since the PTX is treating booleans as an 8-bit type anyway, this is very close to the correct behaviour.
3. Compile targetting sm_13, since that causes NVCC to emit "s8". There are performance reasons why you probably don't want to do this.

An example of (2), where "flag" used to be of type "bool":

myKernel.cu:
__global__ 
void myKernel(float* out, float const* in, unsigned char const flag) 
{
    if (flag) {
        *out = -(*in);
    } else {
        *out = *in;
    }
}
>> k = parallel.gpu.CUDAKernel('myKernel.ptx', 'myKernel.cu');
>> out = feval(k2,0,1,true)

As and when a permanent fix is available I will update this answer.