CUDA 设备 C++ 类，对象变量存储在什么内存类型中，是否可以更改它？

Question

所以我在玩 CUDA C++ 编程。我试图创建一个在设备上运行的类，如下所示：

class DeviceClass {
    int deviceInt = 5;

    __DEVICE__ void DeviceFunc()
    {
        printf("Value of device var: %d\n", deviceInt);
    }
}

现在，除非我错过了示例中的某些内容，否则它实际上会在设备上正确运行。我可以使用 __global__ 方法中的 new 对其进行初始化，然后从那里运行 DeviceFun() 方法。

变量deviceInt存储在设备的哪个内存中？

我可以强制将其分配到默认内存以外的另一种内存类型中吗？例如，我可能想将非常大的数据数组放入全局内存中，而将其他一些更本地的东西放入更快的内存中。

像这样进行 CUDA 设备类设计是个好主意，还是我以后会遇到更大的设计问题？

Answer 1

对象存储在定义它们时指定的任何内存空间中（这也是为什么在类或结构定义中使用内存空间规范是非法的）。

对象模型对何时可以使用具有非默认构造函数或具有虚拟成员的对象有一些硬性限制，但可以在__global__、__shared__、__constant__ 或本地/堆栈内存。

“现代”（计算能力 >= 2.0）GPU 支持 ABI，这意味着编译时的静态指针自省也可以正常工作。

因此可以像这样定义一个带有空构造函数的普通类：

struct foo
{
    float x, y;

    __device__ float f() const
    {
        return x*x + y*y;
    };

    __device__ bool operator< (const foo& x) const
    {
        return (f() < x.f());
    };
};

然后像这样在设备代码中使用它：

__device__ foo foo_min(const foo& x, const foo& y)
{
    return (x < y) ? x : y;
}

__global__ void kernel(foo *indata, foo *outdata, int N)
{
    int idx = threadIdx.x + blockIdx.x  *blockDim.x;
    foo localmin = indata[idx];
    for(; idx < N; idx += blockDim.x * gridDim.x)
    {
        localmin = foo_min(localmin, indata[idx]);
    };

    /* simple shared memory reduction */
    __shared__ foo buff[128];
    buff[threadIdx.x] = localmin;
    __syncthreads();

    if (threadIdx.x < 64) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+64]);
    }
    __syncthreads();

    if (threadIdx.x < 32) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+32]);
    }
    __syncthreads();
    if (threadIdx.x < 16) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+16]);
    }
    __syncthreads();

    if (threadIdx.x < 8) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+8]);
    }
    __syncthreads();

    if (threadIdx.x < 4) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+4]);
    }
    __syncthreads();

    if (threadIdx.x < 2) {
        buff[threadIdx.x] = foo_min(buff[threadIdx.x], buff[threadIdx.x+2]);
    }
    __syncthreads();

    if (threadIdx.x == 0) {
        outdata[blockIdx.x] = foo_min(buff[0], buff[1]);
    }
}

[警告：此代码从未靠近编译器，不保证可以工作]

所以在内核中，你有

1. 指向全局内存中 foo 的指针
2. foo 对象的静态共享内存数组
3. 线程本地 foo 实例

并且可以将它们全部传递给设备函数foo_min，而无需在代码中做任何特殊的事情，并且编译器可以透明地理解和处理这些情况。