Cuda PTX kernel function boolean argument not supported?

1 次查看（过去 30 天）
Gaszton
Gaszton 2012-9-11
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)
  2 个评论
Ben Tordoff
Ben Tordoff 2012-9-12
I've just tried compiling a variety of kernels for a variety of target architectures and with multiple toolkit versions. Boolean arguments always seem to get mapped to "s8" (signed 8-bit integers) for me. For example:
__global__
void myKernel(float* out, const float* in, bool flag)
becomes
.entry _Z8myKernelPfPKfb (
.param .u64 __cudaparm__Z8myKernelPfPKfb_out,
.param .u64 __cudaparm__Z8myKernelPfPKfb_in,
.param .s8 __cudaparm__Z8myKernelPfPKfb_flag)
This is the mapping that MATLAB is expecting. It looks like you are getting "u8" (unsigned 8-bit integer) as the boolean type. Could you show me the exact NVCC command you used to generate this and also add the PTX ".entry" corresponding to the CUDA function so that I can work out why you're getting this different type mapping?
Thanks.
Gaszton
Gaszton 2012-9-12
编辑：Gaszton 2012-9-12
Hello, i use the following command for ptx compilation:
cmd /k nvcc -I "C:\Program Files\Microsoft Visual Studio 9.0\VC\include" -gencode=arch=compute_20,code=\"sm_21,compute_20\" -use_fast_math -ptx backprop_kernels.cu
the ptx .entry:
.entry _Z14backprop_floatPfPKfiiib( .param .u32 _Z14backprop_floatPfPKfiiib_param_0, .param .u32 _Z14backprop_floatPfPKfiiib_param_1, .param .u32 _Z14backprop_floatPfPKfiiib_param_2, .param .u32 _Z14backprop_floatPfPKfiiib_param_3, .param .u32 _Z14backprop_floatPfPKfiiib_param_4, .param .u8 _Z14backprop_floatPfPKfiiib_param_5 )
If i do not specify the compute architecture version, matlab is able to map the complied ptx, and i see that the entry is s8 in that case.
Thank you!

请先登录，再进行评论。

请先登录，再回答此问题。

采纳的回答

Ben Tordoff
Ben Tordoff 2012-9-13
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.
  2 个评论
Gaszton
Gaszton 2012-9-13
Thank you for your help, i already changed the type from boolean to char, and everything runs fine now.
Ben Tordoff
Ben Tordoff 2012-9-18
A patch for this problem in R2011b and R2012a is now online here:
http://www.mathworks.com/support/bugreports/814923
Please can you try it and let me know if you still have issues?

请先登录，再进行评论。

更多回答（0 个）

类别

在 Help Center 和 File Exchange 中查找有关 GPU Computing 的更多信息

标签

Community Treasure Hunt

Find the treasures in MATLAB Central and discover how the community can help you!

Start Hunting!

Translated by