c++ - cudaMallocPitch 和 cudaMemcpy2D

Question

我在将 C++ 二维数组转换为 CUDA 一维数组时出错。让我展示一下我的源代码。

int main(void)
{
      float h_arr[1024][256];
      float *d_arr;

      // --- Some codes to populate h_arr

      // --- cudaMallocPitch
      size_t pitch;
      cudaMallocPitch((void**)&d_arr, &pitch, 256, 1024);

      // --- Copy array to device
      cudaMemcpy2D(d_arr, pitch, h_arr, 256, 256, 1024, cudaMemcpyHostToDevice);
}

我尝试运行代码，但它弹出错误。

如何正确使用cudaMallocPitch()和cudaMemcpy2D()？

Answer 1

Talonmies 已经很好地回答了这个问题。在这里，一些可能对社区有用的进一步解释。

在 CUDA 中访问二维数组时，如果每一行都正确对齐，内存事务会快得多。

CUDA 提供了 cudaMallocPitch 函数来用额外的字节“填充”二维矩阵行，从而实现所需的对齐。请参阅“CUDA C 编程指南”第 3.2.2 和 5.3.2 节，了解更多信息。

假设我们要分配浮点(单精度)元素的二维填充数组，cudaMallocPitch 的语法如下:

cudaMallocPitch(&devPtr, &devPitch, Ncols * sizeof(float), Nrows);

在哪里

• devPtr 是指向 float 的输出指针 (float *devPtr)。
• devPitch 是一个 size_t 输出变量，表示填充行的长度(以字节为单位)。
• Nrows 和 Ncols 是表示矩阵大小的 size_t 输入变量。

回想一下 C/C++ 和 CUDA 按行存储二维矩阵，cudaMallocPitch 将分配大小为字节的内存空间，等于 Nrows * pitch。但是，只有每行的第一个 Ncols * sizeof(float) 字节包含矩阵数据。因此，cudaMallocPitch 消耗的内存比 2D 矩阵存储严格需要的内存多，但这会在更高效的内存访问中返回。CUDA 还提供了 cudaMemcpy2D 函数，用于将数据从主机内存空间复制到/从设备内存空间复制到使用 cudaMallocPitch 分配的设备内存空间。在上述假设下(单精度二维矩阵)，语法如下:

cudaMemcpy2D(devPtr, devPitch, hostPtr, hostPitch, Ncols * sizeof(float), Nrows, cudaMemcpyHostToDevice)

在哪里

• devPtr 和 hostPtr 是 float 的输入指针(float *devPtr 和 float *hostPtr)指向分别是(源)设备和(目标)主机内存空间；
• devPitch 和 hostPitch 是 size_t 输入变量，表示设备和主机内存空间的填充行的长度(以字节为单位)，分别；
• Nrows 和 Ncols 是表示矩阵大小的 size_t 输入变量。

请注意，cudaMemcpy2D 还允许在主机端分配内存。如果主机内存没有间距，则 hostPtr = Ncols * sizeof(float)。此外，cudaMemcpy2D 是双向的。对于上面的示例，我们正在将数据从主机复制到设备。如果我们想从设备复制数据到主机，那么上面这行就变成了

cudaMemcpy2D(hostPtr, hostPitch, devPtr, devPitch, Ncols * sizeof(float), Nrows, cudaMemcpyDeviceToHost)

访问由 cudaMallocPitch 分配的二维矩阵的元素可以按照以下示例执行:

int    tidx = blockIdx.x*blockDim.x + threadIdx.x;
int    tidy = blockIdx.y*blockDim.y + threadIdx.y;

if ((tidx < Ncols) && (tidy < Nrows))
{
    float *row_a = (float *)((char*)devPtr + tidy * pitch);
    row_a[tidx] = row_a[tidx] * tidx * tidy;
}

在这样的示例中，tidx 和 tidy 分别用作列索引和行索引(请记住，在 CUDA 中，x-线程跨越列，y-threads 跨越行以促进合并)。指向一行第一个元素的指针是通过将初始指针 devPtr 偏移行长度 tidy * pitch 以字节为单位计算的(char *是一个指向字节的指针，sizeof(char) 是 1 字节)，其中每行的长度是使用间距信息计算的。

下面，我将提供一个完整的示例来展示这些概念。

#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>
#include<device_launch_parameters.h>
#include<conio.h>

#define BLOCKSIZE_x 16
#define BLOCKSIZE_y 16

#define Nrows 3
#define Ncols 5

/*****************/
/* CUDA MEMCHECK */
/*****************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort = true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %dn", cudaGetErrorString(code), file, line);
        if (abort) { getch(); exit(code); }
    }
}

/*******************/
/* iDivUp FUNCTION */
/*******************/
int iDivUp(int hostPtr, int b){ return ((hostPtr % b) != 0) ? (hostPtr / b + 1) : (hostPtr / b); }

/******************/
/* TEST KERNEL 2D */
/******************/
__global__ void test_kernel_2D(float *devPtr, size_t pitch)
{
    int    tidx = blockIdx.x*blockDim.x + threadIdx.x;
    int    tidy = blockIdx.y*blockDim.y + threadIdx.y;

    if ((tidx < Ncols) && (tidy < Nrows))
    {
        float *row_a = (float *)((char*)devPtr + tidy * pitch);
        row_a[tidx] = row_a[tidx] * tidx * tidy;
    }
}

/********/
/* MAIN */
/********/
int main()
{
    float hostPtr[Nrows][Ncols];
    float *devPtr;
    size_t pitch;

    for (int i = 0; i < Nrows; i++)
        for (int j = 0; j < Ncols; j++) {
            hostPtr[i][j] = 1.f;
            //printf("row %i column %i value %f \n", i, j, hostPtr[i][j]);
        }

    // --- 2D pitched allocation and host->device memcopy
    gpuErrchk(cudaMallocPitch(&devPtr, &pitch, Ncols * sizeof(float), Nrows));
    gpuErrchk(cudaMemcpy2D(devPtr, pitch, hostPtr, Ncols*sizeof(float), Ncols*sizeof(float), Nrows, cudaMemcpyHostToDevice));

    dim3 gridSize(iDivUp(Ncols, BLOCKSIZE_x), iDivUp(Nrows, BLOCKSIZE_y));
    dim3 blockSize(BLOCKSIZE_y, BLOCKSIZE_x);

    test_kernel_2D << <gridSize, blockSize >> >(devPtr, pitch);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy2D(hostPtr, Ncols * sizeof(float), devPtr, pitch, Ncols * sizeof(float), Nrows, cudaMemcpyDeviceToHost));

    for (int i = 0; i < Nrows; i++) 
        for (int j = 0; j < Ncols; j++) 
            printf("row %i column %i value %f \n", i, j, hostPtr[i][j]);

    return 0;    
}