CUDA全局（如C语言）分配给设备内存的动态数组

Question

所以，我试图编写一些利用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电话，我宁愿动态分配方法。

Answer 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; 
} 

给它一个旋风。

Answer 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++'。

Answer 3

呃，这正是将devPtr移动到全局范围的问题，这是我的问题。

我有一个实现，确实如此，两个内核有一个指向数据传入的指针。我明确不想传入这些指针。

我已经仔细阅读了文档，打了nvidia论坛（和谷歌搜索了一个小时左右），但我还没有找到一个实际运行的全球动态设备数组的实现（我试过几个编译，然后以新的和有趣的方式失败）。

Answer 4

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

Answer 5

我继续尝试分配临时指针并将它传递给类似于kernel1的简单全局函数的解决方案。

的好消息是，它的工作:)

不过，我认为它混淆编译器，我现在得到“咨询：不能说什么指针指向，假设全球内存空间”每当我试图访问全球数据。幸运的是，这个假设是正确的，但警告很烦人。

无论如何，为了记录 - 我已经看过很多例子，并且通过nvidia练习运行，其中的重点是让输出说“正确！”。不过，我还没有看过他们的全部。如果有人知道他们做动态全局设备内存分配的sdk示例，我仍然想知道。

Answer 6

由于此数据用于众多功能，我希望它是全球性的。

-

有两种使用全局几个很好的理由。这绝对不是一个。我将把它作为 练习来扩展此示例，以便将“devPtr”移动到全局范围。

如果内核在由阵列的大常量结构工作是什么？使用所谓的常量内存不是一种选择，因为它的大小非常有限......所以你必须把它放在全局内存中。