CUDA设备代码支持哪些真正的C++语言结构？

Question

3.2版本的CUDA文档的附录D是指CUDA设备代码中的C++支持.
很明显,CUDA支持"计算能力2.x设备的类".但是,我正在使用计算能力1.1和1.3的设备,我可以使用此功能!

例如,此代码有效:

// class definition voluntary simplified
class Foo {
  private:
    int x_;

  public:
    __device__ Foo() { x_ = 42; }
    __device__ void bar() { return x_; }
};


//kernel using the previous class
__global__ void testKernel(uint32_t* ddata) {
    Foo f;
    ddata[threadIdx.x] = f.bar(); 
}

我也可以使用广泛的库,如Thrust :: random random generation classes.我唯一的猜测是,由于__device__标记函数的自动内联,我能够这样做,但这并不能解释成员变量的处理方式.

您是否曾在相同的条件下使用过这些功能,或者您能解释一下为什么我的CUDA代码会以这种方式运行吗？参考指南中有什么问题吗？

Answer 1

最初,CUDA不支持2.0之前的设备上的类.

实际上,根据我的经验,只要可以在编译时解析功能,就可以在所有设备上使用所有C++功能.2.0之前的设备不支持函数调用(所有函数都是内联的),并且没有程序跳转到变量地址(仅跳转到常量地址).

这意味着,您可以使用以下C++构造:

• 可见性(公共/受保护/私人)
• 非虚拟继承
• 整个模板编程和元编程(直到你对nvcc错误进行调查;从版本3.2开始,它们中有相当一部分)
• 构造函数(在__ shared __ memory中声明对象时除外)
• 命名空间

您不能使用以下内容:

• 新的和删除操作符(我相信设备> = 2.0可以做到这一点)
• 虚方法(需要在变量地址跳转)
• 函数递归(需要函数调用)
• 例外

实际上,"CUDA编程指南"第D.6章中的所有示例都可以编译设备<2.0