【发布时间】：2016-01-30 19:53:06
【问题描述】：
我有以下最小的非工作示例：
#include <cstdio>
#include <cuda_runtime_api.h>
/* this declaration would normally be in a header, but it doesn't matter */
template<class T_PREC> __global__ void testKernel( T_PREC );
template<class T_PREC> __global__ void testKernel( T_PREC const x )
{
printf( "%f", x );
}
int main()
{
printf("calling kernel...");
testKernel<<<1,1>>>( 3.0f );
cudaDeviceSynchronize();
printf("OK\n");
return 0;
}
我用它编译和运行
nvcc simple.cu && ./a.out
输出是：
calling kernel...
意味着程序在它既不能打印“OK”也不能打印浮点数之前就崩溃了。这不是分段错误，所以我无法回溯任何东西。我正在使用 CUDA 7.0.27。在gdb 中运行时，消息是：
[Inferior 1 (process 27899) exited with code 01]
上面的例子有四个原因：
-
不要使用 CUDA：
template<class T_PREC> void testKernel( T_PREC ); template<class T_PREC> void testKernel( T_PREC const x ) { printf( "%f", x ); } int main() { printf("calling kernel..."); testKernel( 3.0f ); cudaDeviceSynchronize(); printf("OK\n"); return 0; } -
不要使用模板：
__global__ void testKernel( float ); __global__ void testKernel( float const x ) { printf( "%f", x ); } -
省略声明（如果我需要从库中提供它，则不是一个选项）
//template<class T_PREC> void testKernel( T_PREC ); -
不要在声明中省略
const限定符：template<class T_PREC> __global__ void testKernel( T_PREC const );这是最合理的选择，但我不明白为什么必须这样做。在普通 C++ 中，按值调用的
const限定符不应更改函数签名。即使它这样做了，它也不应该链接并且只会在执行时崩溃。那么为什么 CUDA 的行为会有所不同，为什么只使用模板呢？
附加检查：
由于汇编代码本身对我来说太难了，我查看了创建的可执行文件：
nvcc sameTypeQualifier/main.cu -o same.o
[no warning output whatsoever]
nvcc diffTypeQualifier/main.cu -o diff.o
diffTypeQualifier/main.cu: In instantiation of ‘void __wrapper__device_stub_testKernel(T_PREC* const&) [with T_PREC = float]’:
diffTypeQualifier/main.cu:8:45: required from ‘void testKernel(T_PREC*) [with T_PREC = float]’
diffTypeQualifier/main.cu:15:67: required from here
diffTypeQualifier/main.cu:7:86: warning: unused parameter ‘x’ [-Wunused-parameter]
template<class T_PREC> __global__ void testKernel( T_PREC * const x )
^
diff <(nm -C same.o | sed 's/^[0-9a-f]*//') <(nm -C diff.o | sed 's/^[0-9a-f]*//')
389a390
> t void __wrapper__device_stub_testKernel<float>(float*&)
419c420
< t __sti____cudaRegisterAll_39_tmpxft_000050c8_00000000_9_main_cpp1_ii_main()
---
> t __sti____cudaRegisterAll_39_tmpxft_0000511c_00000000_9_main_cpp1_ii_main()
cudaRegisterAll 中的十六进制数字在相同源代码的两次编译之间甚至不同，因此可以忽略。出于某种原因，非工作示例具有额外的包装功能。
【问题讨论】：
-
看来我的观察（CUDA 7.5，MSVS 2010）与你的略有不同：只要模板声明和模板定义在使用
const时匹配（即两者都有它或两者都没有），内核被调用并按预期打印传递的值。我没有足够的 C++ 语言律师来说明这里应该发生什么；可能这也可能是 CUDA 编译器中的一个错误，但我不愿意这么说，因为我不确定当const在声明和定义之间不匹配时代码是否会调用未定义的行为。 -
@njuffa：我认为这些答案是相关的：stackoverflow.com/questions/18215686/… 顺便说一句：
float *上的__restrict__关键字也会出现同样的问题。只有当声明也有它时它才有效。 -
请注意，问题指的是 C，而 CUDA 是 C++ 的变体。虽然 C 和 C++ 密切相关，但它们之间存在许多差异，