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

给它一个旋转.