CUDA 紋理內存

2013-09-30 10:36 4033人閱讀 評論(0) 收藏 舉報 分類： CUDA（106）

http://www.cnblogs.com/traceorigin/archive/2013/04/11/3015755.html

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 //1數組 2 cudaMalloc((void**)&dev_A, data_size); 3 cudaMemcpy(dev_A, host_A, data_size, cudaMemcpyHostToDevice); 4 cudaFree(dev_A); 5 6 //2維數組 7 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>() 8 cudaArray *cuArray; 9 cudaMallocArray(&cuArray, &channelDesc, 64, 32); //64x32 10 cudaMemcpyToArray(cuArray, 0, 0, h_data, sizeof(float)*width*height, cudaMemcpyHostToDevice); 11 cudaFreeArray(cuArray); 12 13 //3維數組 64x32x16 14 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>(); 15 cudaArray *d_volumeArray; 16 cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumSize); 17 18 cudaMemcpy3DParms copyParams = {0}; 19 copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height); 20 copyParams.dstArray = d_volumeArray; 21 copyParams.extent = volumeSize; 22 copyParams.kind = cudaMemcpyHostToDevice; 23 cudaMemcpy3D(&copyParams); 24 25 tex.normalized = true; 26 tex.filterMode = cudaFilterModeLinear; 27 tex.addressMode[0] = cudaAddressModeWrap; 28 tex.addressMode[1] = cudaAddressModeWrap; 29 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維紋理

1 cudaError_t cudaBindTexture( 2 size_t * offset, 3 const struct textureReference * texref, 4 const void * devPtr, 5 const struct cudaChannelFormatDesc * desc, 6 size_t size = UINT_MAX 7 )

例：

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

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

1 cudaError_t cudaBindTexture2D( 2 size_t * offset, 3 const struct textureReference * texref, 4 const void * devPtr, 5 const struct cudaChannelFormatDesc * desc, 6 size_t width, 7 size_t height, 8 size_t pitch 9 )

　例：

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

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

1 cudaError_t cudaBindTextureToArray ( 2 const struct textureReference * texref, 3 const struct cudaArray * array, 4 const struct cudaChannelFormatDesc * desc 5 )

例：

1 void initCudaTexture(const uchar *h_volume, cudaExtent volumeSize) 2 { 3 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>(); 4 5 cutilSafeCall(cudaMalloc3DArray(&d_volumeArray, &channelDesc, volumeSize)); 6 7 cudaMemcpy3DParms copyParams = {0}; 8 copyParams.srcPtr = make_cudaPitchedPtr((void*)h_volume, volumeSize.width*sizeof(uchar), volumeSize.width, volumeSize.height); 9 copyParams.dstArray = d_volumeArray; 10 copyParams.extent = volumeSize; 11 copyParams.kind = cudaMemcpyHostToDevice; 12 cutilSafeCall(cudaMemcpy3D(&copyParams)); 13 14 tex.normalized = true; 15 tex.filterMode = cudaFilterModeLinear; 16 tex.addressMode[0] = cudaAddressModeWrap; 17 tex.addressMode[1] = cudaAddressModeWrap; 18 tex.addressMode[2] = cudaAddressModeWrap; 19 20 cutilSafeCall(cudaBindTextureToArray(tex, d_volumeArray, channelDesc)); 21 }

?(5)紋理拾取

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

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

總結

以上是生活随笔為你收集整理的CUDA 纹理内存的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。