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

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

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

3、y)。CUDA數(shù)組則為紋理訪問進(jìn)行了優(yōu)化，并且在Device端中只能通過紋理拾取訪問。綁定到紋理的線性內(nèi)存和數(shù)組中的元素被稱為像元(texels)，是textureelements的縮寫。像元的數(shù)據(jù)類型可以是其中的元素可以是CUDA中規(guī)定的1,2或者4元組(不能是3元組)的有符號(hào)或者無符號(hào)8-,16-，32-bit整型或者16-bit(目前只能通過driverAPI支持)整型，以及32-bit浮點(diǎn)型數(shù)據(jù)。與CUDA數(shù)組綁定的紋理參照系中的元素使用的N-元組數(shù)據(jù)中的組件數(shù)量必須與CUDA數(shù)組相同。紋理緩存有兩個(gè)作用。首先，紋理緩存中的數(shù)據(jù)可以被重復(fù)利用，當(dāng)一次訪問需要的數(shù)據(jù)已經(jīng)存在于紋理緩存中時(shí)

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

5、的值也與數(shù)據(jù)在線性內(nèi)存中的偏移量相同；而CUDA數(shù)組可以與一維、二維或者三維紋理綁定，紋理拾取坐標(biāo)是浮點(diǎn)型，并且支持許多特殊功能。紋理存儲(chǔ)器的特殊功能有：浮點(diǎn)型紋理拾取坐標(biāo)：使用浮點(diǎn)型的紋理拾取坐標(biāo)對(duì)紋理進(jìn)行尋址，只對(duì)與CUDA數(shù)組綁定的存儲(chǔ)器有效。地址映射的方式可以是歸一化或者非歸一化的：使用歸一化紋理時(shí)，紋理在每個(gè)維度上的坐標(biāo)被映射到浮點(diǎn)數(shù)0.0,1.0)范圍內(nèi)；使用非歸一化紋理坐標(biāo)時(shí)，各個(gè)維度上的坐標(biāo)則被映射到浮點(diǎn)數(shù)0.0,N)的范圍內(nèi)，其中N是紋理在該維度上像元的數(shù)量。由于在GPU中通常用浮點(diǎn)計(jì)算點(diǎn)的坐標(biāo)，因此使用浮點(diǎn)數(shù)作為紋理拾取坐標(biāo)更加自然；使用歸一化的紋理拾取坐標(biāo)可以不用關(guān)心紋理

6、的實(shí)際尺寸，簡化了渲染程序的編寫。尋址模式：尋址模式規(guī)定了紋理拾取的輸入坐標(biāo)超出紋理尋址范圍時(shí)的行為，有鉗位模式和循環(huán)模式兩種。使用鉗位模式時(shí)，當(dāng)輸入的坐標(biāo)超出了尋址范圍，輸入的值將被“鉗位”到尋址范圍的最大值或者最小值；循環(huán)模式只對(duì)歸一化坐標(biāo)有效，此時(shí)要對(duì)超出尋址范圍的紋理坐標(biāo)作求模等處理。例如，對(duì)映射到0.0,1.0)的歸一化紋理坐標(biāo)，輸入拾取坐標(biāo)1.25，鉗位模式會(huì)將輸入按照0.99999處理，而循環(huán)模式會(huì)將輸入0.25處理。類型轉(zhuǎn)換：如果像元中的數(shù)據(jù)是8-bit或者16-bit定點(diǎn)型，類型轉(zhuǎn)換功能對(duì)拾取的返回值進(jìn)行類型轉(zhuǎn)換，將其映射到歸一化的浮點(diǎn)范圍O.Of,1.0f(對(duì)無符號(hào)整型)或

7、者-1.0f,1.0f(對(duì)有符號(hào)整型)。濾波：如果將返回類型是浮點(diǎn)型的CUDA數(shù)組與紋理綁定，那么就可以對(duì)返回的值進(jìn)行濾波。濾波模式可以是最近點(diǎn)取樣模式或者線性濾波模式兩種。最近點(diǎn)模式返回與浮點(diǎn)型的紋理抓取坐標(biāo)最近像元的值，而線性濾波模式則會(huì)先取出附近幾個(gè)像元，然后按照抓取坐標(biāo)與這幾個(gè)像元的距離進(jìn)行線性插值，返回線性插值得到的值。線性濾波可以使紋理渲染得到的畫面更加平滑自然。線性濾波需要的插值計(jì)算不需要可編程單元參與，提供了額外的浮點(diǎn)處理能力，但精度較低。使用線性濾波模式返回的值經(jīng)過了插值處理，適合用于圖像處理；使用最近點(diǎn)取樣模式的返回值不會(huì)改變紋理中像元的值，適合用于實(shí)現(xiàn)查找表。關(guān)于紋理拾取

8、模式的詳細(xì)描述，可以參考附錄F。使用紋理存儲(chǔ)器時(shí)，首先要在主機(jī)端聲明需要綁定到紋理的線性存儲(chǔ)器或CUDA數(shù)組,并設(shè)置好紋理參照系，然后將紋理參照系與線性內(nèi)存或者CUDA數(shù)組綁定。在主機(jī)端完成配置工作后，就可以在內(nèi)核函數(shù)中通過紋理抓取函數(shù)訪問紋理存儲(chǔ)器了。1CUDA數(shù)組在顯存中可以分配的空間有兩種：CUDA數(shù)組和線性內(nèi)存。此外，常數(shù)存儲(chǔ)器中通過緩存加速讀取的數(shù)據(jù)實(shí)際也存在于顯存中。CUDA數(shù)組和線性內(nèi)存都可以與紋理參照系綁定，但CUDA數(shù)組對(duì)紋理拾取訪問進(jìn)行了優(yōu)化，在設(shè)備端也只能通過紋理拾取訪問。聲明CUDA數(shù)組之前，必須先以結(jié)構(gòu)體channelDesc描述CUDA數(shù)組中的數(shù)據(jù)類型。struct

9、cudaChannelFormatDescintx,y,z,w;enumcudaChannelFormatKindf;其中，x,y,z和w分別是每個(gè)返回值成員的位數(shù)，而f是一個(gè)枚舉變量，可以取一下幾個(gè)值：cudaChannelFormatKindSigned，如果這些成員是有符號(hào)整型；cudaChannelFormatKindUnsigned，如果這些成員是無符號(hào)整型；cudaChannelFormatKindFloat，如果這些成員是浮點(diǎn)型；然后，我們要確定CUDA數(shù)組的維度和尺寸。CUDA數(shù)組可以通過cudaMalloc3DArray()或cudaMallocArray()函數(shù)分配。用cu

10、daMalloc3DArray可以分配一維、二維或者三維的CUDA數(shù)組，而cudaMallocArray()般用于分配二維CUDA數(shù)組。在使用完CUDA數(shù)組后，要使用cudaFreeArray函數(shù)釋放顯存。由cudaMalloc3DArray分配的CUDA數(shù)組使用cudaMemcpy3D()完成與其他CUDA數(shù)組或者線性內(nèi)存的數(shù)據(jù)傳輸。CUDAAPI中使用結(jié)構(gòu)體cudaExtent描述3DArray和3D線性內(nèi)存在三個(gè)維度上的尺寸，在描述一維、二維和三維數(shù)組分別用以下的形式：cudaextentextent=make_cudaextent(1,8192,0,0);cudaextentexten

11、t=make_cudaextent(1,65535,1,32768,0);cudaextentextent=make_cudaextent(1,2048,1,2048,1,2048);其中方括號(hào)內(nèi)為允許的尋址范圍。注意到二維CUDA數(shù)組的第一個(gè)維度的尋址范圍大于一維CUDA數(shù)組的尋址范圍，因此在一維CUDA數(shù)組的尺寸不夠用時(shí)，將二維CUDA數(shù)組的第二個(gè)維度設(shè)為1代替一維CUDA數(shù)組，獲得更大的尋址范圍。下面是聲明一個(gè)數(shù)據(jù)類型為char2型，寬x高x深為64x32x16的CUDA3D數(shù)組，對(duì)其初始化，最后釋放數(shù)組的示例代碼：cudaChannelFormatDescchannelDesc=cud

12、aCreateChannelDesc(8,8,0,0,cudaChannelFormatKindunsigned);/每個(gè)像元由兩個(gè)charDOcudaExtentextent=make_cudaextent(64,32,16);/建立cudaExtent結(jié)構(gòu)體，描述CUDA數(shù)組的維度和尺寸cudaArray*cuArray;cudaMalloc3DArray(&cuArray,&channelDesc,extent);/為cuArray開辟空間缺cudaFreeArray(cuArray);下面則是使用cudaMallocArray聲明一個(gè)由float型構(gòu)成，尺寸為64x32的CUDA數(shù)組，

13、對(duì)其賦值，并最后釋放的示例代碼：cudaChannelFormatDescchannelDesc=cudaCreateChannelDesc(32,0,0,0,cudaChannelFormatKindunsigned);/每個(gè)像元由一個(gè)floatDocudaArray*cuArray;cudaMallocArray(&cuArray,&channelDesc,64,32);/為cuArrayDOODcudaMemcpyToArray(cuArray,0,0,h_data,&channelDesc);/第二和第三個(gè)參數(shù)分別表示在寬度和高度上的偏移量，假設(shè)h_data中的數(shù)據(jù)已經(jīng)初始化cudaF

14、reeArray(cuArray);用于在CUDA數(shù)組和主機(jī)端或者設(shè)備端線性內(nèi)存，以及在CUDA數(shù)組間傳輸數(shù)據(jù)的函數(shù)還有很多，這些還是還有一些異步調(diào)用版本，關(guān)于這些函數(shù)的具體使用方法請(qǐng)參考CUDAReferencemannual。2聲明紋理參照系紋理參照系中的一些屬性必須在編譯時(shí)之前被顯示聲明。紋理參照系通過一個(gè)作用范圍為全文件的texture型變量聲明：texturetexRef;其中，Type確定了由紋理拾取返回的數(shù)據(jù)類型；Type可以是B3.1節(jié)中描述的任意一種由基本整型或者單精度浮點(diǎn)型組成能的1-，2-或者4-元組向量類型。Dim確定了紋理參照系的維度，默認(rèn)為1。ReadMode可以是

15、cudaReadModeNormalizedFloat或者cudaReadModeElementType。如果ReadMode是cudaReadModeNomalizedFloat，并且Type是16-或者8-bit整型，那么返回的值將是一個(gè)浮點(diǎn)數(shù)。此時(shí)，原來整形的值域會(huì)被映射到0.0，1.0（對(duì)無符號(hào)整型），或者-1.0,1.0（對(duì)有符號(hào)整型）。例如，一個(gè)值為Oxff的8-bit無符號(hào)整型會(huì)被映射為1.0f。如果使用cudaReadModeElementType，那么就不會(huì)對(duì)輸出進(jìn)行轉(zhuǎn)換。ReadMode是一個(gè)可選參數(shù)，如果不寫，那么默認(rèn)就是cudaReadModeElementType。例

16、如，下面的代碼聲明了一個(gè)二維，像元數(shù)據(jù)為unsignedchar型，但將返回值轉(zhuǎn)換為float型的紋理參照系：texturetexRef;2.1設(shè)置運(yùn)行時(shí)紋理參照系屬性紋理參照系中的其它屬性可以不必聲明，并在運(yùn)行時(shí)進(jìn)行修改。這些參數(shù)規(guī)定了紋理的尋址模式，是否進(jìn)行歸一化，以及紋理濾波runtimeAPI擁有底層的C風(fēng)格和高層的C+風(fēng)格兩種接口。高層API中的texture類型是從底層的textureReference中派生而來的。TextureReference是一個(gè)下面的代碼描述的結(jié)構(gòu)體。structtextureReferenceintnormalized;enumcudaTextureFi

17、lterModefilterMode;enumcudaTextureAddressModeaddressMode3;structcudaChannelFormatDescchannelDesc;normalized設(shè)置是否對(duì)紋理坐標(biāo)是否進(jìn)行歸一化。如果normalized是一個(gè)非零值，那么就會(huì)使用歸一化到0,1）的坐標(biāo)進(jìn)行尋址，否則對(duì)尺寸為width,height,depth的紋理使用坐標(biāo)0,width-1,0,height-1,0,depth-1尋址。例如,一個(gè)尺寸為64x32的紋理可以通過x維度范圍為0,63,y維度范圍0,31的坐標(biāo)尋址。如果采用歸一化方式對(duì)尺寸為64x32的紋理進(jìn)行尋址

18、，在x和y維度上的坐標(biāo)就都是0.0,1.0）。這樣就可以保證紋理的坐標(biāo)與紋理的尺寸無關(guān)。filterMode用于設(shè)置紋理的濾波模式，即如何根據(jù)坐標(biāo)計(jì)算返回的紋理值。濾波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。濾波模式為CudaFilterModePoint時(shí)，返回值是與坐標(biāo)最接近的像元的值。CudaFilterModeLinear模式只能對(duì)返回值為浮點(diǎn)型的紋理使用，啟用這一種模式時(shí)將拾取紋理坐標(biāo)周圍的像元，然后根據(jù)坐標(biāo)與這些像元之間的距離進(jìn)行插值計(jì)算。對(duì)一維紋理可以使用線性濾波，對(duì)二維紋理可以使用雙線性濾波。返回值會(huì)是對(duì)最接近紋理坐標(biāo)的兩個(gè)

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

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

21、們只適用于與CUDA數(shù)組綁定的紋理參照系。附錄D中列出了關(guān)于紋理拾取的更多信息。紋理綁定在kerne1能用紋理參照系從紋理內(nèi)存中讀數(shù)據(jù)前，紋理參照系必須通過cudaBindTexture()或cudaBindTextureToArray()綁定到紋理上。cudaUnbindTexture()用于解除紋理參照系的綁定。以下代碼示例綁定一個(gè)紋理參照系到devPtr指向的線性內(nèi)存：使用低級(jí)API：texturetexRef;textureReference*texRefPtr;cudaGetTextureReference(&texRefPtr,“texRef”);cudaChannelFormat

22、DescchannelDesc=cudaCreateChannelDesc();cudaBindTexture2D(0,texRefPtr,devPtr,&channelDesc,width,height,pitch);使用高級(jí)APItexturetexRef;cudaChannelFormatDescchannelDesc=cudaCreateChannelDesc();cudaBindTexture2D(0,texRef,devPtr,&channelDesc,width,height,pitch);以下代碼示例綁定紋理參照系到一個(gè)CUDA數(shù)組cuArray：使用低級(jí)API：texture

23、texRef;textureReference*texRefPtr;cudaGetTextureReference(&texRefPtr,“texRef”);cudaChannelFormatDescchannelDesc;cudaGetChannelDesc(&channelDesc,cuArray);cudaBindTextureToArray(texRef,cuArray,&channelDesc);使用高級(jí)APItexturetexRef;cudaBindTextureToArray(texRef,cuArray);當(dāng)綁定一個(gè)紋理到紋理參照系時(shí)，格式必須與聲明紋理參照系時(shí)的參數(shù)匹配否則

24、，紋理獲取的結(jié)果是undefined的。紋理拾取紋理拾取函數(shù)采用紋理坐標(biāo)對(duì)紋理存儲(chǔ)器進(jìn)行訪問。對(duì)與線性內(nèi)存綁定的紋理，使用texfetchlD函數(shù)訪問，采用的紋理坐標(biāo)是整型。由cudaMallocPitch或者cudaMalloc3D分配的線性空間實(shí)際上仍然是經(jīng)過填充、對(duì)齊的一維線性空間，因此也用texfetchlD()函數(shù)訪問。對(duì)與一維、二維和三維CUDA數(shù)組綁定的問哪里，分別使用texlD()、tex2D()和tex3D()函數(shù)訪問，并且使用浮點(diǎn)型紋理坐標(biāo)。關(guān)于紋理拾取函數(shù)的更多討論，請(qǐng)見本書附錄D.82.4例子分析：Simpletexture/2Dfloattexturetexturet

25、exRef;/Simpletransformationkernel_global_voidtransformKernel(float*output,intwidth,intheight,floattheta)/根據(jù)tidbid計(jì)算歸一化的拾取坐標(biāo)unsignedintx=blockIdx.x*blockDim.x+threadIdx.x;unsignedinty=blockIdx.y*blockDim.y+threadIdx.y;floatu=x/(float)width;floatv=y/(float)height;/旋轉(zhuǎn)拾取坐標(biāo)u-=0.5f;v-=0.5f;floattu=u*cosf(theta)-v*sinf(theta)+0.5f;floattv=v*cosf(theta)+u*sinf(theta)+0.5f;/從紋理存儲(chǔ)器中拾取數(shù)據(jù)，并寫入顯存outputy*width+x=tex2D(tex,tu,tv);/Hostcodeintmain()/分配CUDA數(shù)組cudaChannelFormatDescchannelDesc=cudaCreateChannelDesc(3