cuda设备变量在不同函数中的分配和使用

Question

我对 CUDA 很陌生，我有一个关于对象内存管理的问题。我有一个对象函数来将数据加载到设备中，如果调用另一个对象函数，则执行计算。

我已阅读 NVIDIA 编程指南的某些部分和一些 SO 问题，但它们在单个函数中进行数据复制和计算，因此不需要多个函数。

更多规格： 数据被读取一次。我不知道编译时的数据大小，因此我需要动态分配。我当前的设备具有 2.1 的计算能力（将很快更新到 6.1）。

我想在第一个函数中复制数据并在不同的函数中使用数据。例如：

__constant__ int dev_size;
__device__ float* dev_data; //<- not sure about this

/* kernel */
__global__ void computeSomething(float* dev_output)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < dev_size)
    {
        dev_output[idx] = dev_data[idx]*100; // some computation;
    }
}

// function 1
void OBJECT::copyVolumeToGPU(int size, float* data)
{
    cudaMalloc(&dev_data, size * sizeof(float));
    cudaMemcpy(dev_data, data, size * sizeof(float), cudaMemcpyHostToDevice );
    cudaMemcpyToSymbol(dev_size, size, sizeof(int));
}

// function 2
void OBJECT::computeSmthOnDevice(int size)
{
    // allocate output array
    auto host_output =  new float[size];
    float* dev_output;
    cudaMalloc(&dev_output, size * sizeof(float));

    int block = 256;
    int grid = ceil(size/block);
    computeSomething<<<grid,block>>>(dev_output);

    cudaMemcpy(host_output, dev_data, size * sizeof(float), cudaMemcpyDeviceToHost);

    /* ... do something with output ... */

    delete[] host_output;
    cudaFree(dev_output);
}

gpuErrChk 是这样执行的：https://stackoverflow.com/a/14038590/3921660 但在本例中省略了。

我可以使用__device__pointer（如__device__ float* dev_data;）复制数据吗？

Answer 1

一般来说，你的想法是可行的，但是这个：

cudaMalloc(&dev_data, size * sizeof(float));

不合法。在主机代码中获取__device__ 项目的地址是不合法的。因此，如果您在编译时知道大小，最简单的方法是将其转换为静态分配，例如

__device__ float dev_data[1000]; 

如果你真的想让它成为一个动态分配的__device__ 指针，那么你需要使用here 中描述的方法，它涉及在主机代码中的典型设备指针上使用cudaMalloc，这是一个“临时”，然后通过 cudaMemcpyToSymbol 将该“临时”指针复制到 __device__ 指针。然后，当您想通过cudaMemcpy 将数据复制到/从该特定分配中复制数据时，您将使用cudaMemcpy 到/从主机代码的临时 指针。

请注意，为了将数据从一个函数“通信”到下一个函数，或者一个内核到下一个函数，没有理由不能只使用来自cudaMemcpy 的动态分配的指针，并传递该指针到任何你需要的地方。你甚至可以通过一个全局变量将它传递给任何需要它的宿主函数，就像一个普通的全局指针。然而，对于内核，您仍然需要通过内核参数将这样一个全局指针传递给内核。