CUDA全局(如C中)分配给设备内存的动态数组

jiusto 发布于 2018-05-02 cuda 最后更新 2018-05-02 08:36 338 浏览

所以,我试图编写一些利用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调用,但我更喜欢动态分配方法。
已邀请:

funde

赞同来自:

花一些时间专注于NVIDIA提供的丰富文档。 从编程指南:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));
这是如何分配内存的一个简单例子。现在,在你的内核中,你应该接受一个像这样的float的指针:
__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);
As this data is used in numerous functions, I would like it to be global.
有几个很好的理由使用全局变量。这绝对不是一个。我将把它作为一个练习来扩展这个例子,包括将“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 ++'。

jiusto

赞同来自:

呃,这正是将devPtr移到全球范围的问题,这是我的问题。 我有一个实现,就是这样做的,两个内核有一个指向数据的指针。我明确地不想传入这些指针。 我已经仔细阅读了文档,并打开了nvidia论坛(和谷歌搜索了一个小时左右),但我还没有找到实际运行的全局动态设备数组的实现(我已经尝试了几个编译和然后以新的和有趣的方式失败)。

funde

赞同来自:

像这样的东西应该可能工作。

#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;
}
给它一个旋转。

peum

赞同来自:

查看SDK附带的示例。这些示例项目中的很多都是一个体面的学习方式。

jiusto

赞同来自:

我继续尝试了分配临时指针并将它传递给类似于kernel1的简单全局函数的解决方案。 好消息是它确实工作:) 然而,我认为它会让编译器感到困惑,因为我现在每次尝试访问全局数据时都会得到“咨询:无法知道指向哪个指针,假设全局内存空间”。幸运的是,这个假设是正确的,但警告很烦人。 无论如何,为了记录 - 我已经看了很多例子,并且通过了nvidia练习,其中的重点是让输出结果表示“正确!”。但是,我没有看到所有这些。如果有人知道他们做动态全局设备内存分配的sdk示例,我仍然想知道。

taut

赞同来自:

As this data is used in numerous functions, I would like it to be global.
-
There are few good reasons to use globals. This definitely is not one. I'll leave it as an exercise to expand this example to include moving "devPtr" to a global scope.
内核如何在由数组组成的大型const结构上运行?使用所谓的常量内存不是一种选择,因为它的大小非常有限......所以你必须把它放在全局内存中。