C / C ++语言支持:
如使用NVCC编译中所述,使用nvcc编译的CUDA源文件可以包含主机代码和设备代码的混合。 CUDA前端编译器旨在模拟主机编译器相对于C ++输入代码的行为。 输入源代码根据C ++ ISO / IEC 14882:2003,C ++ ISO / IEC 14882:2011或C ++ ISO / IEC 14882:2014规范进行处理,CUDA前端编译器旨在模拟任何主机编译器与 ISO规范。 此外,受支持的语言受以下限制:
C ++ 11语言功能和C ++ 14语言功能分别为C ++ 11和C ++ 14功能提供支持矩阵。 限制列出了语言限制。 多态函数包装和实验特征:扩展Lambda描述了附加特征。 代码示例提供了代码示例。
C ++ 11语言特性:
下表列出了已被C ++ 11标准接受的新语言功能。 “Proposal”列提供了描述该功能的ISO C ++委员会提案的链接,而“Available in nvcc(设备代码)”列指示包含此功能实现的第一版nvcc(如果已实施 )为设备代码。
C ++ 11语言特性表:
C ++ 14语言特性:
下表列出了已被C ++ 14标准接受的新语言功能:
限制:
主机编译器扩展:
设备代码不支持主机编译器特定的语言扩展。 主机和设备代码都不支持__float128
和__float80
内建类型。
预处理符号:__CUDA_ARCH__
1.下列实体的类型签名不取决于__CUDA_ARCH__
是否被定义,或者取决于__CUDA_ARCH__
的特定值:
-
__global__
函数和函数模板 -
__device__
和__constant__
变量 - 纹理和表面
例子:
#if !defined(__CUDA_ARCH__)
typedef int mytype;
#else
typedef double mytype;
#endif
__device__ mytype xxx; // error: xxx's type depends on __CUDA_ARCH__
__global__ void foo(mytype in, // error: foo's type depends on __CUDA_ARCH__
mytype *ptr)
{
*ptr = in;
}
2.如果__global__
函数模板被实例化并从主机启动,那么函数模板必须使用相同的模板参数实例化,而不管__CUDA_ARCH__
是否被定义以及__CUDA_ARCH__
的值如何。
例子:
__device__ int result;
template <typename T>
__global__ void kern(T in)
{
result = in;
}
__host__ __device__ void foo(void)
{
#if !defined(__CUDA_ARCH__)
kern << <1, 1 >> >(1); // error: "kern<int>" instantiation only
// when __CUDA_ARCH__ is undefined!
#endif
}
int main(void)
{
foo();
cudaDeviceSynchronize();
return 0;
}
3.在单独的编译模式下,是否存在具有外部链接的函数或变量的定义不应取决于__CUDA_ARCH__
是定义的还是取决于__CUDA_ARCH__
的特定值。
例子:
#if !defined(__CUDA_ARCH__)
void foo(void) { } // error: The definition of foo()
// is only present when __CUDA_ARCH__
// is undefined
#endif
4.在单独编译中,__CUDA_ARCH__不能用于标题中,以便不同的对象可以包含不同的行为。 或者,必须保证所有对象都会为相同的compute_arch编译。 如果在头中定义了弱函数或模板函数,并且其行为取决于__CUDA_ARCH__,那么如果对象是针对不同的计算拱编译的,则对象中该函数的实例可能会发生冲突。
例子:
template<typename T>
__device__ T* getptr(void)
{
#if __CUDA_ARCH__ == 200
return NULL; /* no address */
#else
__shared__ T arr[256];
return arr;
#endif
}
然后,如果a.cu和b.cu都包含a.h并为相同类型实例化getptr,并且b.cu期望非NULL地址,并且编译时使用:
nvcc –arch=compute_20 –dc a.cu
nvcc –arch=compute_30 –dc b.cu
nvcc –arch=sm_30 a.o b.o
在链接时,只使用getptr的一个版本,所以行为取决于选择哪个版本。 为避免这种情况,a.cu和b.cu必须为同一个计算拱编译,或者__CUDA_ARCH__
不应该用在共享头文件函数中。
编译器不保证将为上述不支持的`__CUDA_ARCH__用途生成诊断。