二維數組 cudaMallocPitch() 和三維數組 cudaMalloc3D() 的使用

▶ 使用函數 cudaMallocPitch() 和配套的函數 cudaMemcpy2D() 來使用二維數組。C 中二維數組內存分配是轉化為一維數組，連貫緊湊，每次訪問數組中的元素都必須從數組首元素開始遍歷；而 cuda 中這樣分配的二維數組內存保證了數組每一行首元素的地址值都按照 256 或 512 的倍數對齊，提高訪問效率，但使得每行末尾元素與下一行首元素地址可能不連貫，使用指針尋址時要注意考慮尾部。

// cuda_rumtime_api.h
extern __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr, size_t *pitch, size_t widthByte, size_t height);
 
extern __host__ cudaError_t CUDARTAPI cudaMemcpy2D(void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, enum cudaMemcpyKind kind);

● cudaMAllocPitch() 傳入存儲器指針 **devPtr，偏移值的指針 *pitch，數組行字節數 widthByte，數組行數 height。函數返回后指針指向分配的內存（每行地址對齊到 AlignByte 字節，為 256B 或 512B），偏移值指針指向的值為該行實際字節數（= sizeof(datatype) * width + alignByte - 1) / alignByte）。

● cudaMemcpy2D() 傳入目標存儲器的指針 *dst，目標存儲器行字節數 dpitch，源存儲器指針 *src，源存儲器行字節數 spitch，數組行字節數 width，數組行數 height，拷貝方向 kind。這里要求存儲器行字節數不小於數組行字節數，多出來的部分就是每行尾部空白部分。

● 整個測試代碼。

#include <stdio.h>
#include <malloc.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"

__global__ void myKernel(float* devPtr, int height, int width, int pitch)
{
    int row, col;
    float *rowHead;
    
    for (row = 0; row < height; row++)
    {
        rowHead = (float*)((char*)devPtr + row * pitch);
        
        for (col = 0; col < width; col++)
        {
            printf("\t%f", rowHead[col]);// 逐個打印並自增 1
            rowHead[col]++;
        }
        printf("\n");
    }
}

int main()
{
    size_t width = 6;
    size_t height = 5;
    float *h_data, *d_data;
    size_t pitch;

    h_data = (float *)malloc(sizeof(float)*width*height);
    for (int i = 0; i < width*height; i++)
        h_data[i] = (float)i;

    printf("\n\tAlloc memory.");
    cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);
    printf("\n\tPitch = %d B\n", pitch);
                                
    printf("\n\tCopy to Device.\n");
    cudaMemcpy2D(d_data, pitch, h_data, sizeof(float)*width, sizeof(float)*width, height, cudaMemcpyHostToDevice);
    
    myKernel << <1, 1 >> > (d_data, height, width, pitch);
    cudaDeviceSynchronize();

    printf("\n\tCopy back to Host.\n");
    cudaMemcpy2D(h_data, sizeof(float)*width, d_data, pitch, sizeof(float)*width, height, cudaMemcpyDeviceToHost);

    for (int i = 0; i < width*height; i++)
    {
        printf("\t%f", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
    }               

    free(h_data);
    cudaFree(d_data);

    getchar();
    return 0;
}

 ● 輸出結果：

        Alloc memory.
        Pitch = 512 B

        Copy to Device.
        0.000000        1.000000        2.000000        3.000000        4.000000        5.000000
        6.000000        7.000000        8.000000        9.000000        10.000000       11.000000
        12.000000       13.000000       14.000000       15.000000       16.000000       17.000000
        18.000000       19.000000       20.000000       21.000000       22.000000       23.000000
        24.000000       25.000000       26.000000       27.000000       28.000000       29.000000

        Copy back to Host.
        1.000000        2.000000        3.000000        4.000000        5.000000        6.000000
        7.000000        8.000000        9.000000        10.000000       11.000000       12.000000
        13.000000       14.000000       15.000000       16.000000       17.000000       18.000000
        19.000000       20.000000       21.000000       22.000000       23.000000       24.000000
        25.000000       26.000000       27.000000       28.000000       29.000000       30.000000

 

▶ 使用函數 cudaMalloc3D() 和配套的函數 cudaMemcpy3D() 來使用三維數組。因為涉及的參數較多，需要定義一些用來傳參的結構，形式上和二維數組的使用有較大差距，不好看。

● 涉及的相關代碼

// driver_types.h
struct cudaArray;                      // cuda 數組
typedef struct cudaArray * cudaArray_t;// cuda 指針

struct __device_builtin__ cudaPitchedPtr
{
    void   *ptr;      // 實際數組指針（用完后要用 cudaFree() 釋放掉）
    size_t  pitch;    // 數組行字節數
    size_t  xsize;    // 數組列數 
    size_t  ysize;    // 數組行數
};

struct __device_builtin__ cudaExtent
{
    size_t width;     // 數組行字節數 
    size_t height;    // 數組行數 
    size_t depth;     // 數組層數
};

struct __device_builtin__ cudaPos
{
    size_t x;
    size_t y;
    size_t z;
};

struct __device_builtin__ cudaMemcpy3DParms
{
    cudaArray_t            srcArray;  // 原數組指針
    struct cudaPos         srcPos;    // 原數組偏移 
    struct cudaPitchedPtr  srcPtr;    // ？Pitched source memory address 

    cudaArray_t            dstArray;  // 目標數組指針 
    struct cudaPos         dstPos;    // 目標數組偏移
    struct cudaPitchedPtr  dstPtr;    // ？Pitched destination memory address 

    struct cudaExtent      extent;    // 數組實際尺寸（去掉對齊用的空白部分） 
    enum cudaMemcpyKind    kind;      // 拷貝類型 
};

// driver_functions.h
static __inline__ __host__ struct cudaPitchedPtr make_cudaPitchedPtr(void *d, size_t p, size_t xsz, size_t ysz)
{                                     // 簡單生成 cudaPitchedPtr 結構的方法
    struct cudaPitchedPtr s;

    s.ptr = d;
    s.pitch = p;
    s.xsize = xsz;
    s.ysize = ysz;

    return s;
}

static __inline__ __host__ struct cudaPos make_cudaPos(size_t x, size_t y, size_t z)
{                                     // 簡單的生成 cudaPos 結構的方法
    struct cudaPos p;

    p.x = x;
    p.y = y;
    p.z = z;

    return p;
}

static __inline__ __host__ struct cudaExtent make_cudaExtent(size_t w, size_t h, size_t d)
{                                     // 簡單的生成 cudaExtent 結構的方法
    struct cudaExtent e;

    e.width = w;
    e.height = h;
    e.depth = d;

    return e;
}

// cuda_runtime_api.h
extern __host__ cudaError_t CUDARTAPI cudaMalloc3D(struct cudaPitchedPtr* pitchedDevPtr, struct cudaExtent extent);

extern __host__ cudaError_t CUDARTAPI cudaMemcpy3D(const struct cudaMemcpy3DParms *p);

 

● 完整的測試程序

#include <stdio.h>
#include <malloc.h>
#include <cuda_runtime_api.h>
#include "device_launch_parameters.h"
#include <driver_functions.h>

__global__ void myKernel(cudaPitchedPtr devPitchedPtr, cudaExtent extent)
{
    float * devPtr = (float *)devPitchedPtr.ptr;
    float *sliceHead, *rowHead;
        // 可以定義為 char * 作面、行遷移的時候直接加減字節數，取行內元素的時候再換回 float *

    for (int z = 0; z < extent.depth; z++)
    {
        sliceHead = (float *)((char *)devPtr + z * devPitchedPtr.pitch * extent.height);
        for (int y = 0; y < extent.height; y++)
        {
            rowHead = (float*)((char *)sliceHead + y * devPitchedPtr.pitch);
            for (int x = 0; x < extent.width / sizeof(float); x++)// extent 存儲的是行有效字節數，要除以元素大小 
            {
                printf("\t%f",rowHead[x]);// 逐個打印並自增 1
                rowHead[x]++;
            }
            printf("\n");
        }
        printf("\n");
    }
}

int main()
{
    size_t width = 2;
    size_t height = 3;
    size_t depth = 4;
    float *h_data;

    cudaPitchedPtr d_data;
    cudaExtent extent;
    cudaMemcpy3DParms cpyParm;

    h_data = (float *)malloc(sizeof(float) * width * height * depth);
    for (int i = 0; i < width * height * depth; i++)
        h_data[i] = (float)i;

    printf("\n\tAlloc memory.");
    extent = make_cudaExtent(sizeof(float) * width, height, depth);
    cudaMalloc3D(&d_data, extent);
                                
    printf("\n\tCopy to Device.\n");
    cpyParm = {0};
    cpyParm.srcPtr = make_cudaPitchedPtr((void*)h_data, sizeof(float) * width, width, height);
    cpyParm.dstPtr = d_data;
    cpyParm.extent = extent;
    cpyParm.kind = cudaMemcpyHostToDevice;
    cudaMemcpy3D(&cpyParm);

    myKernel << <1, 1 >> > (d_data, extent);
    cudaDeviceSynchronize();

    printf("\n\tCopy back to Host.\n");
    cpyParm = { 0 };
    cpyParm.srcPtr = d_data;
    cpyParm.dstPtr = make_cudaPitchedPtr((void*)h_data, sizeof(float) * width, width, height);
    cpyParm.extent = extent;
    cpyParm.kind = cudaMemcpyDeviceToHost;
    cudaMemcpy3D(&cpyParm);

    for (int i = 0; i < width*height*depth; i++)
    {
        printf("\t%f", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
        if ((i + 1) % (width*height) == 0)
            printf("\n");
    }               

    free(h_data);
    cudaFree(d_data.ptr);
    getchar();
    return 0;
}

 ● 輸出結果：

        Alloc memory.
        Copy to Device.
        0.000000        1.000000
        2.000000        3.000000
        4.000000        5.000000

        6.000000        7.000000
        8.000000        9.000000
        10.000000       11.000000

        12.000000       13.000000
        14.000000       15.000000
        16.000000       17.000000

        18.000000       19.000000
        20.000000       21.000000
        22.000000       23.000000


        Copy back to Host.
        1.000000        2.000000
        3.000000        4.000000
        5.000000        6.000000

        7.000000        8.000000
        9.000000        10.000000
        11.000000       12.000000

        13.000000       14.000000
        15.000000       16.000000
        17.000000       18.000000

        19.000000       20.000000
        21.000000       22.000000
        23.000000       24.000000