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

[a5000ml (咖啡裡的海洋藍)]

終於要步入最佳化的部份, 實在令人興奮... 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)
※ a5000ml:轉錄至看板 C_and_CPP                                    12/11 01:47
推 vip82:有看有推
推 yayax:推!!
推 VictorTom:推;   您才說下的差不多, 這兩天就有新人加入下載了XD
推 CDavid:推
推 tamato9519:看不懂推 XD
推 sellor:猛..(裝懂)推....(暈！謎！)
推 VictorTom:OpenCL 1.0 spec出來了, 好想哭....Q_Q~
推 lavatar:openCL跟CUDA沒衝突喔,請見 http://0rz.tw/d25fg
→ VictorTom:小弟沒有說衝突啦, 只是spec出來就會被assign去看了Q_Q~
推 damnc:只看標題就先推了...
→ a5000ml:Tom 大加油 ^_^
推 wch6858:推