如何在 cuda 上创建全局可访问的变量？答案

【问题标题】：How to create globally-accessible variable on cuda?如何在 cuda 上创建全局可访问的变量？
【发布时间】：2012-10-08 13:50:44
【问题描述】：

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

由于 Cuda 实际上是在两台计算机上运行，当你在设备上时，指向主机的内存是无效的，这意味着如果结构（或对象）有指针成员，你就不能将它们复制到设备。

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

使用整数而不是指针。整数是内存池中的偏移量。整数被包装在一个类中（重载“->”和“*”），使其看起来像一个指针。
内存池对象管理着一个连续的对象数组，可以很容易地转移到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？

【问题讨论】：

您是否尝试过使用固定分配？
单身人士：解决你从未遇到过的问题。
您的类（无论是否为单例），可以驻留在 CPU 上，并拥有指向使用普通 cudaMalloc 分配的主机和设备内存的指针等。然后您可以检索设备指针（根据需要偏移）在内核调用时从类中提取并将它们传递给内核。我认为这里没有问题...
@jiandingzhe：您能否将您的解决方案添加为该问题的答案？您稍后将能够接受自己的答案（这是允许的），并且该问题将被标记为已解决，以便其他人更容易通过搜索找到。谢谢。
@jiandingzhe：如果您接受我添加到此问题的社区 wiki 答案，这将很有帮助，以便从 CUDA 未回答的问题列表中删除

标签： c++ cuda memory-pool

【解决方案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 标签的未回答队列中删除]

【讨论】：