GPU程式設計之進擊的優化系列四--GPU程式設計優化技術總結

第四章 GPU程式設計優化技術總結

4.1.0 CUDA裝置上的優化技術

4.1.1 訪存優化

4.1.2 指令優化

4.1.3 核心調用優化

4.2.0 GCN裝置上的優化技術

4.2.1 訪存優化

4.2.2 指令優化

4.2.3 核心調用優化

4.3 建構性能可移植的程式

我們在兩個章節分别講述針對CUDA和GCN這兩大目前主流的GPU并行計算的裝置。但是諸如合并通路,如何避免共享記憶體的bank conflicts以及簡單的指令優化等基本内容這裡不再叙述，有需要的可以參考<<CUDA Programming Guide>>和<<AMD Accelerated Parallel ProcessingOpenCL Programming Guide>>，這裡僅給出一些不常見的優化技巧。

4.1.0  CUDA裝置的優化技術

     1 在計算能力為2.0或以上的CUDA裝置上，當一個warp内的所有線程通路同一個位址時，可以使用統一加載操作将一個資料通過緩存廣播到warp内的所有線程中，進而提升性能。雖然CUDAprogramming guide上提到當通路資料的位址和線程号無關且是隻讀資料時，編譯器會自動使用LDU加載指令，但有時編譯器并不能得到我們想要的結果。比如：

__global__ void add( float* d_a, const float* d_b, int n)

            {

                        intwarpid=(blockDim.x>>5)*blockIdx.x+(threadIdx.x>>5);

                        if(warpid>=n)return;

                        d_a[threadIdx.x]+=d_b[warpid];

            }

檢視PTX代碼，我們發現編譯器并未使用LDU，是以我們就需要顯示的使用内聯PTX彙編來達到我們的目的：

#if defined(_WIN64)|| defined(__x86_64)||defined(_M_X64)|| defined(_M_IA64)||defined(_M_AMD64)

#define PTX_PTR “l”

#else

#define PTX_PTR “r”

#endif

            __device____forceinline__ float __ldu(const float* p )

{

                        float val;

                        asm volatile("ldu.global.f32 {%0}, [%1];" : "=f"(val) : PTX_PTR(p));

                        returnval;

}

  根據測試，在滿足使用LDU的情況下的所有裝置中均能獲得性能提升，即使是在計算能力3.5+的裝置上其效果也要略好于使用LDG(使用紋理緩存)。例如卷積神經網絡的計算中，當每個通道中對應的是一個标量的偏置值，那麼在卷積計算後對通道施加偏置的操作就可以通過LDU操作高效的完成（雖然也可以通過共享記憶體，但是使用LDU具有更簡潔的實作，并具有輕微的性能優勢）。

4.1 指令優化

      1 對于可以完成相同計算的指令集合，應盡可能選擇具有更低延遲以及更高混合比例的指令集合，比如在某些裝置上雙精度可以和記憶體加載存儲指令雙發，但是卻無法和單精度以及整數指令雙發。

2 同時對同一個數組進行多次等距尋址時盡量将不變的索引在開始處加到數組的基址上，這樣可以減少位址的計算或是便于基址+常量尋址，進而減少指令數量。

3 對于存在大量計算的循環中如果某些指令，如資料存取指令無需複雜的尋址計算，那麼考慮對每個存儲操作使用斷定，便于編譯器将計算和存儲指令混合排列進而利用指令的雙發(dualissue)機制。

4 使用某些特定的常量，可以将資料融入指令碼中，進而具有更小的代碼體積。在kepler和maxwell裝置上的32位浮點數和整數的雙操作數（輸入）指令支援全精度的常量，比如

            c=a+128.f   對應的SASS指令為 FADD   R2，R0，128,

            c=a+10007.f對應的SASS指令為 FADD32I R2，R0，10007，

    這些立即數會被嵌入指令的編碼序列中，但是對于三操作數指令（如FMA）則會将常數放入常量記憶體的第2個bank中（猜測原因是受限于指令編碼的長度，因為多出的一個操作數需要額外的位數表示寄存器索引）,是以當一個計算序列中使用FMA不能減少指令數量時（亦即和使用FMUL,FADD數量相同），如果涉及到立即數，則盡量不要使用FMA代替FMUL和FADD，因為操作數直接嵌入指令編碼具有更小的代碼體積，除非出于精度考慮。如将

temp.x=c*b.x+(-s)*b.y;            

temp.y=c*b.y+s*b.x;   

b.x=a.x-temp.x;                       

b.y=a.y-temp.y;                       

a.x+=temp.x;               

a.y+=temp.y;               

轉換為

temp.x=b.x;     

temp.y=b.y;     

b.x=a.x+(-c)*temp.x+s*temp.y;                       

b.y=a.y+(-c)*temp.y+(-s)*temp.x;                    

a.x+=c*temp.x+(-s)*temp.y;               

a.y+=c*temp.y+s*temp.x;                    

并不能減少指令數量，也不會帶來性能提升，除非是出于精度考慮。對于雙精度資料，有規律的常量也可被嵌入到指令碼中，比如0.5,0.25,0.125,0.0625,0.03125,…,1.0,1.5, 1.25,…, 64.0, 128.0, 65536.0, …;但是無規則的常量會被放入常量記憶體的第2個bank中，比如

            c=a+128.0      對應的SASS指令為 DADD R2,  R7,  128

            c=a+790045.7對應的SASS指令為 DADD  R2, R7,  c[0x2][0x0]

可以包含在雙精度指令碼中的常量的具體規則為：

            …

            +-512.0

            +-256.0

            +-128.0,+-128.5

            +-64.0,+-64.5, 64.25

            +-32.0,+-32.5, +-32.25, +-32.125

            +-16.0,+-16.5, +-16.25, +-16.125, +-16.0625

4.1.2 分支優化

      1 使用小的局部數組消除多分支或是簡化複雜的條件代碼計算。

      2 巧妙的利用位操作和局部數組消除分支，例如第三章中通過局部數組簡化了主分割面的選擇，而通過巧妙的位操作減少了确定下個待周遊節點的分支。

3 分析算法看是否能将不同的路徑配置設定到不同的warp中或block中，同時保證warp或block中的指令路徑相同；或是将問題進行拆分成多個核心進行處理。

      4 在多分支結構中将判據按照命中的機率從高到低進行排列。

      5 某些情況下使用對函數指針清單的尋址代替switch邏輯已消除對大批量分支判斷的周遊，同時可以生成更小的代碼。         

    通常不是限制性能的地方，但是當很多核心在一個循環中被比較長時間的跨距調用時（是以裝置驅動的熱身會被過長的間隔抹消或是被其它核心的調用覆寫掉核心參數緩存），尤其是當核心具有很多參數時，每次核心參數都需要從記憶體到裝置上的核心參數緩存的複制過程，有時這也會給效率帶來較大影響，這裡根據作者經驗總結了幾個方法來優化核心的啟動時間：

    1 如果核心參數很多，對于指針類型的參數，考慮合并多個指針變量，并在核心内部解引用，這樣做有時也會減輕寄存器壓力（但也不要想當然，任何時候都應該試着檢視編譯後的寄存器使用情況）。例如，假設四個長度均為1024的數組(指針合并并不要求每個指針指向的數組大小一樣)：

            __global__ void …( …,const int * d_a, const int * d_b, const int * d_c, const int * d_d, … )

                        …

                        d_a+=tidx;

                        d_b+=tidx;

                        d_c+=tidx;

                        d_d+=tidx;

          可以改成如下形式：

            __global__ void … ( …, const int * d_a, … )

                        const int *d_b=d_a+1024;

                        const int *d_c=d_b+1024;

                        const int *d_d=d_c+1024;

                    /*

                    或者通過對d_a的常量偏移分别通路各個數組：

                      d_a,  d_a+1024，d_a+2048,  d_a+3072

                    */

     2  從PTX指令到本地SASS彙編指令并不是嚴格一一對應的，在這個翻譯的過程中ptxas會進行實際的寄存器配置設定，指令的替換和重排等優化，是以很多時候你無法通過使用PTX達到控制指令執行順序和寄存器配置設定的目的（ptxas做的并不夠好，一個實際的例子就是對于矩陣乘法，如果想要達到接近峰值的效率，必須直接對SASS指令進行重拍以及對寄存器進行細緻的配置設定以最小化指令計算延遲和寄存器bankconflicts引起的指令流水線停頓。但是很可惜，NVIDIA并未開放本地彙編的程式設計環境，甚至連SASS ISA的指令編碼格式都未公開，是以需要程式員自己繞開種種限制開發自己的第三方GPU彙編器），但是PTX仍然能在一定程度上影響最終得到的SASS結果，這需要開發者耐心的對指令的順序和邏輯進行調整并觀察最終編譯出來的SASS代碼。

    3  建立CUDA上下文時使用CU_CTX_LMEM_RESIZE_TO_MAX标志，以避免那些具有寄存器溢出的核心在下次啟動時重新在裝置記憶體上為寄存器溢出配置設定局部記憶體，這樣會造成目前線程中的CUDA上下文中所包含的的所有流上的資料傳輸和核心計算操作中斷（即使操作是異步的）。

     CUDA裝置上的分支優化和核心調用優化方法同樣可以用在GCN裝置上，是以本節不再做重複的叙述。

      1  雖然GCN裝置上一個wavefront對應的連續256位元組對齊資料具有最高的傳輸效率（每個線程4位元組），但是當遇到計算密集型的問題時，如第一章中所講的那樣使用寬向量加載和存儲操作可以具有更高的效率。

2  GCN裝置上的緩存結構并不具備在wavefront線程間的廣播機制，是以如果多個線程通路同一個或少數幾個資料，更好的方式是通過局部記憶體，例如：

          #if(get_local_id(0)<4){

                        l_data=g_data[get_local_id(0)];

            }barrier(CLK_LOCAL_MEM_FENCE);

          而不是

          data=g_data[get_local_id(0)&3]

3  将不同block内的全局資料通路盡量分散到不同的全局記憶體channel和bank中，如果多個同時進行全局記憶體資料通路的不同block通路的資料位于同一個channel或bank中，則記憶體操作會串行執行，對效率的影響很大，必要時顯式的對block進行排程。

4   GCN裝置上的共享記憶體可以不經過寄存器直接通路（有點類似fermi之前的CUDA裝置），是以可以省去volatile關鍵字。

5.2.2 指令優化

1   由于GCN裝置具有獨立的标量計算單元，是以支援整數計算和浮點計算指令的雙發，合理排程指令的順序可以更好的隐藏指令的發射和計算延遲，比如将浮點計算指令和預取資料的位址計算指令交叉排列。

2  當指令中包含了一些特定的常量值時，編譯器可以生成更小的代碼，因為這些特殊的常量對應了指令的二進制編碼中特定的幾個比特位。這些值是

                              0，1~64,， -1~-16，+-0.5， +-1.0，+-2.0， +-4.0，1.0/(2*PI)

對1.0/(2*PI)内嵌常量的支援更多的是考慮到諸如FFT等圖像計算方面的應用，但是隻有矢量指令才支援在指令碼中内嵌1.0/(2*PI)常量。同時自定義的PI值可能無法比對指令支援的值，是以最好通過使用OpenCL中的内置的定義。

3  和比fermi更早期的CUDA裝置類似，目前所有GCN（1.0~1.3）裝置上對24位整數乘法提供原生支援，是以使用24位整數乘法具有更高的效率。

4  對于GCN1.1，GCN1.2的裝置,在諸如歸約和掃描的應用中盡量使用OpenCL内置的歸約和掃面函數，這樣可以幫助編譯器生成DPP指令進而可以使用硬體上的資料并行引擎執行跨通道計算（無需通過LDS中轉）。

OpenCL是為跨平台的高性能并發程式設計而制定的開放式規範，雖然理論上使用OpenCL開發的程式可以在任何支援OpenCL的平台上運作，但是實際上受限于不同平台對OpenCL支援的力度以及不同硬體架構上得差異，使得同一個OpenCL程式在兩個不同的裝置上的性能表現可能差别很大(甚至這兩個裝置在理論上的技術名額很接近)。很多實際的應用，僅僅擁有代碼的可移植性是不夠的，是以本章主要讨論如何利用OpenCL的運作時編譯系統建構性能可移植的程式。為了寫出性能可移植的程式，不僅僅需要對同一廠商的不同架構的裝置做針對性的優化，同時還要針對不同廠商的裝置給出不同的優化實作；這一過程雖然增加了開發的時間和難度，但是從給與使用者更佳體驗的角度來說是完全值得的。下面我們以并行規約為例講解如何開發性能可移植的OpenCL程式。

4.3.1 CUDA裝置上的并行規約

#ifdef CUDA_DEVICE

#if CUDA_SM<30

#define SMEM_SIZE264

inline voidwarp_reduce_add( double& s, __local volatile double* sptr, int lane )

            if(lane<16)

                        *sptr=s; s+=*(sptr+16);

                        *sptr=s; s+=*(sptr+ 8);

                        *sptr=s; s+=*(sptr+ 4);

                        *sptr=s; s+=*(sptr+ 2);

                        *sptr=s; s+=*(sptr+ 1);

#define SMEM_SIZE 8

inline double__shfl( double val, int mask )

            double out;

            asm volatile ("{                                                                        \n\t"

                        ".reg.b32 slo, shi,dlo, dhi        ;                       \n\t"

                        "mov.b64 { slo, shi}, %1                     ;                       \n\t"

                        "shfl.down.b32 dlo,slo, %2, 0x1f        ;           \n\t"

                        "shfl.down.b32 dhi,shi, %2, 0x1f        ;           \n\t"

                        "mov.b64 %0, { dlo,dhi }                     ;                       \n\t"

                        "}" :"=d"(out) : "d"(val), "r"(mask) );

            return out;

inline voidwarp_reduce_add( double& s )

            s+=__shfl(s,16);

            s+=__shfl(s, 8);

            s+=__shfl(s, 4);

            s+=__shfl(s, 2);

            s+=__shfl(s, 1);

#else defined(GCN_DEVICE)

#define SMEM_SIZE 260

inline voidblock_reduce_add( double& s, __local double* smem, int lane, int warpid )

            __local volatile double*sptr=&smem[get_local_id(0)];

            warp_reduce_add( s, sptr, lane );

            if( lane==0 ){

                        smem[256+warpid]=s;

            } barrier(CLK_LOCAL_MEM_FENCE);

            sptr+=256;

            if(get_local_id(0)<4){

                        s=*sptr; s+=*(sptr+4);

                        *sptr=s; s+=*(sptr+2);

                        *sptr=s; s+=*(sptr+1);

            warp_reduce_add( s );

            if(lane==0){ smem[warpid]=s; }

            barrier(CLK_LOCAL_MEM_FENCE);

            if(get_local_id(0)<8)

                        s=smem[get_local_id(0)];

                        s+=__shfl(s,4);

                        s+=__shfl(s,2);

                        s+=__shfl(s,1);

__kernel voidkReduceAdd( __global double              *g_mapped,

__global double                       *g_temp,

__globalunsigned int  * g_mutex,

__globalconst double * g_a, int n )

            __local unsigned int l_mutex;

            __local double l_temp[SMEM_SIZE];

            double c=0;

            inti=(get_group_id(0)<<8)+get_local_id(0);   

            unsigned intstride=get_num_groups(0)<<8;

            while(i<n){ c+=g_a; i+=stride;}

            const intlane=get_local_id(0)&31;

            const intwarpid=get_local_id(0)>>5;

            block_reduce_add( c, l_temp, lane,warpid );

            if(get_local_id(0)==0){

                        __global double*g_out=(get_num_groups(0)>1)?&g_temp[get_group_id(0)]:g_mapped;

                        *g_out=c;

            if(get_num_groups(0)>1)

               barrier(CLK_GLOBAL_MEM_FENCE);

               if(get_local_id(0)==0){

                            l_mutex=atom_add( &g_mutex, 1 );

               } barrier(CLK_LOCAL_MEM_FENCE);

               if(l_mutex==(get_num_groups(0)-1))

               {

                            c=(get_local_id(0)<get_num_groups(0))?g_temp[get_local_id(0)]:0;

                            block_reduce_add( c, l_temp, lane, warpid);

                            if(get_local_id(0)==0){                                 

                                        g_mapped[0]=c; g_mutex=0;

                                    }

                        }

首先看warp_reduce_add函數，我們使用指針操作，并把通路共享記憶體的次數降到了最小，如果在主函數内

改成 smem[threadIdx.x]=c然後将warp_reduce_add和block_reduce_add改成如下形式

          inline void warp_reduce_add( __local volatiledouble* sptr )

                        if(lane<16)

                        {

                                    sptr[0]+=sptr[16];

                                    sptr[0]+=sptr[8];

                                    sptr[0]+=sptr[4];

                                    sptr[0]+=sptr[2];

                                    sptr[0]+=sptr[1];

            inline double block_reduce_add(double* smem, int lane, int warpid )

            volatile double*sptr=&smem[threadIdx.x];

                        warp_reduce_add( s,sptr, lane );

                        if( lane==0 ){

                                    smem[256+warpid]=s;

                        }barrier(CLK_LOCAL_MEM_FENCE);

                        sptr+=256;

                        if(get_local_id(0)<4){

                        return smem[0];

那麼會多出9次共享記憶體的通路操作。同時注意到在不支援warpshuffle操作的裝置上每個block我們多配置設定了64位元組（SMEM_SIZE=256+8）的共享記憶體，這是為了減少快内同步而做的優化，如果不加上這額外的共享記憶體，那麼我們必須像下面這樣多加一次同步：

warp_reduce_add(s, sptr, lane );

barrier(CLK_LOCAL_MEM_FENCE);

                        smem[warpid]=s;

如果warp_reduce_add之後不加同步的話，那麼就無法保證來自其他warp的第一個線程對共享記憶體寫入之前第一個warp的計算已經完成，就會有寫沖突的問題。這一優化政策對于GCN裝置上的實作同樣适用，這其實也是對雙緩沖技術的一個變相應用。

4.3.2 GCN裝置上的并行規約

     GCN裝置上的實作和CUDA裝置上的實作很相似，稍微的不同在CUDA裝置上鎖步計算的粒度是32，而GCN裝置上的鎖步粒度則是64(和CUDA裝置上基于顯式的warp并發一樣，單個wavefront内的線程是以鎖步的方式并行執行的，是以無需同步)，是以為了在GCN架構上具有更好的親和性，我們需要将鎖步并發粒度改成64。

inline voidwavefront_reduce_add( double& s, __local double* sptr, int lane )

            if(lane<32)

*sptr=s; s+=*(sptr+32);

inline voidblock_reduce_add( double& s, __local double* smem, int lane, intwavefront_id )

            wavefront_reduce_add( s, sptr, lane);

                        smem[256+wavefront_id]=s;

            if(get_local_id(0)<2){

__kernel voidkReduceAdd(    __global double                       * g_mapped,

__global unsigned int  * g_mutex,

__global const double * g_a, int n )

            int i=(get_group_id(0)<<8)+get_local_id(0);   

            const intlane=get_local_id(0)&63;

            const intwavefront_id=get_local_id(0)>>6;

            block_reduce_add( c, l_temp, lane,wavefront_id );

__globaldouble* g_out=(get_num_groups(0)>1)?&g_temp[get_group_id(0)]:g_mapped;

                           c=(get_local_id(0)<get_num_groups(0))?g_temp[get_local_id(0)]:0;

                            block_reduce_add( c, l_temp, lane,wavefront_id );

現在，我們需要将所有的版本整合進同一個OpenCL核心程式中，這個可以簡單的通過預處理指令來實作

            #define CUDA_GPU   0

            #define GCN_GPU      1

#pragmaOPENCL EXTENSION cl_khr_fp64:enable

#if DEVICE == CUDA_DEVICE

包含CUDA裝置版本的代碼

#elif DEVICE == GCN_GPU

包含GCN裝置版本的代碼

#elif DEVICE == …

包含針對其它裝置優化的版本

…

          ‘DEVICE’以及’CUDA_SM’并沒有在裝置代碼端定義，是以我們需要将它們作為指令行參數傳遞給OpenCL運作時編譯系統，這樣就可以讓OpenCL驅動程式在運作時根據目前的裝置選擇合适的版本進行編譯，進而在多個不同裝置上都可以獲得很高的性能

                        sprintf ( options, “-DDEVICE=dev_type –DCUDA_SM=cc”, …);

                        clBuildProgram(prog, options, … );

小結

          GPU計算優化技術和方法種類繁多，每個人也可能習慣于自己的方法。但總體來說不會有太多不同，即使對于不同架構的裝置很多優化技術也依然是通用的。開發性能可移植的程式從某種意義上來說不僅僅是一種挑戰，也是對程式員技能的一種考驗和磨練；從使用者的角度考慮，他們是性能可移植性的程式的最終受益者，是以這樣做也就更具有現實意義。有些優化技術在一些情況下可能适用，而另一些情況則可能适得其反，是以實際的驗證必不可少。

參考資料

  1《CUDAProgramming Guide》

  2《AMD GCN Architecture Whitepaper》

  3《AMDSouthern Island Series Instruction Set Architecture》

  4《Graphics Core Next Architecture，Generation3》

  5《AMD Parallel Processing OpenCL Programming Guide》

原文釋出時間為：2016-6-24 11:43:44

原文由：nvadmin釋出，版權歸屬于原作者 

本文來自雲栖社群合作夥伴NVIDIA，了解相關資訊可以關注NVIDIA官方網站