CUDA 紋理內存

1、概述

　　紋理存儲器中的數據以一維、二維或者三維數組的形式存儲在顯存中，可以通過緩存加速訪問，並且可以聲明大小比常數存儲器要大的多。

　　在kernel中訪問紋理存儲器的操作稱為紋理拾取(texture fetching)。將顯存中的數據與紋理參照系關聯的操作，稱為將數據與紋理綁定(texture binding).

　　顯存中可以綁定到紋理的數據有兩種，分別是普通的線性存儲器和cuda數組。

　　注：線性存儲器只能與一維或二維紋理綁定,采用整型紋理拾取坐標，坐標值與數據在存儲器中的位置相同；

    　　 CUDA數組可以與一維、二維、三維紋理綁定，紋理拾取坐標為歸一化或者非歸一化的浮點型，並且支持許多特殊功能。

2、紋理緩存：

　　（1）、紋理緩存中的數據可以被重復利用

　　（2）、紋理緩存一次預取拾取坐標對應位置附近的幾個象元，可以實現濾波模式。

3、紋理存儲器的特殊功能

4、紋理存儲器的使用

　　使用紋理存儲器時，首先要在主機端聲明要綁定到紋理的線性存儲器或CUDA數組

　　(1)聲明紋理參考系

texture<Type, Dim, ReadMode> texRef;
//Type指定數據類型，特別注意：不支持3元組
//Dim指定紋理參考系的維度，默認為1
//ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默認)

　　　　注：紋理參照系必須定義在所有函數體外

　　(2) 聲明CUDA數組，分配空間

　　　　CUDA數組可以通過cudaMalloc3DArray()或者cudaMallocArray()函數分配。前者可以分配1D、2D、3D的數組，后者一般用於分配2D的CUDA數組。使用完畢，要用       　　　　　　　cudaFreeArray()函數釋放顯存。 

//1數組
cudaMalloc((void**)&dev_A, data_size);
cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice);
cudaFree(dev_A);

//2維數組
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>()
cudaArray *cuArray;
cudaMallocArray(&cuArray, &channelDesc, 64, 32); //64x32
cudaMemcpyToArray(cuArray, 0, 0, h_data, sizeof(float)*width*height, cudaMemcpyHostToDevice);
cudaFreeArray(cuArray);

//3維數組 64x32x16
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();
cudaArray *d_volumeArray;
cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumSize);

cudaMemcpy3DParms copyParams = {0};
copyParams.srcPtr   = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
copyParams.dstArray = d_volumeArray;
copyParams.extent   = volumeSize;
copyParams.kind     = cudaMemcpyHostToDevice;
cudaMemcpy3D(&copyParams);

tex.normalized = true;
tex.filterMode = cudaFilterModeLinear;
tex.addressMode[0] = cudaAddressModeWrap;
tex.addressMode[1] = cudaAddressModeWrap;
tex.addressMode[2] = cudaAddressModeWrap;

（3）設置運行時紋理參照系屬性

struct textureReference{
    int normalized;
    enum cudaTextureFilterMode filterMode;
    enum cudaTextureAddressMode addressMode[3];
    struct cudaChannelFormatDesc channelDesc;
}

　　normalized設置是否對紋理坐標歸一化

　　filterMode用於設置紋理的濾波模式

　　addressMode說明了尋址方式

 (4)紋理綁定

　　通過cudaBindTexture() 或 cudaBindTextureToArray()將數據與紋理綁定。

　　通過cudaUnbindTexture()用於解除紋理參照系的綁定

　　注：與紋理綁定的數據的類型必須與聲明紋理參照系時的參數匹配

　　(I).cudaBindTexture() //將1維線性內存綁定到1維紋理

cudaError_t cudaBindTexture(    
    size_t *     offset,
    const struct textureReference *     texref,
    const void *     devPtr,
    const struct cudaChannelFormatDesc *     desc,
    size_t     size = UINT_MAX     
)    

例：

cudaMalloc((void**)&data.dev_inSrc, imageSize);
cudaBindTexture(NULL, tex, data.dev_inSrc, imageSize);

　　(II).cudaBindTexture2D //將1維線性內存綁定到2維紋理

cudaError_t cudaBindTexture2D(    
    size_t *     offset,
    const struct textureReference *     texref,
    const void *     devPtr,
    const struct cudaChannelFormatDesc *     desc,
    size_t     width,
    size_t     height,
    size_t     pitch     
)    

　例：

cudaMalloc((void**)&data.dev_inSrc, imageSize);
cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
cudaBindTexture2D(NULL, tex, data.dev_inSrc, desc, DIM, DIM, sizeof(float)*DIM);

　　(III). cudaBindTextureToArray() //將cuda數組綁定到紋理

cudaError_t cudaBindTextureToArray    (    
    const struct textureReference *     texref,
    const struct cudaArray *     array,
    const struct cudaChannelFormatDesc *     desc     
)

例：

void initCudaTexture(const uchar *h_volume, cudaExtent volumeSize)
{
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>();

    cutilSafeCall(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize));

    cudaMemcpy3DParms copyParams = {0};
    copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height);
    copyParams.dstArray = d_volumeArray;
    copyParams.extent   = volumeSize;
    copyParams.kind     = cudaMemcpyHostToDevice;
    cutilSafeCall(cudaMemcpy3D(&copyParams));

    tex.normalized = true;
    tex.filterMode = cudaFilterModeLinear;
    tex.addressMode[0] = cudaAddressModeWrap;
    tex.addressMode[1] = cudaAddressModeWrap;
    tex.addressMode[2] = cudaAddressModeWrap;

    cutilSafeCall(cudaBindTextureToArray(tex, d_volumeArray, channelDesc));
}

 (5)紋理拾取

　　對於線性存儲器綁定的紋理，使用tex1Dfetch()訪問，采用的紋理坐標是整型。由cudaMallocPitch() 或者 cudaMalloc3D()分配的線性空間實際上仍然是經過填充、對齊的一維線性空　　　間，因此也用tex1Dfetch()

　　對與一維、二維、三維cuda數組綁定的紋理，分別使用tex1D(), tex2D() 和 tex3D()函數訪問，並且使用浮點型紋理坐標。