CUDA C Programming Guide 在線教程學習筆記 Part 3

▶ 表面內存使用

● 創建 cuda 數組時使用標志 cudaArraySurfaceLoadStore 來創建表面內存，可以用表面對象（surface object）或表面引用（surface reference）來對其進行讀寫。

● 使用 Surface Object API

■ 涉及的結構定義、接口函數。

// vector_types.h
struct __device_builtin__ __align__(4) uchar4
{
    unsigned char x, y, z, w;
};

// surface_types.h
typedef __device_builtin__ unsigned long long cudaSurfaceObject_t;

■ 完整的測試代碼，使用表面內存進行簡單的讀寫。

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

#define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)

__global__ void myKernel(cudaSurfaceObject_t inputSurfObj, cudaSurfaceObject_t outputSurfObj, int width, int height)
{
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
    if (idx < width && idy < height)
    {
        uchar4 data;
        // 簡單的表面內存讀寫，使用了字節地址，而不是簡單的線程編號
        surf2Dread(&data, inputSurfObj, sizeof(float) * idx, idy);
        surf2Dwrite(data, outputSurfObj, sizeof(float) * idx, idy);
    }
    cudaBindSurfaceToArray();
}

int main()
{
    // 基本數據
    int i;
    float *h_data, *d_data; 
    int width = 32;
    int height = 32;

    int size = sizeof(float)*width*height;
    h_data = (float *)malloc(size);
    cudaMalloc((void **)&d_data, size);

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

    printf("\n\n");
    for (i = 0; i < width*height; i++)
    {
        printf("%6.1f ", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
    }
    
    // 申請 cuda 數組
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
    cudaArray* cuInputArray;
    cudaMallocArray(&cuInputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore); 
    cudaArray* cuOutputArray;
    cudaMallocArray(&cuOutputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);
    cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,cudaMemcpyHostToDevice);

    // 指定表面內存
    struct cudaResourceDesc resDesc;
    memset(&resDesc, 0, sizeof(resDesc));
    resDesc.resType = cudaResourceTypeArray;

    // 創建表面對象
    resDesc.res.array.array = cuInputArray;
    cudaSurfaceObject_t inputSurfObj = 0;
    cudaCreateSurfaceObject(&inputSurfObj, &resDesc);
    resDesc.res.array.array = cuOutputArray;
    cudaSurfaceObject_t outputSurfObj = 0;
    cudaCreateSurfaceObject(&outputSurfObj, &resDesc);

    // 運行核函數
    dim3 dimBlock(16, 16);
    dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
    myKernel << <dimGrid, dimBlock >> > (inputSurfObj, outputSurfObj, width, height);

    // 結果回收和檢查結果
    memset(h_data,0,size);// 刷掉原來的 h_data，再用 cuOutputArray 的數據寫入     
    cudaMemcpyFromArray(h_data, cuOutputArray, 0, 0, size, cudaMemcpyDeviceToHost);

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

    // 回收工作
    cudaDestroySurfaceObject(inputSurfObj);
    cudaDestroySurfaceObject(outputSurfObj);
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);

    getchar();
    return 0;
}

 

● 使用 Surface Reference API。

■ 表面引用的一些只讀屬性需要在聲明的時候指定，以便編譯時提前確定，只能在全局作用域內靜態指定，不能作為參數傳遞給函數。使用 surface 指定紋理引用屬性，Datatype 為數據類型，Type 為紋理引用類型，有 7 種，默認 cudaSurfaceType1D。

surface<void, Type> surfRef;

// cuda_texture_types.h
template<class T, int dim = 1>
struct __device_builtin_surface_type__ surface : public surfaceReference
{
#if !defined(__CUDACC_RTC__)
    __host__ surface(void)
    {
        channelDesc = cudaCreateChannelDesc<T>();
    }

    __host__ surface(struct cudaChannelFormatDesc desc)
    {
        channelDesc = desc;
    }
#endif /* !__CUDACC_RTC__ */  
};

//surface_types.h
#define cudaSurfaceType1D              0x01
#define cudaSurfaceType2D              0x02
#define cudaSurfaceType3D              0x03
#define cudaSurfaceTypeCubemap         0x0C
#define cudaSurfaceType1DLayered       0xF1
#define cudaSurfaceType2DLayered       0xF2
#define cudaSurfaceTypeCubemapLayered  0xFC

// 訪問邊界模式
enum __device_builtin__ cudaSurfaceBoundaryMode
{
    cudaBoundaryModeZero = 0,    // 0 邊界模式
    cudaBoundaryModeClamp = 1,   // 擠壓模式
    cudaBoundaryModeTrap = 2     // 陷阱模式
};

// ？表面格式模式
enum __device_builtin__  cudaSurfaceFormatMode
{
    cudaFormatModeForced = 0,   // 強制模式
    cudaFormatModeAuto = 1      // 自動模式
};

// 表面引用的通道描述
struct __device_builtin__ surfaceReference
{
    struct cudaChannelFormatDesc channelDesc;
};

// cuda_runtime_api.h
extern __host__ cudaError_t CUDARTAPI cudaBindSurfaceToArray(const struct surfaceReference *surfref, cudaArray_const_t array, const struct cudaChannelFormatDesc *desc);

 

■ 表面引用使用字節地址來定位訪問（而不是像紋理那樣使用 fetch 函數），如以上代碼中 surf1Dread(surfRef, sizeof(float) * idx) 或是 surf1Dread(surfRef, sizeof(float) * idx) 。

■ 表面引用必須用函數 cudaBindSurfaceToArray() 綁定到 cuda 數組上才能使用，要求表面引用的維度、數據類型與該數組匹配，否則操作時未定義的，使用完后不需要特殊函數來解除綁定。

■ 將表面引用綁定到 cuda 數組上的范例代碼。

// 准備工作
surface<void, Type>surfRef;

...

int width, height;
size_t pitch;
float *d_data;
cudaMallocPitch((void **)&d_data, &pitch, sizeof(float)*width, height);

// 第一種方法，低層 API
surfaceReference* surfRefPtr;
cudaGetSurfaceReference(&surfRefPtr, "surfRef");
cudaChannelFormatDesc channelDesc;
cudaGetChannelDesc(&channelDesc, cuArray);
cudaBindSurfaceToArray(surfRef, cuArray, &channelDesc);

// 第二種方法，高層 API
cudaBindSurfaceToArray(surfRef, cuArray);

 

■ 完整的應用樣例代碼。與前面表面對象代碼的功能相同。

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

#define CEIL(x,y) (((x) + (y) - 1) / (y) + 1)

// 聲明表面引用
surface<void, 2> inputSurfRef;
surface<void, 2> outputSurfRef;

__global__ void myKernel(int width, int height)
{
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;
    if (idx < width && idy < height)
    {
        uchar4 data;
        // 簡單的表面內存讀寫，使用了字節地址，而不是簡單的線程編號
        surf2Dread(&data, inputSurfRef, sizeof(float) * idx, idy);
        surf2Dwrite(data, outputSurfRef, sizeof(float) * idx, idy);
    }
}

int main()
{
    // 基本數據
    int i;
    float *h_data, *d_data; 
    int width = 32;
    int height = 32;

    int size = sizeof(float)*width*height;
    h_data = (float *)malloc(size);
    cudaMalloc((void **)&d_data, size);

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

    printf("\n\n");
    for (i = 0; i < width*height; i++)
    {
        printf("%6.1f ", h_data[i]);
        if ((i + 1) % width == 0)
            printf("\n");
    }
    
    // 申請 cuda 數組
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
    cudaArray* cuInputArray;
    cudaMallocArray(&cuInputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore); 
    cudaArray* cuOutputArray;
    cudaMallocArray(&cuOutputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);
    cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size,cudaMemcpyHostToDevice);

    // 綁定表面引用，注意與表面對象的使用不一樣
    cudaBindSurfaceToArray(inputSurfRef, cuInputArray);
    cudaBindSurfaceToArray(outputSurfRef, cuOutputArray);

    // 運行核函數
    dim3 dimBlock(16, 16);
    dim3 dimGrid(CEIL(width, dimBlock.x), CEIL(height, dimBlock.y));
    myKernel << <dimGrid, dimBlock >> > (width, height);

    // 結果回收和檢查結果
    memset(h_data,0,size);// 刷掉原來的 h_data，再用 cuOutputArray 的數據寫入     
    cudaMemcpyFromArray(h_data, cuOutputArray, 0, 0, size, cudaMemcpyDeviceToHost);

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

    // 回收工作
    cudaFreeArray(cuInputArray);
    cudaFreeArray(cuOutputArray);

    getchar();
    return 0;
}

 

▶ 立方體表面 Cubemap Surface 。 （想象成一個正方體的外表面）

● 一種特殊的二維分層表面。函數 surfCubemapread() 和函數 surfCubemapwrite() 來對其進行讀寫，使用一個整數下標和兩個浮點數有序組來定義層號和表面坐標。

 

▶ 分層立方體表面 Cubemap Layered Surfaces 。（想象成一個多層的正方體的各外表面）

● 一種特殊的二維分層表面。函數 surfCubemapread() 和函數 surfCubemapwrite() 來對齊進行讀寫。使用一個整數下標和兩個浮點數有序組來定義層號和表面坐標。

● 分層立方體貼圖紋理只能使用函數 cudaMAlloc3DArray() 加上 cudaArrayLayered 和 cudaArrayCubemap 標志來聲明，使用函數 texCubemapLayered() 來進行訪問濾波只在同一層內部進行，不會跨層執行。

 

▶ cuda 數組。

● cuda 優化的數組類型，可以有一維或二維或三維，每個元素可以有 1 個或 2 個或 4 個分量，各分量可以是 1 B 或 2 B 或 4 B 尺寸的有符號或無符號整數，或 2 B 或 4 B 尺寸的浮點數。cuda 數組只能用紋理訪問函數來訪問，或表面函數來進行讀寫。

● 紋理內存和表面內存都是可緩存的，且不能保證緩存和內存的一致性。同一個核函數中，用紋理訪問或表面訪問來讀取“已經全局寫入或表面寫入的內存”是未定義的。

 

▶ 壓縮版的 surface_types.h

#if !defined(__SURFACE_TYPES_H__)
#define __SURFACE_TYPES_H__

#include "driver_types.h"

#define cudaSurfaceType1D              0x01
#define cudaSurfaceType2D              0x02
#define cudaSurfaceType3D              0x03
#define cudaSurfaceTypeCubemap         0x0C
#define cudaSurfaceType1DLayered       0xF1
#define cudaSurfaceType2DLayered       0xF2
#define cudaSurfaceTypeCubemapLayered  0xFC

//CUDA Surface boundary modes
enum __device_builtin__ cudaSurfaceBoundaryMode
{
    cudaBoundaryModeZero = 0,   // Zero boundary mode */
    cudaBoundaryModeClamp = 1,  // Clamp boundary mode */
    cudaBoundaryModeTrap = 2    // Trap boundary mode */
};

//CUDA Surface format modes
enum __device_builtin__  cudaSurfaceFormatMode
{
    cudaFormatModeForced = 0,   // Forced format mode */
    cudaFormatModeAuto = 1      // Auto format mode */
};

//CUDA Surface reference
struct __device_builtin__ surfaceReference
{
    // Channel descriptor for surface reference
    struct cudaChannelFormatDesc channelDesc;
};

//An opaque value that represents a CUDA Surface object
typedef __device_builtin__ unsigned long long cudaSurfaceObject_t;


#endif