cudaMallocManaged 用于 2D 和 3D 阵列

Question

如果想要将阵列从主机复制到设备，请执行 cudamalloc 和 cudaMemcpy。但是为了减少麻烦，只需做 cudaMallocManaged 没有前两件事，生活从未如此简单。 代码看起来像这样（或多或少）

__global__ void convert(float kelvin[], float celsius[])  //can pass 
arrays in kernel
{
     int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i<N)
    kelvin[i]=celsius[i]+273.15;
}

int main()
{
    float *celsius =(float *)malloc(N*sizeof(float));
    float *kelvin =(float *)malloc(N*sizeof(float));
    cudaMallocManaged(&celsius, N*sizeof(float));
    cudaMallocManaged(&kelvin, N*sizeof(float));

// init celsius here

dim3 blocksPerGrid(1,1,1); //use only one block
dim3 threadsPerBlock(N,1,1); //use N threads in the block
convert<<<blocksPerGrid, threadsPerBlock>>>(kelvin,celsius);
cudaDeviceSynchronize();

//Doing stuff with the output here

return 0;
}

前面的例子对我来说似乎很清楚。 但是，怎么办 cudaMallocManaged 用于 2D 和 3D 数组？我一直在努力

__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}

int main()
{   // I thonk, 2D arrays can be passed as pointer to pointers
    float **A = (float **)malloc(N*N*sizeof(float));
    float **B = (float **)malloc(N*N*sizeof(float));
    float **C = (float **)malloc(N*N*sizeof(float));
    cudaMallocManaged(&A, N*N*sizeof(float));
    cudaMallocManaged(&B, N*N*sizeof(float));
    cudaMallocManaged(&C, N*N*sizeof(float));


A[N][N]={{1,0,0},{0,1,0},{0,0,1}};
B[N][N]={{1,0,0},{0,1,0},{0,0,1}};
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
//outputs and all

}

但是，它显示以下错误

matrix_add.cu(22): error: too many initializer values

matrix_add.cu(25): error: argument of type "float **" is incompatible with parameter of type "float (*)[3]"

非常感谢您的帮助。

Answer 1

您在尝试中犯了很多错误，以至于编写一个工作版本比在您的问题中列出代码中的所有单个问题要快。因此，这是您尝试做的工作的工作版本：

#include <algorithm>
#include <iostream>

const int N = 3;

__global__ void MatAdd(float A[][N], float B[][N], float C[][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    float* A; cudaMallocManaged(&A, N*N*sizeof(float));
    float* B; cudaMallocManaged(&B, N*N*sizeof(float));
    float* C; cudaMallocManaged(&C, N*N*sizeof(float));

    const float A_vals[N][N]={{1,0,0},{0,1,0},{0,0,1}};
    const float B_vals[N][N]={{1,0,0},{0,1,0},{0,0,1}};
    float (*C_vals)[N] = reinterpret_cast<float (*)[N]>(C);

    std::copy(&A_vals[0][0], &A_vals[0][0] + N*N, A);
    std::copy(&B_vals[0][0], &B_vals[0][0] + N*N, B);

    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(1, 1);
    MatAdd<<<numBlocks, threadsPerBlock>>>( reinterpret_cast<float (*)[N]>(A),
                                            reinterpret_cast<float (*)[N]>(B),
                                            C_vals );

    cudaDeviceSynchronize();

    for(int i=0; i<N; i++) {
        for(int j=0; j<N; j++) {
            std::cout << C_vals[i][j] << "  ";
        }
        std::cout << std::endl;
    }

    return 0;
}

一些重点：

1. 托管内存分配取代了标准的主机内存分配，并生成可在主机和设备上直接访问的内存。
2. 当按值作为参数传递给函数时，所有数组都会衰减为指针。这种衰减不是递归的。有关详细信息，请参阅here。
3. 您可以（并且需要）强制转换，以便在运行时动态分配的线性内存上使用[][] 访问语法（这适用于malloc、new 或任何CUDA 主机内存分配API . 详情请见here）。
4. 数组的初始化语法和赋值语法不可互换。

我只能建议你彻底研究它，直到你了解它是如何工作的。