CUDA中多維數(shù)組以及多維紋理內(nèi)存的使用

1、CUDA中多維數(shù)組以及多維紋理內(nèi)存的使用紋理存儲器(texture memory)是一種只讀存儲器，由GPU用于紋理渲染的的圖形專用單元發(fā)展而來，因此也提供了一些特殊功能。紋理存儲器中的數(shù) 據(jù)位于顯存，但可以通過紋理緩存加速讀取。在紋理存儲器中可以綁定的數(shù)據(jù) 比在常量存儲器可以聲明的 64K 大很多，并且支持一維、二維或者三維紋理。 在通用計算中，紋理存儲器十分適合用于實現(xiàn)圖像處理或查找表，并且對數(shù)據(jù) 量較大時的隨機(jī)數(shù)據(jù)訪問或者非對齊訪問也有良好的加速效果。紋理存儲器在硬件中并不對應(yīng)一塊專門的存儲器，而實際上是牽涉到顯 存、兩級紋理緩存、紋理抓取單元的紋理流水線。紋理存儲器提供了地址映 射、數(shù)

2、據(jù)濾波、緩存等功能，這些功能都是圍繞著紋理渲染的需求設(shè)計的。關(guān) 于GPU紋理流水線的介紹可以參考本書333節(jié)。在CUDA編程模型中，紋理緩存是透明的，編程人員不用去了解 它的實現(xiàn)機(jī)制。從CUDA的內(nèi)核函數(shù)訪問紋理存儲器的操作被稱為紋理拾取 (texture fetching)。紋理拾取使用的坐標(biāo)與數(shù)據(jù)在顯存中的地址可以不同，兩者通過紋理 參照系(texture referenee)約定從數(shù)據(jù)的地址到紋理坐標(biāo)的映射方式。將顯存 中的數(shù)據(jù)與紋理參照系關(guān)聯(lián)的操作，稱為將數(shù)據(jù)與紋理綁定 (texture binding)。 顯存中可以綁定到紋理的數(shù)據(jù)有兩種，分別是普通的線性內(nèi)存( Linear Mem

3、roy)禾口 CUDA數(shù)組(CUDA Array。CUDA數(shù)組則為紋理訪問進(jìn)行了優(yōu)化， 并且在Device端中只能通過紋理拾取訪問。綁定到紋理的線性內(nèi)存和數(shù)組中的元素被稱為像元(texels)，是texture elements的縮寫。像元的數(shù)據(jù)類型可以是其中的元素可以是 CUDA中規(guī)定的1， 2或者 4 元組(不能是 3 元組)的有符號或者無符號 8-,16-， 32-bit 整型或者 16- bit(目前只能通過driver API支持)整型，以及32-bit浮點(diǎn)型數(shù)據(jù)。與CUDA數(shù)組 綁定的紋理參照系中的元素使用的 N-元組數(shù)據(jù)中的組件數(shù)量必須與 CUDA數(shù)組 相同。紋理緩存有兩個作用。

4、首先，紋理緩存中的數(shù)據(jù)可以被重復(fù)利用，當(dāng)一次 訪問需要的數(shù)據(jù)已經(jīng)存在于紋理緩存中時，就可以避免對顯存進(jìn)行讀取。數(shù)據(jù) 重用過濾了一部分對顯存的訪問，節(jié)約了帶寬，也不必按照顯存對齊的要求讀 取。第二，紋理緩存可以緩存拾取坐標(biāo)附近幾個像元的數(shù)據(jù)，可以實現(xiàn)濾波模 式，也能提高具有一定局部性的訪問的效率。紋理存儲器是只讀的，不需要關(guān)心緩存數(shù)據(jù)一致性問題。這意味著如果更 改了綁定到紋理存儲器的數(shù)據(jù)，紋理緩存中的數(shù)據(jù)可能并沒有被更新，此時通 過紋理拾取得到的數(shù)據(jù)可能是錯誤的。因此，在每次修改了綁定到紋理的數(shù)據(jù) 以后，都需要對紋理進(jìn)行重新綁定。由于不能從設(shè)備端修改CUDA數(shù)組，因此只有在對綁定到紋理的線性內(nèi)存

5、進(jìn)行修改時才需要注意這一點(diǎn)。線性內(nèi)存中的數(shù)據(jù)只能與一維紋理綁定，并且紋理拾取坐標(biāo)是定點(diǎn)型，坐 標(biāo)的值也與數(shù)據(jù)在線性內(nèi)存中的偏移量相同；而CUDA數(shù)組可以與一維、二維或者三維紋理綁定，紋理拾取坐標(biāo)是浮點(diǎn)型，并且支持許多特殊功能。紋理存 儲器的特殊功能有：浮點(diǎn)型紋理拾取坐標(biāo)：使用浮點(diǎn)型的紋理拾取坐標(biāo)對紋理進(jìn)行尋址，只對與CUDA數(shù)組綁定的存儲器有效。地址映射的方式可以是歸一化或者非歸一化的：使用歸一化紋理時，紋理在每個維度上的坐標(biāo)被映射到浮點(diǎn)數(shù) 0.0,1.0）范圍內(nèi)；使用非歸一化紋理坐標(biāo)時，各個維度上的坐標(biāo)則被映射到浮點(diǎn) 數(shù)0.0, N）的范圍內(nèi)，其中N是紋理在該維度上像元的數(shù)量。由于在GPU中

6、通常用浮點(diǎn)計算點(diǎn)的坐標(biāo)，因此使用浮點(diǎn)數(shù)作為紋理拾取坐 標(biāo)更加自然；使用歸一化的紋理拾取坐標(biāo)可以不用關(guān)心紋理的實際尺寸，簡化 了渲染程序的編寫。尋址模式：尋址模式規(guī)定了紋理拾取的輸入坐標(biāo)超出紋理尋址范圍時的行為，有鉗位 模式和循環(huán)模式兩種。使用鉗位模式時，當(dāng)輸入的坐標(biāo)超出了尋址范圍，輸入的值將被 “鉗位 ”到尋址范圍的最大值或者最小值；循環(huán)模式只對歸一化坐標(biāo)有 效，此時要對超出尋址范圍的紋理坐標(biāo)作求模等處理。例如，對映射到 0.0,1.0）的歸一化紋理坐標(biāo)，輸入拾取坐標(biāo)1.25，鉗位模式會將輸入按照0.999,處理，而循環(huán)模式會將輸入0.25 處理。類型轉(zhuǎn)換：如果像元中的數(shù)據(jù)是 8-bit 或者

7、 16-bit 定點(diǎn)型，類型轉(zhuǎn)換功能對拾取的返回 值進(jìn)行類型轉(zhuǎn)換，將其映射到歸一化的浮點(diǎn)范圍 0.0f,1.0f （對無符號整型）或者 -1.0f,1.0f（對有符號整型）。濾波：如果將返回類型是浮點(diǎn)型的 CUDA數(shù)組與紋理綁定，那么就可以對返回的 值進(jìn)行濾波。濾波模式可以是最近點(diǎn)取樣模式或者線性濾波模式兩種。最近點(diǎn) 模式返回與浮點(diǎn)型的紋理抓取坐標(biāo)最近像元的值，而線性濾波模式則會先取出 附近幾個像元，然后按照抓取坐標(biāo)與這幾個像元的距離進(jìn)行線性插值，返回線 性插值得到的值。線性濾波可以使紋理渲染得到的畫面更加平滑自然。線性濾 波需要的插值計算不需要可編程單元參與，提供了額外的浮點(diǎn)處理能力，但精

8、度較低。使用線性濾波模式返回的值經(jīng)過了插值處理，適合用于圖像處理；使 用最近點(diǎn)取樣模式的返回值不會改變紋理中像元的值，適合用于實現(xiàn)查找表。關(guān)于紋理拾取模式的詳細(xì)描述，可以參考附錄 F。使用紋理存儲器時，首先要在主機(jī)端聲明需要綁定到紋理的線性存儲器或 CUDA數(shù)組，并設(shè)置好紋理參照系，然后將紋理參照系與線性內(nèi)存或者CUDA數(shù)組綁定。在主機(jī)端完成配置工作后，就可以在內(nèi)核函數(shù)中通過紋理抓取函數(shù)訪 問紋理存儲器了。1 CUDA數(shù)組在顯存中可以分配的空間有兩種：CUDA數(shù)組和線性內(nèi)存。此外，常數(shù)存儲器中通過緩存加速讀取的數(shù)據(jù)實際也存在于顯存中。CUDA數(shù)組和線性內(nèi)存都可以與紋理參照系綁定，但CUDA數(shù)組

9、對紋理拾取訪問進(jìn)行了優(yōu)化，在設(shè)備端也只能通過紋理拾取訪問。聲明CUDA數(shù)組之前，必須先以結(jié)構(gòu)體 channelDesc描述CUDA數(shù)組中的數(shù) 據(jù)類型。structcudaCha nn elFormatDesc ;intx, y,乙 w;enu mcudaCha nn elFormatK ind f;其中，x, y, z和w分別是每個返回值成員的位數(shù)，而 f是一個枚舉變量，可 以取一下幾個值：cudaChannelFormatKindSigned如果這些成員是有符號整型；cudaChannelFormatKindUnsigned如果這些成員是無符號整型；cudaChannelFormatKind

10、Float如果這些成員是浮點(diǎn)型；然后，我們要確定CUDA數(shù)組的維度和尺寸。CUDA數(shù)組可以通過 cudaMalloc3DArray()或 cudaMallocArray()函數(shù)分配。用 cudaMalloc3DArray 可以 分配一維、二維或者三維的 CUDA數(shù)組，而cudaMallocArray()般用于分配二維 CUDA數(shù)組。在使用完CUDA數(shù)組后，要使用cudaFreeArray函數(shù)釋放顯存。由cudaMalloc3DArray分配的CUDA數(shù)組使用cudaMemcpy3D()完成與其他 CUDA數(shù)組或者線性內(nèi)存的數(shù)據(jù)傳輸。CUDA AP中使用結(jié)構(gòu)體cudaExtent描述3D Arr

11、ay和3D線性內(nèi)存在三個維度上的尺寸，在描述一維、二維和三維數(shù)組分 別用以下的形式：cudaextent extent = make_cudaextent(1,8192,0,0);cudaextent extent = make_cudaextent(1,65535,1,32768,0);cudaextent extent = make_cudaextent(1,2048,1,2048,1,2048);其中方括號內(nèi)為允許的尋址范圍。注意到二維 CUDA數(shù)組的第一個維度的 尋址范圍大于一維CUDA數(shù)組的尋址范圍，因此在一維 CUDA數(shù)組的尺寸不夠用 時，將二維CUDA數(shù)組的第二個維度設(shè)為1代替一

12、維CUDA數(shù)組，獲得更大的尋 址范圍。下面是聲明一個數(shù)據(jù)類型為 char2型，寬 高 深為64X 32X 1的CUDA 3D數(shù) 組，對其初始化，最后釋放數(shù)組的示例代碼：cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc(8, 8, 0,O,cudaChannelFormatKindunsigned); /海個像元由兩個 char 構(gòu)成cudaExtent extent = make_cudaextent(64,32,16);/健立 cudaExtent結(jié)構(gòu)體， 描述CUDA數(shù)組的維度和尺寸cudaArray* cuArray;cud

13、aMalloc3DArray(&cuArray, &channelDesc, extent); 為 cuArray 開辟空間缺 cudaFreeArray(cuArray);下面則是使用cudaMallocArray聲明一個由float型構(gòu)成，尺寸為64X 32的 CUDA數(shù)組，對其賦值，并最后釋放的示例代碼：cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc(32, 0, 0,0,cudaChannelFormatKindunsigned); /海個像元由一個 float 構(gòu)成cudaArray* cuArray;cudaMall

14、ocArray(&cuArray, &channelDesc, 64, 32); /為 cuArray 開辟空間cudaMemcpyToArray(cuArray, 0, 0, h_data, & cha nn elDesc);/第二和第三個參 數(shù)分別表示在寬度和高度上的偏移量，假設(shè)h_data中的數(shù)據(jù)已經(jīng)初始化cudaFreeArray(cuArray);用于在CUDA數(shù)組和主機(jī)端或者設(shè)備端線性內(nèi)存，以及在CUDA數(shù)組間傳輸數(shù)據(jù)的函數(shù)還有很多，這些還是還有一些異步調(diào)用版本，關(guān)于這些函數(shù)的具體使用方法請參考CUDA Referenee mannuaJ 2聲明紋理參照系紋理參照系中的一些屬性必須

15、在編譯時之前被顯示聲明。紋理參照系通過一個作用范圍為全文件的texture型變量聲明：texturevType, Dim, ReadModetexRef;其中，Type確定了由紋理拾取返回的數(shù)據(jù)類型；Type可以是B3.1節(jié)中描述的任意一種由基本整型或者單精度浮點(diǎn)型組成能的1-，2-或者4-元組向量類型。Dim確定了紋理參照系的維度，默認(rèn)為 1。ReadMode可以是 cudaReadModeNormalizedFloat或者 cudaReadModeElementType 如果 ReadMode 是 cudaReadModeNomalizedFloat,并且 Type是 16-或者 8-bi

16、t 整型，那么返回的值將是一個浮點(diǎn)數(shù)。此時，原來整形的值域會被映射到0.0，1.0(對無符號整型)，或者-1.0，1.0(對有符號整型)。例如，一個值為0xff的8-bit無符號整型會被映射為1.0f。如果使用cudaReadModeElementType那么就不會對輸出進(jìn)行轉(zhuǎn)換。ReadMode是一個可選參數(shù)，如果不寫，那么默認(rèn)就是cudaReadModeEleme ntType例如，下面的代碼聲明了一個二維，像元數(shù)據(jù)為un sig ned char型，但將返回值轉(zhuǎn)換為float型的紋理參照系：texturev un sig ned char, 2, cudaReadModeNormaliz

17、edFloattexRef;2.1設(shè)置運(yùn)行時紋理參照系屬性紋理參照系中的其它屬性可以不必聲明，并在運(yùn)行時進(jìn)行修改。這些參數(shù) 規(guī)定了紋理的尋址模式，是否進(jìn)行歸一化，以及紋理濾波。runtimeAPI擁有底層的C風(fēng)格和高層的C+風(fēng)格兩種接口。高層API中的texture類型是從底層的 textureReferenee中派生而來的。TextureReferenee是一個下面的代碼描述的結(jié) 構(gòu)體。structtextureRefere nee i ntn ormalized;enu mcudaTextureFilterModefilterMode;enu mcudaTextureAddressMode

18、addressMode3;structcudaCha nn elFormatDesccha nn elDesc;normalized設(shè)置是否對紋理坐標(biāo)是否進(jìn)行歸一化。如果normalized是一個非零值，那么就會使用歸一化到0，1）的坐標(biāo)進(jìn)行尋址，否則對尺寸為 width,height,depth 的紋理使用坐標(biāo)O,width-1,O,height-1,O,depth-1尋址。例 如，一個尺寸為64X32的紋理可以通過x維度范圍為0, 63, y維度范圍0,31 的坐標(biāo)尋址。如果采用歸一化方式對尺寸為64X 32的紋理進(jìn)行尋址，在x和y維度上的坐標(biāo)就都是0.0,1.0）。這樣就可以保證紋理的坐

19、標(biāo)與紋理的尺寸無關(guān)。filterMode用于設(shè)置紋理的濾波模式，即如何根據(jù)坐標(biāo)計算返回的紋理 值。濾波模式可以是 cudaFilterModePoint或者cudaFilterModeLinear。濾波模式 為CudaFilterModePoint時，返回值是與坐標(biāo)最接近的像元的值。CudaFilterModeLi near模式只能對返回值為浮點(diǎn)型的紋理使用，啟用這一種 模式時將拾取紋理坐標(biāo)周圍的像元，然后根據(jù)坐標(biāo)與這些像元之間的距離進(jìn)行 插值計算。對一維紋理可以使用線性濾波，對二維紋理可以使用雙線性濾波。返回值會是對最接近紋理坐標(biāo)的兩個像元（對一維紋理），四個像元（對 二維紋理）或者八個像元

20、（對三維紋理）進(jìn)行插值后得到的值。addressmode說明了尋址模式，即如何處理超出尋址范圍的紋理坐標(biāo)； addressmode是一個大小為3的數(shù)組，三個元素分別說明對第一、二、三個紋理 坐標(biāo)的取址模式；取址模式可以是 cudaAddressModeClamp或 cudaAddressModeWrap中的一種，前者將超出尋址范圍的紋理坐標(biāo) ”鉗位”到尋 址范圍內(nèi)的最大或最小值，后者將超出尋址范圍的紋理坐標(biāo)折疊”進(jìn)合理范圍。cudaAddressModeWrap只支持歸一化的紋理坐標(biāo)。對非歸一化的坐標(biāo)，如果尋址的坐標(biāo)超過了范圍0, N，大于N的坐標(biāo)將被鉗位，設(shè)為N-1。對歸一化的坐標(biāo)，有鉗位和

21、循環(huán)兩種處理方式，在鉗位方式下，超過0.0，1.0）范圍的坐標(biāo)將被鉗位到0.0，1.0）；循環(huán)方式一般用于周期循環(huán)紋理，它只使用了紋理坐標(biāo)中有用的小數(shù) 部分，例如1.25會被當(dāng)作0.25處理，而-1.25則會被當(dāng)成0.75處理cha nn elDesc描述紋理獲取返回值類型，我們已經(jīng)在324.1小節(jié)講解CUDA array時介紹過這個結(jié)構(gòu)體。紋理參照系的返回值類型描 述必須和與之綁定的CUDA array的數(shù)據(jù)類型描述相同，或者和與之綁定的線性 內(nèi)存中的元素類型相同。normalized, addressMode和filterMode可以直接在主機(jī)端代碼中修改。它們 只適用于與CUDA數(shù)組綁定

22、的紋理參照系。附錄 D中列出了關(guān)于紋理拾取的更 多信息。2.2紋理綁定在kernel能用紋理參照系從紋理內(nèi)存中讀數(shù)據(jù)前，紋理參照系必須通過cudaBindTexture()或 cudaBindTextureToArray(綁定到紋理上。 cudaU nbin dTexture()用于解除紋理參照系的綁定。以下代碼示例綁定一個紋理參照系到devPtr指向的線性內(nèi)存：texturevfloat, 2, cudaReadModeEleme ntTypetexRef;textureRefere nee* texRefPtr;cudaGetTextureReference(&texRfPtr,“ tex

23、Ref ” );cudaCha nn elFormatDesccha nn elDesc = cudaCreateCha nn elDesc();cudaBi ndTexture2D(0, texRefPtr, devPtr, & cha nn elDesc, width, height, pitch);使用低級API:使用高級APItexturetexRef;cudaCha nn elFormatDesccha nn elDesc = cudaCreateCha nn elDesc();cudaBi ndTexture2D(0, texRef, devPtr, &cha nn elDesc,

24、 width, height, pitch);以下代碼示例綁定紋理參照系到一個CUDA數(shù)組cuArray:使用低級API:texturevfloat, 2, cudaReadModeEleme ntTypetexRef;textureRefere nee* texRefPtr;cudaGetTextureRefererce(&texRefPtr,“ texRef ” );cudaCha nn elFormatDesccha nn elDesc;cudaGetCha nn elDesc (&cha nn elDesc, cuArray);cudaB in dTextureToArray(texR

25、ef, cuArray, &cha nn elDesc);使用高級APItexturevfloat, 2, cudaReadModeEleme ntTypetexRef;cudaB in dTextureToArray(texRef, cuArray);當(dāng)綁定一個紋理到紋理參照系時，格式必須與聲明紋理參照系時的參數(shù)匹配；否則，紋理獲取的結(jié)果是 undefined的。2.3紋理拾取紋理拾取函數(shù)采用紋理坐標(biāo)對紋理存儲器進(jìn)行訪問。對與線性內(nèi)存綁定的紋理，使用texfetchlD函數(shù)訪問，采用的紋理坐標(biāo)是整型。由cudaMallocPitch或者cudaMalloc3D分配的線性空間實際上仍然是經(jīng)過填

26、充、對齊的一維線性空 間，因此也用texfetch1D()函數(shù)訪問。對與一維、二維和三維 CUDA數(shù)組綁定的問哪里，分別使用tex1D()、 tex2D()和tex3D()函數(shù)訪問，并且使用浮點(diǎn)型紋理坐標(biāo)。關(guān)于紋理拾取函數(shù)的更多討論，請見本書附錄D.82.4 例子分析：Simple texture/ 2D float texturetexturetexRef;/ Simple transformation kernel_global_voidtransformKernel(float* output, intwidth,intheight,floattheta)/ Host codeintma

27、in()/ 分配 CUDA數(shù)組cudaChannelFormatDescchannelDesc = cudaCreateChannelDesc(32, 0, 0, cudaArray* cuArray;cudaMallocArray(&cuArray, &channelDesc, width, height);/ Copy to device memory some data located at address h_data/ in host memorycudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);/ Set texture parameters/根據(jù) tid bid 計算歸一化的拾取坐標(biāo)unsignedintx = blockIdx.x * blockDim.x + threadIdx.x;unsignedinty = blockIdx.y * blockDim.