CUDA学习(六十八)

简介:

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语言特性表:
1
2
3
4
C ++ 14语言特性:
下表列出了已被C ++ 14标准接受的新语言功能:
5
限制:
主机编译器扩展:
设备代码不支持主机编译器特定的语言扩展。 主机和设备代码都不支持__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__用途生成诊断。
u_3741838079_1934910147_fm_27_gp_0

目录
相关文章
|
并行计算 C语言 编译器
|
并行计算 索引
|
并行计算 API C语言
|
并行计算 API 编译器
CUDA学习(六十五)
很早之前就发现云栖社区的编辑器有一个Bug,往草稿箱存博客,当草稿箱博客数超过十篇时,无法再选择十篇前的博客进行编辑
2403 0
|
存储 并行计算 C语言
|
并行计算 前端开发
|
并行计算 API
|
存储 并行计算 API