cuda支援使用一小部分gpu用於圖形顯示的紋理硬體來對紋理記憶體和表面記憶體進行訪問。相比於從全域性記憶體,從紋理或者表面記憶體中讀取資料的在之前的裝置記憶體章節已經進行了介紹。
有兩類api 用來完成對紋理和表面記憶體進行使用: 紋理引用api在所有的裝置都能使用; 紋理物件api只支援在計算力大於3.0的裝置上使用。 紋理引用api有一些紋理物件api所有沒有的限制,使用時需要有所注意。
核可以使用裝置端的函式來讀紋理記憶體。呼叫乙個這樣的函式來讀取紋理的過程稱之為texture fetch。每次fetch都會指明乙個引數,texture object 或者 texture reference給相對應的api。
紋理物件或者紋理引用:
-紋理texture,是指fetch的一小片紋理記憶體。 紋理物件是在執行時建立的且同時對紋理進行明確。紋理引用是在編譯時建立並在執行時明確。一些間隔較遠的紋理引用可能會界定相同的紋理或者有記憶體重疊的紋理。紋理可以是一段線性的記憶體區域或者是乙個cuda陣列。
紋理物件的建立可以利用cudacreatetextureobject來實現,這個介面在呼叫的時候需要使用乙個稱為cudaresourcedesc的結構體, 它指明了紋理的一些屬性。
struct cudatexturedesc
;
下示**是紋理物件的應用場景之一:
// ****** transformation kernel
__global__ void transformkernel(float* output,
cudatextureobject_t texobj,
int width, int height,
float theta)
// host code
int main()
紋理引用的一些屬性是不可更改的,因此必須在編譯的時候就確定。它們將指明何時宣告紋理引用。可以使用下列方式宣告檔案範圍內(file scope)的紋理型變數:
texturetexref;
-datatype指明紋理元的型別
-type指明紋理引用的型別:cudatexture1d、cudatexturetype2d、cudatexturetype3d分別表示一維、二維和三維紋理;cudatexturetype1dlayered、cudatexturetype1dlayered分別表示一維和二維的layered紋理。這個引數的預設選項是cudatexture1d。
-readmode指明了讀取模式,預設值為cudareadmodeelementtype。
注意紋理引用只能被生命為靜態的全域性變數且並不能夠作為引數傳遞給函式使用。
紋理引用的其他變數時可變的, 可以通過通過主機的執行時在執行的時候進行修改。如參考手冊中所描述,執行時api有低階別的c風格介面和高階別的c++風格介面。紋理型別texture定義的乙個高階別api結構體,它派生與低級別api 中的紋理引用型別texturereference。
struct texturereference
struct cudachannelformatdesc ;
在核函式使用紋理引用來讀取紋理記憶體之前,需要使用相應的介面來將紋理和紋理引用進行繫結。對於線性記憶體使用cudabindtexture()或cudabindtexture2d(), 對於cuda 陣列使用cudabindtexturetoarray()。cudaunbindtexture()用於解綁乙個紋理引用。當乙個紋理引用解綁之後才能重新繫結其他陣列,即使之前的繫結在核函式中並沒有使用完畢也可以。
下列分別使用低階api 和高階api來展示如何進行繫結。
// 低階api
// example1
texturetexref;
texturereference* texrefptr;
cudagettexturereference(&texrefptr, &texref);
cudachannelformatdesc channeldesc =
cudacreatechanneldesc();
size_t offset;
cudabindtexture2d(&offset, texrefptr, devptr, &channeldesc,
width, height, pitch);
// example2
texturetexref;
texturereference* texrefptr;
cudagettexturereference(&texrefptr, &texref);
cudachannelformatdesc channeldesc;
cudagetchanneldesc(&channeldesc, cuarray);
cudabindtexturetoarray(texref, cuarray, &channeldesc);
//高階api
// example1
texturetexref;
cudachannelformatdesc channeldesc =
cudacreatechanneldesc();
size_t offset;
cudabindtexture2d(&offset, texref, devptr, channeldesc,
width, height, pitch);
// example2
texturetexref;
cudabindtexturetoarray(texref, cuarray);
注意, 在將乙個紋理繫結至紋理引用的時候其指定的格式必須與宣告紋理引用的指定是相匹配的; 否則紋理fetch的結果是未定義的。同時核函式所能同時繫結的紋理個數是有限制的。
下示例**中對紋理使用了一些簡單的變化操作。
// 2d float texture
texturetexref;
// ****** transformation kernel
__global__ void transformkernel(float* output,
int width, int height,
float theta)
// host code
int main()
cuda 陣列支援16位浮點數, 即半精度。cuda c 並不提供相匹配的資料型別,但是提供了內部函式來利用無符號短整型unsigned short進行從/到32位浮點數的轉換:__float2half_rn(float)、__half2float(unsigned short)。這些函式只能在裝置端函式內使用。相應的主機端等價函式可以在openexr庫中找到。
在紋理fetching的時候, 執行filtering之前會將16位浮點分量提公升至32位浮點型別。16位浮點形式的通道描述符可以通過呼叫cudacreatechanneldeschalf*()來進行建立。
一維或二維的layered紋理(在direct3d中陳偉紋理陣列texture array, 在opengl中稱為陣列紋理array texutre)是有連續的層構成,每層都是具有相同維度、大小和資料型別的矩形紋理構成。
一維的layered紋理的定址是通過乙個整數索引和乙個浮點紋理座標, 索引表示層的位置而座標表示在該層中紋理元的位置; 二維的layered紋理的定址是通過乙個整數索引和兩個浮點座標來進行的, 索引表示層的位置,而兩個浮點數座標表示紋理元在該層中的位置。
layered 紋理只能通過在呼叫cudamalloc3darray()的時候傳遞標誌cudaarraylayered建立。layered紋理的fetching需要通過裝置端函式tex1dlayered() 和tex2dlayered()來進行。紋理的fetching只能在層中進行,不會發生跨層。
layered紋理只在計算力高於2.0的裝置上才能後使用。
cubemap紋理是一種特殊的二維layered紋理。它由6層組成,代表了乙個cube的各個面。
cubemap紋理只能通過在呼叫cudamalloc3darray()的時候傳遞標誌cudaarraycubemap建立。cubemap紋理的fetching需要通過裝置端函式texcubemap() 和texcubemap()來進行。
cubemap紋理只在計算力高於2.0的裝置上才能後使用。
紋理聚合使用特殊的紋理fetch, 它只適用於二維的紋理。通過呼叫tex2dgather()來實現,它和tex2d()具有相同的引數 ,再加上乙個額外的值域為的comp引數。它返回四個32位的數字,對應於四個紋理元的分量值, 用於在常規紋理fetch時進行雙線性filtering。 例如, 如果紋理元的值為(253, 20, 31, 255), (250, 25, 29, 254), (249, 16, 37, 253), (251, 22, 30, 250), comp的值是2,那麼tex2dgather返回的值是(31, 29, 37, 30)。
注意,紋理座標的計算只有8位的小數精度。tex2dgather()因此可能會返回 與tex2d()不同的結果。
CUDA C程式設計手冊 程式設計介面(六)
對於計算力高於2.0的裝置,使用cudaarraysu celoadstore標誌建立的cuda陣列,可以通過su ce object.或者su ce reference進行讀寫。不同的裝置,所支援的表面記憶體的大小是不一樣的。表面物件的描述可以使用struct cudaresourcedesc來表...
Qt介面程式設計(五)
使用網路通訊模組前,要在.pro檔案中新增network模組。1 建立qudpsocket物件 2 繫結ip的埠號 3 傳送資料 qint64 writedatagram const char data,qint64 len,const qhostaddress host,quint16 port ...
《CUDA C程式設計權威指南》 2 6 習題
1.在檔案sumarraysongpu timer.cu中,設定block.x 1 023,重新編譯並執行。與執行配置為block.x 1 024的執行結果進行比較,試著解釋其區別和原因。2.參考檔案sumarraysongpu timer.cu,設定block.x 256。新建乙個核心,使得每個執...