c++ - 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 支持，这意味着编译时的静态指针自省(introspection)也能正常工作。

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

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，而无需在代码中做任何特殊的事情，并且编译器可以透明地理解和处理这些情况。