分配给设备内存的 CUDA 全局（如 C 语言）动态数组
Posted
技术标签:
【中文标题】分配给设备内存的 CUDA 全局（如 C 语言）动态数组【英文标题】：CUDA global (as in C) dynamic arrays allocated to device memory 【发布时间】：2010-09-09 21:32:05 【问题描述】：所以，我尝试编写一些利用 Nvidia 的 CUDA 架构的代码。我注意到在设备之间进行复制确实会损害我的整体性能，所以现在我正尝试将大量数据移动到设备上。
由于这些数据用于许多功能，我希望它是全局的。是的，我可以传递指针，但我真的很想知道在这种情况下如何使用全局变量。
所以，我有想要访问设备分配数组的设备函数。
理想情况下，我可以这样做：
__device__ float* global_data;
main()
cudaMalloc(global_data);
kernel1<<<blah>>>(blah); //access global data
kernel2<<<blah>>>(blah); //access global data again
但是，我还没有弄清楚如何创建动态数组。我通过如下声明数组找到了解决方法：
__device__ float global_data[REALLY_LARGE_NUMBER];
虽然这不需要 cudaMalloc 调用，但我更喜欢动态分配方法。
【问题讨论】：
看看使用共享内存，全局是设备内存层中最慢的。 为什么要使用全局变量而不是将设备指针作为参数传递给内核？这样做只会给您带来与在 CPU 代码中使用全局内存相同的限制，几乎没有优势。 【参考方案1】：这样的事情应该可以工作。
#include <algorithm>
#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do \
cudaThreadSynchronize(); \
cudaError_t err = cudaGetLastError(); \
if( cudaSuccess != err) \
fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n", \
errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
exit(EXIT_FAILURE); \
while (0)
__device__ float *devPtr;
__global__
void kernel1(float *some_neat_data)
devPtr = some_neat_data;
__global__
void kernel2(void)
devPtr[threadIdx.x] *= .3f;
int main(int argc, char *argv[])
float* otherDevPtr;
cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));
kernel1<<<1,128>>>(otherDevPtr);
CUT_CHECK_ERROR("kernel1");
kernel2<<<1,128>>>();
CUT_CHECK_ERROR("kernel2");
return 0;
试一试。
【讨论】：
这很有趣。我看到 cudaMalloc 不是直接在 devPtr 上调用的，而是在第一个内核调用中设置的。稍后我会试一试，如果成功了再告诉你，非常感谢:D【参考方案2】：花一些时间专注于 NVIDIA 提供的大量文档。
来自编程指南：
float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));
这是一个如何分配内存的简单示例。现在，在您的内核中，您应该接受一个指向浮点数的指针，如下所示：
__global__
void kernel1(float *some_neat_data)
some_neat_data[threadIdx.x]++;
__global__
void kernel2(float *potentially_that_same_neat_data)
potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
所以现在你可以像这样调用它们：
float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));
kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);
由于这些数据被用于许多 功能，我希望它是 全球。
使用全局变量的充分理由很少。这绝对不是一个。我将把它作为一个练习来扩展这个例子，包括将“devPtr”移动到全局范围。
编辑：
好的，基本问题是这样的：您的内核只能访问设备内存，并且它们可以使用的唯一全局范围指针是 GPU 指针。从 CPU 调用内核时，在幕后发生的事情是指针和原语在内核执行之前被复制到 GPU 寄存器和/或共享内存中。
所以我可以建议的最接近的是：使用 cudaMemcpyToSymbol() 来实现您的目标。但是，在后台，请考虑另一种方法可能是正确的方法。
#include <algorithm>
__constant__ float devPtr[1024];
__global__
void kernel1(float *some_neat_data)
some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
__global__
void kernel2(float *potentially_that_same_neat_data)
potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
int main(int argc, char *argv[])
float some_data[256];
for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
some_data[i] = i * 2;
cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
float* otherDevPtr;
cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));
kernel1<<<1,128>>>(otherDevPtr);
kernel2<<<1,128>>>(otherDevPtr);
return 0;
不要忘记此示例中的“--host-compilation=c++”。
【讨论】：
是的 - 这是我最初的解决方案。只是，不在常量内存中，因为数组相当大：constant float* devPtr 的判断是什么？ （或者在我的情况下 device float* devPtr;）我怀疑有一个很好的理由为什么你不能有一个指向设备数据的全局指针 另外 - 没有看到您的编辑。但是，我仍然不确定为什么在数组正常时指向设备内存的 指针 无效。【参考方案3】：我继续尝试了分配临时指针并将其传递给类似于 kernel1 的简单全局函数的解决方案。
好消息是它确实有效:)
但是，我认为这会使编译器感到困惑，因为我现在在尝试访问全局数据时得到“建议：无法判断指针指向什么，假设是全局内存空间”。幸运的是，这个假设恰好是正确的，但是警告很烦人。
无论如何，为了记录 - 我已经查看了许多示例，并且确实运行了 nvidia 练习，其中的重点是让输出说“正确！”。但是，我没有看过所有。如果有人知道他们进行动态全局设备内存分配的 sdk 示例，我仍然想知道。
【讨论】：
【参考方案4】：呃，我的问题正是将 devPtr 移动到全局范围的问题。
我有一个完全可以做到这一点的实现，两个内核都有一个指向传入数据的指针。我明确不想传入这些指针。
我已经相当仔细地阅读了文档，并访问了 nvidia 论坛（谷歌搜索了一个小时左右），但我还没有找到实际运行的全局动态设备数组的实现（我尝试了几个编译然后以新的有趣方式失败）。
【讨论】：
【参考方案5】：查看 SDK 中包含的示例。其中许多示例项目都是通过示例学习的不错方式。
【讨论】：
【参考方案6】：由于这些数据用于许多功能，我希望它是全局的。
-
使用全局变量的充分理由很少。这绝对不是一个。我会把它作为一个 练习扩展此示例以包括将“devPtr”移动到全局范围。
如果内核在一个由数组组成的大型 const 结构上运行会怎样？使用所谓的常量内存不是一种选择，因为它的大小非常有限..所以你必须把它放在全局内存中..？
【讨论】：
以上是关于分配给设备内存的 CUDA 全局（如 C 语言）动态数组的主要内容，如果未能解决你的问题，请参考以下文章