c++ - 如何在cuda上创建全局可访问的变量？

Question

这是一个相当复杂的问题，我的母语不是英语，所以如果您有足够的耐心阅读我的问题，我将非常感谢。

由于 Cuda 实际上是在两台计算机上运行的，因此在设备上指向主机内存是无效的，这意味着如果结构(或对象)具有指针成员，则无法将其复制到设备。

我尝试制作以下系统来解决这个问题:

1. 使用整数代替指针。整数是内存池中的偏移量。整数包装在一个类中(重载“->”和“*”)，使其看起来像一个指针。
2. 内存池对象管理一个连续的对象数组，这些对象可以很容易地传输到Cuda设备。池的内容在主机和设备之间同步，因此整数偏移量在两侧具有相同的含义。

为了方便使用偏移量，应该将其包裹起来。在主机端，包装器看起来像这样:

template<typename T>
class MemPoolPointer {
public:
    inline T* operator -> () const
    {
        return &( MemPool<T>::get_instance.get_object(_p) );
    }
    uint64_t _p;
}

我们可以看到，指针类需要全局访问内存池。这通常是通过使内存池成为单例来实现的。但是，Cuda 不允许静态成员，并且它将 __device__ 变量限制在文件范围内。我该如何解决这些限制？或者我应该试试 OpenCL？

Answer 1

OP 能够通过使用如下静态类方法包装全局范围 __device__ 变量来解决此问题:

class FooBar;
__device__ FooBar* FOOBAR_DEVICE_POOL;
class FooBar
{
    __device__ static FooBar& DEVICE_GET(uint64_t p);
}

template<typename T>
class MemPoolPointer {
public:
    inline T* operator -> () const
    {
#ifdef __CUDA_ARCH__
        return &( T::DEVICE_GET(_p) );
#else
        return &( MemPool<T>::get_instance.get_object(_p) );
#endif
    }
    uint64_t _p;
}

[此答案作为社区 wiki 条目添加，以将问题从 CUDA 标记的未回答队列中移除]