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

Question

所以我正在玩 CUDA C++ 编程。我试图在设备上创建一个 class，它是运行，像这样：

class DeviceClass {
    int deviceInt = 5;

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

现在，除非我在示例中遗漏了某些内容，否则这实际上会运行在设备上正确。我可以使用 __global__ 方法中的 new 和运行中的 DeviceFun() 方法对其进行初始化。

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

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

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

Answer 1

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

对象模型对何时可以使用具有 non-default 构造函数或虚拟成员的对象有一些硬性限制，但可以在 __global__、__shared__、__constant__ 或 local/stack 内存。

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

所以可以定义一个简单的 class 和一个像这样的空构造函数：

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]);
    }
}

[警告：此代码从未在编译器附近使用过，不能保证有效]

所以在内核中，你有

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

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

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

CUDA device C++ class, in what memory type are object variables stored and is it possible to change it?

c++

memory

cuda

gpgpu

class