CUDA優化例項（四）紋理記憶體

本文參考：

本節的內容可能和標題不服，本節主要將紋理記憶體的使用，它到底有什麼有什麼速度的提公升，優化體現在**，我下節會寫，本節主要寫乙個紋理記憶體的例子。

紋理記憶體的使用有兩種方式，分別是texture object和texture reference，前者可在程式中動態生成，後者的部分在編譯期間靜態生成（確定了），有點像靜態陣列和動態陣列的意思。如當事先不知道要處理的資料的型別時，使用texture object是個不錯的選擇。本文以texture reference為例。

使用紋理記憶體步驟較複雜簡單介紹：

宣告texture referece：

texturetexref;申請裝置記憶體，可以是線性記憶體或cuda記憶體，其不同[cuda紋理筆記]

( 這裡以2維cuda array為例：

// 申請 cuda 陣列並拷貝資料

cudachannelformatdesc channeldesc = cudacreatechanneldesc(8, 0, 0, 0, cudachannelformatkindunsigned);
cudaarray* cuarray;
cudamallocarray(&cuarray, &channeldesc, width, height);
cudamemcpytoarray(cuarray, 0, 0, h_data, size, cudamemcpyhosttodevice);

cudamemcpytoarray是將host段的資料拷貝到cuda array。

3. 設定texture reference引數：

// 指定紋理引用引數

//一維的超越邊界取0
texref.addressmode[0] = cudaaddressmodeborder;
//二維的超越邊界取0
texref.addressmode[1] = cudaaddressmodeborder;
//不插值，取整
texref.filtermode = cudafiltermodepoint;
//不使用歸一化紋理座標
texref.normalized = 0;

// 繫結紋理引用

cudabindtexturetoarray(texref, cuarray, channeldesc);

texture memory.cu：

#include 

#include 
#include 
#include 
#include "device_launch_parameters.h"
#define degre_to_radian(x) ((x) * 3.1416f / 180)
#define ceil(x,y) ((((x) + (y) - 1))/ (y) )
typedef unsigned char uchar;
// 宣告紋理引用
texturetexref;
// 簡單的線性變換
__global__ void transformkernel(uchar* output, int width, int height, float theta)
intmain
()    // 申請 cuda 陣列並拷貝資料
cudachannelformatdesc
channeldesc = cudacreatechanneldesc
(8, 0, 0, 0, cudachannelformatkindunsigned);
cudaarray* cuarray;
cudamallocarray
(&cuarray, &channeldesc, width, height);
cudamemcpytoarray
(cuarray, 0, 0, h_data, size, cudamemcpyhosttodevice);
// 指定紋理引用引數，注意與紋理物件的使用不一樣
texref.addressmode[0] = cudaaddressmodeborder;
texref.addressmode[1] = cudaaddressmodeborder;
texref.filtermode = cudafiltermodepoint;
texref.normalized = 0;
// 繫結紋理引用
cudabindtexturetoarray
(texref, cuarray, channeldesc);
// 執行核函式
dim3
dimblock
(16, 16);
dim3
dimgrid
(ceil(width, dimblock.x), ceil(height, dimblock.y));
printf
("x:  %d , y:  %d  \n",dimgrid.x,dimgrid.y);
transformkernel
<< > > (d_data, width, height, angle);
cudadevicesynchronize
();    // 結果**和檢查結果
cudamemcpy
(h_data, d_data, size, cudamemcpydevicetohost);
printf
("\n\n");
for(i = 0; i < width*height; i++)
// **工作
cudafreearray
(cuarray);
cudafree
(d_data);
return 0;
}

當angle = 30時

當angle = 0時

180旋轉出現0是因為texref.addressmode[1] = cudaaddressmodeborder;和#define degre_to_radian(x) ((x) * 3.1415926 / 180)pi精度不夠，

#define degre_to_radian(x) ((x) * acos(-1) / 180)會好一些。

本節介紹個紋理記憶體的例子，下節會介紹紋理記憶體比全域性記憶體的優越性。

CUDA分支優化

標籤 cuda 分支優化 warpsm 2015 07 16 10 24 293人閱讀收藏 舉報 cuda 26 在cuda中，分支會極大的減弱效能，因為 沒有分支 因此只能讓束內線程在每個分支上都執行一遍，當然如果某個分支沒有執行緒執行，就可以忽略，因此要減少分支的數目。可以簡單的說 1.同乙個w...

CUDA 學習優化思路

參考 cpu中memory l3 cache傳輸頻寬為20gb s,除以64bytes line得到傳輸記錄速度約300m line s，約為300m 8 2.4g double s.一般地，浮點數操作需要兩個輸入 1個輸出，那麼loading 3個數 3 lines 的代價為 100mflops。...

cuda程式設計 效能優化

做了幾個月的gpu的效能優化，一直沒來得及總結，先把大概要點總結一下，後面有空進行進一步完善。tesla v100 cuda 9.0 linux c 資料格式要進行設計，把每個部分精簡至最小，減少對gpu視訊記憶體的占用，同時縮小資料的比對時間。從記憶體load到快取，減少load的時間消耗，可採用...