Как создать глобально доступную переменную на CUDA?
Это довольно сложный вопрос, и я не являюсь носителем английского языка, поэтому я буду благодарен, если вы достаточно терпеливы, чтобы прочитать мой вопрос.
Поскольку 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?
Решение
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;
}
Другие решения
Других решений пока нет …