[分享] CUDA 程式設計(13) -- 材質快取 (一)

看板C_and_CPP (C/C++)作者a5000ml (咖啡裡的海洋藍)時間16年前 (2008/12/11 01:47)推噓2(2推 0噓 1→)

留言3則, 3人參與討論串1/1

※ [本文轉錄自 VideoCard 看板] 作者: a5000ml (咖啡裡的海洋藍) 看板: VideoCard 標題: [分享] CUDA 程式設計(13) -- 材質快取 (一) 時間: Thu Dec 11 01:38:43 2008 終於要步入最佳化的部份, 實在令人興奮... XDD =========================================================================== 材質參考 (part.1) =========================================================================== ◆簡介 -------------------------------------------------------------------- 「材質參考」(Texture Reference) 其實就是開啟 GPU 的「快取機制」(cache)。傳統 CPU 在記憶體的讀取上，預設是有快取的，並提供一些指令來控制後續的指令資料流的快取 (例如: x86 的預先拮取 prefetch)，然而這種方式，程式設計師必需事先計算資料從 DRAM 預先被拮取所需的時間，並在適當的地方穿插 prefetch 指令，可是硬體的進展使得預先拮取的時間變得不確定，所以 prefetch 的效果並不具備向下相容性，使得在某型號的機器上好不容易調到最佳效能，換新機器後就不見得。而「材質參考」是完全不一樣的設計，源自於繪圖所需的材質貼圖管線 (texture pipeline)，因為在貼材質時需要快速的位址計算，以及快速拮取 2D 相鄰的資料，所以顯示卡中早已內建多個位址計算單元和 2D 的硬體拮取來應付材質貼圖的需求，到了 CUDA 時代才把它化為快取機制而己，在硬體的指令設計上，也和傳統 CPU 完全不同，採用讀取材質和直接讀取使用不同指令集的方式，去除了 prefetch 這種效果不佳的指令，改由程式執行之初的快取方式設定，支援多維度和歸一化浮點數的定址，且讓多個 multiprocessors 可以共用，增加資料在快取中的運用率，這些設計保留未來硬體調整的彈性，使得「材質參考」效果具備向下相容性。 ◆兩種不同的最佳化風格 (合併存取 vs. 材質參考) -------------------------------------------------------------------- 「合併存取」可視為是 RISC 的延伸，因為資料要合併存取才能最佳化，會迫使程式設計師先把一塊連續資料載入處理器中，處理完後再把這些資料做一次輸出，形成了所謂 Load/Save 的程式風格，以及大量暫存器的需求，而這種風格與需求正是 RISC 的本意。「材質參考」本質為快取機制，如果資料剛好在快取中，那就 Lucky 載入很快，如果沒有，那就要靠快取處理單元，將資料從記憶體中讀入，然後再餵給指令，因為使用上太便利了，程式設計師不用管這些細節，會養成在需要資料時才載入，那就是 CISC 的架構所推崇的方式。事實上 GPU 本身是 RISC 架構，具備可控制的快取機制，只是在 CUDA 這個以 C/C++ 為基礎的延伸語言中，提供了兩種不同的最佳化方式讓使用者自行調整，可完全使用「材質參考」寫程式 (會比較便利，並具備向下相容性)，也可選擇「合併存取」來做進階最佳化，或是兩者的混合達成平衡。 ◆效率因素 -------------------------------------------------------------------- 使用材質參考時最重要的是避免 __syncthreads()。單純的材質參考會比直接讀取慢，因為材質位址單元還要經過一些運算，以及載入時會一併載入相鄰的資料，但合併讀取是由執行緒分開發出的，在許多情況下 (執行緒非獨立串流) 還要透過共享記憶體來交換相鄰的資料，所以必需 __syncthreads()，這樣會使得執行緒停擺率增加，且資料載入無法和運算重疊 (除非能塞入足夠的運算在 __syncthreads() 之前)，失去了使用多執行緒來隱藏資料載入延遲的好處。而材質快取是跨執行緒的(甚至跨區塊)，不需透過共享記憶體交換相鄰的資料，理論上可以完全不用 __syncthreads()，使得資料載入可以很好的和運算重疊，多執行緒也因此可以有效的隱藏資料載入延遲，而且少用共享記憶體，也可以增加多處理器同時執行的區塊數。實際使用時，兩者要配合，因為 cache 空間有限，有些不需和相鄰執行緒交換的資料，那就用直接載入，原則是儘量不觸發 __syncthreads()，如此可避免過多的材質參考使得快取效能大幅下降。 ◆注意事項 -------------------------------------------------------------------- (1)材質記憶體是「唯讀快取」，本身在硬體上並不是的獨立的記憶晶片，而是顯示卡的 DRAM (global memory)，配合 GPU 上的 cache 來存取而己。 (2)在 CUDA 中預設的全域記憶體存取是沒有快取的，只有透過材質參考來存取，才能使用 GPU 的讀取快取，兩者的讀取指令並不相同。 (3)材質參考可被設定成較適合的方式以增進存取效率，現階段 nv 提供 1D、2D、3D 三種模式，1D 的可以直接綁在全域記憶體上，其餘 2D 和 3D 必需綁 CUDA Array，這東西其實和也和全域記憶體一樣位在顯示卡的 DRAM 上，只是有特殊的位址對齊符合 2D 和 3D 的要求，使得材質快取可以有效率的進行而己。 (4)材質參考提供多種定址方式，除了整數位址之外，還提供歸一化後的浮點數位址 (速度和整數位址差不多)，在某些應用上可節省位址計算的時間，並達到 fuzzy 定址的效果。 (5)材質參考可以跨多個 multiprocessors 共享，增加資料的利用率 (和硬體版本有關 G80/G90 系列的為兩個多處理器共用，而 G200 為 3 個共用)。 ◆使用方式 -------------------------------------------------------------------- 使用材質快取只需三個步驟： (1)宣告材質的參考方式 (2)在呼叫 kernel 前的 host 程序中，做材質參考的綁定 (3)在 kernel 中使用材質參考。先介紹 1D 對全域記憶體的材質參考，然後在範例中測試它的效能： (1)宣告材質參考的物件【語法】 texture<資料型態, 維度, 定址模式> 材質物件名稱; 資料型態：必需是基本的資料型態，或位址對齊的基本向量 (含 2 或 4 個元素) 例如 float, int, int4, float2, uint4 等維度：現階段可以指定 1、2、3 維。全域記憶體只能指定 1 維。定址模式：可以是 cudaReadModeElementType 或 cudaReadModeNormalizedFloat 前者代表傳統的整數位址，後者代表歸一化過的浮點數位址，全域記憶體只能用整數位址。【範例】 texture<float, 1, cudaReadModeElementType> tex; 宣告材質快取對應的陣列元素之資料型態為 float，快取的方式為 1D，使用整數位址存取，材質物件名稱為 tex (2)在 host 程序中使用 API 對材質參考做綁定【語法】 cudaBindTexture(偏移量, 材質物件, 全域記憶體位址, 記憶體大小); 偏移量：傳回所需的偏移，和記憶體對齊有關。(這裡可以先不管它) 材質物件：填入前面宣告的材質物件全域記憶體位址：填入要綁定的記憶體位址記憶體大小：填入要綁定的記憶體大小 (byte) 【範例】 float* aaa; cudaMalloc((void**)&aaa,N*sizeof(float)); ... cudaBindTexture(0, tex, aaa, N*sizeof(float)); 把全域記憶體 aaa 綁定 N 個 float 到 tex 這個材質物件。 (3)在 kernel 中使用材質參考。【語法】 tex1Dfetch(材質名稱, x); //全域記憶體只能用一維整數位址 tex2Dfetch(材質名稱, x, y); tex3Dfetch(材質名稱, x, y, z); 【範例】比較直接讀取全域記憶體 b=aaa[k]; 透過材質參考讀取 (tex 綁定到 aaa) b=tex1Dfetch(tex, x); ◆範例一：1D Laplace (d/dx)^2 的差分算符 -------------------------------------------------------------------- #include <stdlib.h> #include <stdio.h> #include <math.h> #include <cuda.h> //對照函數. void laplace(float* y, float* x, int n){ for(int i=1; i<n-1; i++){ y[i]=x[i+1]-2*x[i]+x[i-1]; } //periodic boundary y[0]=x[1]-2*x[0]+x[n-1]; y[n-1]=x[0]-2*x[n-1]+x[n-2]; } //使用合併讀取+共享記憶體來做. __global__ void ker_laplace_shared(float* y, float* x, int n){ int t=threadIdx.x; int b=blockIdx.x*blockDim.x; int i=b+t; __shared__ float sm[512+2]; float* s=sm+1; if(i<n){ s[t]=x[i]; } //載入區塊邊界點 (使用兩個 warp 來分散). if(t==0){ if(blockIdx.x==0){ s[-1]=x[n-1]; } else{ s[-1]=x[b-1]; } } if(t==32){ if(n-b<=blockDim.x){ s[n-blockIdx.x*blockDim.x]=x[0]; } else{ s[blockDim.x]=x[(blockIdx.x+1)*blockDim.x]; } } __syncthreads(); if(i<n){ y[i]=s[t+1]-2*s[t]+s[t-1]; } } //最原始的 kernel. __global__ void ker_laplace_naive(float* y, float* x, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; if(i==0){ y[0]=x[1]-2*x[0]+x[n-1]; } else if(i<n-1){ y[i]=x[i+1]-2*x[i]+x[i-1]; } else if(i==n-1){ y[n-1]=x[0]-2*x[n-1]+x[n-2]; } } //宣告材質物件. texture<float, 1, cudaReadModeElementType> texRefX; //使用材質快取的 kernel. __global__ void ker_laplace_texture(float* y, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; //用 macro 節省篇打字. #define xx(k) tex1Dfetch(texRefX, k) if(i==0){ y[0]=xx(1)-2*xx(0)+xx(n-1); } else if(i<n-1){ y[i]=xx(i+1)-2*xx(i)+xx(i-1); } else if(i==n-1){ y[n-1]=xx(0)-2*xx(n-1)+xx(n-2); } #undef xx } //亂數向量產生器. void vec_gen(float* vec, int size){ for(int i=0; i<size; i++){ vec[i]=(float)rand()/RAND_MAX*2-1; } } //比對兩向量的相對誤差. double diff(float* v1, float* v2, int size){ double sd=0; double sv=0; for(int i=0; i<size; i++){ double d=v1[i]-v2[i]; double v=v1[i]; sd+=d*d; sv+=v*v; } return sqrt(sd/sv); } //計萛時間差. double diff(timespec& t1, timespec& t2){ return (double)(t1.tv_sec-t2.tv_sec) + (double)(t1.tv_nsec-t2.tv_nsec)*1e-9; } //主函數. int main(){ int loops=503; //test loops int size=1024*1024; srand(time(0)); timespec ts1,ts2; float* v1=new float[size]; float* v2=new float[size]; float* v3=new float[size]; printf("------------------------------\n"); printf("1D Laplace Operator (periodic)\n"); printf(" vector size : %dK\n",size/1024); printf(" average loop : %d\n",loops); printf("------------------------------\n"); //---- generate profile in host----- vec_gen(v1,size); memset(v2,0,size*sizeof(float)); memset(v3,0,size*sizeof(float)); //---- set device memory ----- float *gv1,*gv3; cudaMalloc((void**)&gv1,(size+2)*sizeof(float)); cudaMalloc((void**)&gv3,(size+2)*sizeof(float)); cudaMemcpy(gv1,v1, size*sizeof(float), cudaMemcpyHostToDevice); //綁定材質參考 cudaBindTexture(0, texRefX, gv1, size*sizeof(float)); double dtdev,dthost; //---- test host performance ----- clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ laplace(v2, v1, size); } clock_gettime(0,&ts2); dthost=diff(ts2,ts1)*1000/loops; printf("time(host): %g ms\n",dthost); //---- test naive device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_naive<<<size/512+1,512>>>(gv3,gv1,size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(naive): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); //---- test texture device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_texture<<<size/512+1,512>>>(gv3, size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(texture): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); //---- test shared device performance ----- memset(v3,0,size*sizeof(float)); cudaMemset(gv3,0,size*sizeof(float)); cudaThreadSynchronize(); clock_gettime(0,&ts1); for(int k=0; k<loops; k++){ ker_laplace_shared<<<size/512+1,512>>>(gv3,gv1,size); } cudaThreadSynchronize(); clock_gettime(0,&ts2); dtdev=diff(ts2,ts1)*1000/loops; cudaMemcpy(v3,gv3, size*sizeof(float), cudaMemcpyDeviceToHost); printf("time(shared): %g ms (%g x) error:%g\n",dtdev, dthost/dtdev,diff(v2,v3,size)); cudaFree(gv1); cudaFree(gv3); delete [] v1; delete [] v2; delete [] v3; return 0; } ◆結果一 -------------------------------------------------------------------- 這次測試使用 GTX260 vs. Intel E8400，其結果如下： ------------------------------ 1D Laplace Operator (periodic) vector size : 1024K average loop : 503 ------------------------------ time(host): 5.62482 ms time(naive): 0.246039 ms (22.8615 x) error:4.02879e-08 time(texture): 0.117252 ms (47.972 x) error:4.02879e-08 time(shared): 0.166139 ms (33.856 x) error:4.02879e-08 ------------------------------ ker_laplace_texture() 基本上只把 ker_laplace_naive() 改成材質快取而己，裡面沒有 __syncthread()，大家可以看到它飆得超快，將近 host 的 50x，而 shared 版本寫得超辛苦，結果還是比開啟快取的版本慢上許多。 ◆範例二：在 ker_laplace_texture() 插入 __syncthreads() -------------------------------------------------------------------- __global__ void ker_laplace_texture(float* y, int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; #define xx(k) tex1Dfetch(texRefX, k) //先載入資料. float a,b,c; if(i==0){ a=xx(1); b=xx(0); c=xx(n-1); } else if(i<n-1){ a=xx(i+1); b=xx(i); c=xx(i-1); } else if(i==n-1){ a=xx(0); b=xx(n-1); c=xx(n-2); } //像 shared 版一樣在資料讀取和運算中間插入一個同步指令. __syncthreads(); //計算. y[i]=a-2*b+c; #undef xx } ◆結果二：在 ker_laplace_texture() 插入 __syncthreads() -------------------------------------------------------------------- 結果如下，我們可看出 texture 效果顯著變差了，幾乎和 shared 一樣慢，可見 __syncthread() 的確會對材質的效能造成很大的衝擊，這裡會稍微比 shared 快是因為 texture+sync 版少了載入邊界的這個動作。 ------------------------------ 1D Laplace Operator (periodic) vector size : 1024K average loop : 503 ------------------------------ time(host): 5.62484 ms time(naive): 0.24602 ms (22.8634 x) error:4.02361e-08 time(texture): 0.156428 ms (35.958 x) error:4.02361e-08 (sync 版) time(shared): 0.166084 ms (33.8675 x) error:4.02361e-08 ------------------------------ 這個測試証實了我們前面的理論分析，所以要儘量避免 __syncthreads()，才能讓 texture 發揮最大功效。 -- 好像國網光碟大家都下載得差不多了, 現在用的人好像還不多, 沒關係~~ 讓我們繼續聊, 好東西值得繼續發展 ^^y -- 。o O　○。o O　○。o O　○。o O　○。o O　○。o 國網 CUDA 中文教學 DVD 影片 (免費線上版) 請至國網的教育訓練網登入 https://edu.nchc.org.tw BT 牌的種子下載點 http://www.badongo.com/file/12156676 http://rapidshare.de/files/41036559/NCHC_CUDA_video.torrent.html http://www.btghost.com/link/54915319/ -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.209.244 ※ 編輯: a5000ml 來自: 114.45.209.244 (12/11 01:43) -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 114.45.209.244