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

【问题标题】：What are the real C++ language constructs supported by CUDA device code?CUDA 设备代码支持哪些真正的 C++ 语言结构？
【发布时间】：2011-02-27 00:57:44
【问题描述】：

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 随机生成类。我唯一的猜测是，由于__device__ 标记函数的自动内联，我能够做到这一点，但这并不能解释成员变量的处理。

您是否曾经在相同条件下使用过这些功能，或者您能向我解释一下为什么我的 CUDA 代码会这样吗？参考指南有什么问题吗？

【问题讨论】：

标签： c++ class cuda gpu gpgpu

【解决方案1】：

官方上，CUDA 不支持 2.0 之前的设备上的类。

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

这意味着，您可以使用以下 C++ 结构：

可见性（公共/受保护/私有）
非虚拟继承
整个模板编程和元编程（直到您发现 nvcc 错误；从 3.2 版开始就有不少）
构造函数（对象在__共享__内存中声明时除外）
命名空间

您不能使用以下内容：

new & delete 运算符（我相信设备 >=2.0 可以做到这一点）
虚方法（需要跳转到变量地址）
函数递归（需要函数调用）
例外情况

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

【讨论】：

您的回答带来了有关限制的有趣观点。在使用计算能力的设备时，编译时的必要性似乎是要考虑的重点
在调查了问题后，我无法（也可能不会）找到比您更好的答案。因此，我依靠您的经验 :) 无论如何，我现在拥有一个 C2050 \o/，因此我可以比较我的执行以确保我的代码保持可移植性。谢谢你的回答@CygnusX1

【解决方案2】：

一些 C++ 类功能会起作用，但是编程指南基本上是说它不完全受支持，因此并非所有 C++ 类功能都会起作用。如果你能做你想做的事，那么你应该继续！

【讨论】：

好的，但我需要确切地知道什么会起作用。我打算在库中使用这些功能，所以如果它们不安全，我就不能依赖它们。我们可以很容易地想象，这段代码在某些情况下会造成严重破坏。我真正需要的是我应该做与不做之间的确切界限。