我们可以在 Cuda C++ 中对数组的数组使用动态分配吗？

Question

我在设备内存中有一个非常大的数组，我需要将它分成一些较小的部分。现在，我想知道我是否可以使用一个数组数组来通过索引访问它们。我尝试编写以下代码，但是，它 returns 垃圾，我认为这是因为它的未定义行为。它没有错误，我不知道是否可能。任何帮助将不胜感激。

#include <stdio.h>
#include <assert.h>
#include <iostream>

inline
cudaError_t checkCuda(cudaError_t result) {
#if defined(DEBUG) || defined(_DEBUG)
  if (result != cudaSuccess) {
    fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
    assert(result == cudaSuccess);
  }
#endif
  return result;
}

__global__ void cudaVectorFill(int **array, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N)
        array[0][i] = 1;
    else if (i < 2 * N)
        array[1][i - N] = 2;
    else if (i < 3 * N)
        array[2][i - 2 * N] = 3;
}

int main() {

    int N = 100000000;

    int **array = new int*[3];
 
    checkCuda( cudaMalloc(&array[0], N * sizeof(int)) );
    checkCuda( cudaMalloc(&array[1], N * sizeof(int)) );
    checkCuda( cudaMalloc(&array[2], N * sizeof(int)) );
 
    cudaVectorFill<<<(3 * N + 1023) / 1024, 1024>>>(array, N);

    checkCuda( cudaPeekAtLastError() );
 
    auto *host_array0 = new int[1];
    auto *host_array1 = new int[1];
    auto *host_array2 = new int[1];
 
    checkCuda( cudaMemcpy(host_array0, array[0], 1 * sizeof(int), cudaMemcpyDeviceToHost) );
    checkCuda( cudaMemcpy(host_array1, array[1], 1 * sizeof(int), cudaMemcpyDeviceToHost) );
    checkCuda( cudaMemcpy(host_array2, array[2], 1 * sizeof(int), cudaMemcpyDeviceToHost) );
 
    std::cout << *host_array0 << std::endl << *host_array1 << std::endl << *host_array2 << std::endl;

    return 0;
}

输出类似于：

707093096
707093104
707093112

正确的输出应该是：

1
2
3

Answer 1

如评论中所述，如果您将指针传递给 GPU 内核，它们必须可供 GPU 访问。这意味着您要么显式分配主机设备指针数组的副本并将其填充到设备上，要么依赖托管或 GPU 可访问的主机内存。

一种可能适用于这种情况的方法是：

int N = 100000000;

int **array = new int*[3];
 
checkCuda( cudaMalloc(&array[0], N * sizeof(int)) );
checkCuda( cudaMalloc(&array[1], N * sizeof(int)) );
checkCuda( cudaMalloc(&array[2], N * sizeof(int)) );

int **array_d;
checkCuda( cudaMalloc(&array_d, 3 * sizeof(int*)) );
checkCuda( cudaMemcpy(array_d, array, 3 * sizeof(int*), cudaMemcpyHostToDevice) );
 
cudaVectorFill<<<(3 * N + 1023) / 1024, 1024>>>(array_d, N);

[标准免责声明，代码在浏览器中编写，不提供或暗示任何保证，使用风险自负]

即在主机内存中构建 array 之后，在 GPU 内存中制作一个副本并将该 GPU 内存副本传递给您的内核。您的代码可能还有其他问题，我没有进一步分析前六行。

Answer 2

仅供参考，我刚刚发现了另一种在设备内存中进行二维分配的方法。有关详细信息，请参阅此中的方法 3。所以我们可以使用类似的东西：

int N = 100000000;

int **array;
cudaMallocManaged(&array, 3 * sizeof(int *));
cudaMallocManaged(&(array[0]), N * sizeof(int));
cudaMallocManaged(&(array[1]), N * sizeof(int));
cudaMallocManaged(&(array[2]), N * sizeof(int));

cudaVectorFill<<<(3 * N + 1023) / 1024, 1024>>>(array, N);

它也很好用。

我们可以在 Cuda C++ 中对数组的数组使用动态分配吗？

Can we use dynamic allocation for array of arrays in Cuda C++?

cuda