深入淺出GPU優(yōu)化系列：reduce優(yōu)化
本篇文章主要是介紹如何對GPU中的reduce算法進行優(yōu)化。目前針對reduce的優(yōu)化，Nvidia的官方文檔reduce優(yōu)化已經(jīng)說得比較詳細，但是過于精簡，很多東西一筆而過。對于初入該領(lǐng)域的新人而言，理解起來還是較為費勁。因而在官方文檔的基礎(chǔ)，進行更深入地說明和講解，盡可能地讓每一個讀者通過此文都能徹底地了解reduce的優(yōu)化技術(shù)。
前言
首先需要對reduce算法進行介紹。reduce算法本質(zhì)上就是計算
x=x0?x1?x2?x3......?xn?1?xnx=x0?x?1?x?2?x3......?x?n?1?xn?
下面本文將詳細說明如何在GPU中實現(xiàn)reduce算法并進行深入地優(yōu)化。
并行算法設(shè)計
在GPU中，reduce采用了一種樹形的計算方式。如下圖所示。
從上至下，將數(shù)據(jù)不斷地累加，直到得出最后的結(jié)果，即25。但由于GPU沒有針對global數(shù)據(jù)的同步操作，只能針對block的數(shù)據(jù)進行同步。所以，一般而言將reduce分為兩個階段，其示意圖如下：
我們仔細來看看這個事，假設(shè)給定一個長度為N的數(shù)組，需要計算該數(shù)組的所有元素之和。首先需要將數(shù)組分為m個小份。而后，在第一階段中，開啟m個block計算出m個小份的reduce值。最后，在第二階段中，使用一個block將m個小份再次進行reduce，得到最終的結(jié)果。由于第二階段本質(zhì)上是可以調(diào)用第一個階段的內(nèi)核，所以不做單獨說明，本文只是探索第一階段的優(yōu)化技巧。
所以kernel接口為：
__global__ void reduce(T *input, T* output)
其中，input代表輸入的數(shù)組，即一個長度為N的數(shù)組，output代表輸出數(shù)組，即第一階段的結(jié)果，即長度為M的數(shù)組。隨后要開始激動人心的編碼階段，但在CUDA編程中，我們首先需要設(shè)置三個參數(shù)：
BlockNum：即開啟的block數(shù)量，即上面所說的M，代表需要將數(shù)組切分為幾份。
Thread_per_block：每個block中開啟的線程數(shù)，一般而言，取128，256，512，1024這幾個參數(shù)會比較多。
Num_per_block：每個block需要進行reduce操作的長度。
其中，BlockNum* Num_per_block=N
三個參數(shù)的示意圖如下：
Reduce baseline算法介紹
Baseline算法比較簡單，分為三個步驟。第一個步驟是將數(shù)據(jù)加載至共享內(nèi)存中，第二個步驟是在共享內(nèi)存中對數(shù)據(jù)進行reduce操作，第三個步驟是將最后的結(jié)果寫回全局內(nèi)存中。代碼如下：
__global__ void reduce0(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*blockDim.x+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?for(unsigned int s=1; s<blockDim.x; s*=2){
? ? ? ?if(tid%(2*s) == 0){
? ? ? ? ? ?sdata[tid]+=sdata[tid+s];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
在進行優(yōu)化之前，我們需要再來好好地梳理一下這個baseline代碼。優(yōu)化的本質(zhì)是通過軟件榨干硬件資源，所以必須清楚地了解代碼在硬件上的執(zhí)行過程才能更好地進行優(yōu)化。因此，本節(jié)將花較多的篇幅說明代碼和硬件的對應(yīng)關(guān)系，為后續(xù)的優(yōu)化打好基礎(chǔ)。
在第一個步驟中，我們讓Num_per_block與Thread_per_block一致，每個block設(shè)定為256個線程，一個block負責(zé)256個數(shù)據(jù)的reduce工作。假設(shè)需要處理32M的數(shù)據(jù)，則有128K個塊。tid代表線程號，i代表在原始數(shù)組中的索引號。第tid號線程將第i號的數(shù)據(jù)從global中取出，放到sharedmemory的第tid元素中。比如在第0號block中，0號線程將0號元素取出，放到共享內(nèi)存的第0號位置。示意圖見：
從硬件角度來分析一下代碼。為了執(zhí)行代碼，GPU需要分配兩種資源，一個是存儲資源，一個是計算資源。存儲資源包括在global memory中分配的一塊32M
×?sizeof（float）的空間以及在共享內(nèi)存中分配的256×?sizeof（float）的空間。需要注意的是，共享內(nèi)存存在bank沖突的問題，因而需要格外小心。計算資源其實是根據(jù)thread數(shù)量來確定的，一個block中分配256個thread線程，32個線程為一組，綁定在一個SIMD單元。所以256個線程可以簡單地理解為分配了8組SIMD單元。（但實際的硬件資源分配不是這樣，因為一個SM的計算資源有限，不可能真的給每一個block都分配這么多的SIMD單元。）總而言之，在第一個階段，就是tid號線程將i號數(shù)據(jù)從global memory中取出，再放進shared memory中，嚴謹一點的話，中間是走一遍寄存器再到共享內(nèi)存中的。
到了第二個階段，block中需要計算的256個元素已經(jīng)全部被存儲在了共享內(nèi)存中，此時需要對其進行reduce操作。這個過程需要進行多輪迭代，在第一輪迭代中，如果tid%2 ==0， 則第tid號線程將共享內(nèi)存中第tid號位置的值和第tid+1號的值進行相加，而后放在第tid號位置。在第二輪迭代中，如果tid%4==0，則第tid號線程將共享內(nèi)存中第tid號位置的值和第tid+2號的值進行相加，而后放在第tid號位置。不斷迭代，則所有元素都將被累加到第0號位置。其示意圖如下。其中，紅色的線程代表符合if條件的線程，只有它們有任務(wù)，需要干活。
在第三個階段中，block負責(zé)的256個元素之和都放置在共享內(nèi)存的0號位置，此時，只需要將0號位置的元素寫回即可。
實驗結(jié)果
優(yōu)化技巧1：解決翹曲發(fā)散
現(xiàn)有問題
目前reduce0存在的最大問題就是Warp Divergent的問題。對于一個block而言，它所有的thread都是執(zhí)行同一條指令。如果存在if-else這樣的分支情況的話，thread會執(zhí)行所有的分支。只是不滿足條件的分支，所產(chǎn)生的結(jié)果不會記錄下來。可以在上圖中看到，在每一輪迭代中都會產(chǎn)生兩個分支，分別是紅色和橙色的分支。這嚴重影響了代碼執(zhí)行的效率。
解決方式
解決的方式也比較明了，就是盡可能地讓所有線程走到同一個分支里面。
代碼示意如下：
__global__ void reduce1(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*blockDim.x+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?for(unsigned int s=1; s<blockDim.x; s*=2){
? ? ? ?int index = 2*s*tid;
? ? ? ?if(index < blockDim.x){
? ? ? ? ? ?sdata[index]+=sdata[index+s];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
雖然代碼依舊存在著if語句，但是卻與reduce0代碼有所不同。我們繼續(xù)假定block中存在256個thread，即擁有256/32=8個warp。當(dāng)進行第1次迭代時，0-3號warp的index<blockDim.x， 4-7號warp的index>=blockDim.x。對于每個warp而言，都只是進入到一個分支內(nèi)，所以并不會存在warp divergence的情況。當(dāng)進行第2次迭代時，0、1號兩個warp進入計算分支。當(dāng)進行第3次迭代時，只有0號warp進入計算分支。當(dāng)進行第4次迭代時，只有0號warp的前16個線程進入分支。此時開始產(chǎn)生扭曲發(fā)散。通過這種方式，我們消除了前3次迭代的warp divergence。
實驗結(jié)果
優(yōu)化技巧2：解決bank沖突
現(xiàn)有問題
reduce1的最大問題是銀行沖突。我們把目光聚焦在這個for循環(huán)中。并且只聚焦在0號warp。在第一次迭代中，0號線程需要去load shared memory的0號地址以及1號地址的數(shù)，然后寫回到0號地址。而此時，這個warp中的16號線程，需要去load shared memory中的32號地址和33號地址?？梢园l(fā)現(xiàn)，0號地址跟32號地址產(chǎn)生了2路的bank沖突。在第2次迭代中，0號線程需要去load shared memory中的0號地址和2號地址。這個warp中的8號線程需要load shared memory中的32號地址以及34號地址，16號線程需要load shared memory中的64號地址和68號地址，24號線程需要load shared memory中的96號地址和100號地址。又因為0、32、64、96號地址對應(yīng)著同一個bank，所以此時產(chǎn)生了4路的bank沖突?，F(xiàn)在，可以繼續(xù)算下去，8路bank沖突，16路bank沖突。由于bank沖突，所以reduce1性能受限。下圖說明了在load第一個數(shù)據(jù)時所產(chǎn)生的bank沖突。
解決方式
在reduce中，解決bank沖突的方式就是把for循環(huán)逆著來。原來stride從0到256，現(xiàn)在stride從128到0。其偽代碼如下：
__global__ void reduce2(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*blockDim.x+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?for(unsigned int s=blockDim.x/2; s>0; s>>=1){
? ? ? ?if(tid < s){
? ? ? ? ? ?sdata[tid]+=sdata[tid+s];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
那為什么通過這么一個小小的改變就能消除銀行沖突呢，我們繼續(xù)進行分析。
把目光繼續(xù)看到這個for循環(huán)中，并且只分析0號warp。0號線程需要load shared memory的0號元素以及128號元素。1號線程需要load shared memory中的1號元素和129號元素。這一輪迭代中，在讀取第一個數(shù)時，warp中的32個線程剛好負載一行共享內(nèi)存數(shù)據(jù)。再分析第2輪迭代，0號線程load 0號元素和64號元素，1號線程load 1號元素和65號元素。咦，也是這樣，每次加載共享內(nèi)存的一行。再來分析第3輪迭代，0號線程load 0號元素和32號元素，接下來不寫了，總之，一個warp load shared memory的一行。沒有銀行沖突。到了4輪迭代，0號線程load 0號元素和16號元素。那16號線程呢，16號線程啥也不干，因為s=16，16-31號線程啥也不干，跳過去了。示意圖如下：
實驗結(jié)果
優(yōu)化技巧3：解決idle線程
現(xiàn)有問題
reduce2最大的問題就是線程的浪費?？梢钥吹轿覀儐恿?56個線程，但是在第1輪迭代時只有128個線程在干活，第2輪迭代只有64個線程在干活，每次干活的線程都會減少一半。第一輪迭代示意圖如下，只有前128個線程在load數(shù)據(jù)。后128個線程啥也不干，光看著。
解決方式
對于HPC從業(yè)者而言，我們希望變成GPU的資本家，去盡可能地壓榨GPU。但是呢，在這里，每一次迭代有一半的線程不干活。而且，128-255號線程最過分，它娘的，沒有任何貢獻，啥也不干。想來想去，能不能讓它們干點活呢。想來想去，那這樣吧，讓它好歹做一次加法。除了去全局內(nèi)存中取數(shù)外，再做一次加法。當(dāng)然為了實現(xiàn)這個，block數(shù)就得改一改了。Block數(shù)量減少，Num_per_block增加一倍。也就是說原來一個block只需要管256個數(shù)就行，現(xiàn)在得管512個數(shù)了。代碼如下：
__global__ void reduce3(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i] + d_in[i+blockDim.x];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?for(unsigned int s=blockDim.x/2; s>0; s>>=1){
? ? ? ?if(tid < s){
? ? ? ? ? ?sdata[tid]+=sdata[tid+s];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
通過這種方式，將一些空閑的線程給利用起來了。
實驗結(jié)果
這里面多說一句，讓idle線程利用起來的這個加速比其實有點出乎意料。
優(yōu)化技巧4：展開最后一維減少同步
現(xiàn)有問題
對于reduce3來說，性能已經(jīng)算是比較好了。但是依舊沒有達到我們想要的效果。我們再來仔細地看看還有什么可以改進的地方。我們發(fā)現(xiàn)，當(dāng)進行到最后幾輪迭代時，此時的block中只有warp0在干活時，線程還在進行同步操作。這一條語句造成了極大的浪費。
解決方式
由于一個warp中的32個線程其實是在一個SIMD單元上，這32個線程每次都是執(zhí)行同一條指令，這天然地保持了同步狀態(tài)，因而當(dāng)s=32時，即只有一個SIMD單元在工作時，完全可以將__syncthreads（）這條同步代碼去掉。所以我們將最后一維進行展開以減少同步。偽代碼如下：
__device__ void warpReduce(volatile float* cache,int tid){
? ?cache[tid]+=cache[tid+32];
? ?cache[tid]+=cache[tid+16];
? ?cache[tid]+=cache[tid+8];
? ?cache[tid]+=cache[tid+4];
? ?cache[tid]+=cache[tid+2];
? ?cache[tid]+=cache[tid+1];
}
__global__ void reduce4(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i] + d_in[i+blockDim.x];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?for(unsigned int s=blockDim.x/2; s>32; s>>=1){
? ? ? ?if(tid < s){
? ? ? ? ? ?sdata[tid]+=sdata[tid+s];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid<32)warpReduce(sdata,tid);
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
可以通過下面的示意圖更好地了解，（圖畫的有點丑，比例也不太對，大家將就著看看。）warp0會被綁定在一個SIMD單元上，上面有thread0-thread31。warp1會被綁在另外一個SIMD單元上，上面有thread32-thread63。由于在一個 SIMD單元上，然后不管啥時候thread0和thread7肯定是同一狀態(tài)，不需要同步。而thread0和thread34就不能保證同步，必須用__syncthreads（）來保證同步操作。
實驗結(jié)果
在做到這一步時，帶寬已經(jīng)到了756GB/s。這個時候就已經(jīng)優(yōu)化地差不多了。性能也很難做到有效提升了。
優(yōu)化技巧5：完全展開減少計算
現(xiàn)有問題
其實到了這一步，reduce的效率已經(jīng)足夠高了。再進一步優(yōu)化其實已經(jīng)非常困難了。為了探索極致的性能表現(xiàn)，Mharris接下來給出的辦法是對for循環(huán)進行完全展開。我覺得這里主要是減少for循環(huán)的開銷。Mharris的實驗表明這種方式有著1.41x的加速比。但是用的機器是G80，十幾年前的卡。性能數(shù)據(jù)也比較老了，至于能不能真的有這么好的加速比，我們拭目以待。
解決方法
我們將整個for循環(huán)進行展開，非常暴力，代碼如下：
template <unsigned int blockSize>
__device__ void warpReduce(volatile float* cache,int tid){
? ?if(blockSize >= 64)cache[tid]+=cache[tid+32];
? ?if(blockSize >= 32)cache[tid]+=cache[tid+16];
? ?if(blockSize >= 16)cache[tid]+=cache[tid+8];
? ?if(blockSize >= 8)cache[tid]+=cache[tid+4];
? ?if(blockSize >= 4)cache[tid]+=cache[tid+2];
? ?if(blockSize >= 2)cache[tid]+=cache[tid+1];
}
template <unsigned int blockSize>
__global__ void reduce5(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?sdata[tid]=d_in[i] + d_in[i+blockDim.x];
? ?__syncthreads();
? ?// do reduction in shared mem
? ?if(blockSize>=512){
? ? ? ?if(tid<256){
? ? ? ? ? ?sdata[tid]+=sdata[tid+256];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?if(blockSize>=256){
? ? ? ?if(tid<128){
? ? ? ? ? ?sdata[tid]+=sdata[tid+128];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?if(blockSize>=128){
? ? ? ?if(tid<64){
? ? ? ? ? ?sdata[tid]+=sdata[tid+64];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid<32)warpReduce<blockSize>(sdata,tid);
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
實驗結(jié)果
可以看到，還是有所收益，但是并沒有那么地顯著。這主要是因為GPU硬件架構(gòu)的不斷發(fā)展，以及NV在編譯器上面也做了較多的工作。
優(yōu)化技巧6：合理設(shè)置block數(shù)量
現(xiàn)有問題
當(dāng)走到這一步的時候，能調(diào)的東西已經(jīng)基本上調(diào)完了。我們再把眼光放在block和thread的設(shè)置上。之前默認了Num_per_block=Thread_per_block。也就是說，一個block開啟256個線程時，這個block負責(zé)256個元素的reduce操作。那可不可以讓一個block多管點數(shù)。這樣的話，開啟的block數(shù)量少一些。以此對block設(shè)置進行調(diào)整，獲得最優(yōu)block取值，這樣或許能夠帶來一些性能收益？
解決方式
這樣需要再思考一下block的取值。對于GPU而言，block的取值到底是多更好，還是少更好。如此對CUDA編程熟悉的同學(xué)，肯定會毫不猶豫地說：“那肯定是多更好啦。Block數(shù)量多，block可以進行快速地切換，去掩蓋訪存的延時?！斑@個問題按下不表，我們看看Mharris是怎么說的。
如果一個線程被分配更多的工作時，可能會更好地覆蓋延時。這一點比較好理解。如果線程有更多的work時，對于編譯器而言，就可能有更多的機會對相關(guān)指令進行重排，從而去覆蓋訪存時的巨大延時。雖然這句話并沒有很好地說明在某種程度上而言，block少一些會更好。但是，有一點不可否認，block需要進行合理地設(shè)置。嘮嘮叨叨說了很多，現(xiàn)在把代碼貼一下：
template <unsigned int blockSize>
__global__ void reduce6(float *d_in,float *d_out){
? ?__shared__ float sdata[THREAD_PER_BLOCK];
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?unsigned int gridSize = blockSize * 2 * gridDim.x;
? ?sdata[tid] = 0;
? ?while(i<n){
? ? ? ?sdata[tid] +=d_in[i]+d_in[i+blockSize];
? ? ? ?i+=gridSize;
? ?}
? ?__syncthreads();
? ?// do reduction in shared mem
? ?if(blockSize>=512){
? ? ? ?if(tid<256){
? ? ? ? ? ?sdata[tid]+=sdata[tid+256];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?if(blockSize>=256){
? ? ? ?if(tid<128){
? ? ? ? ? ?sdata[tid]+=sdata[tid+128];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?if(blockSize>=128){
? ? ? ?if(tid<64){
? ? ? ? ? ?sdata[tid]+=sdata[tid+64];
? ? ? ?}
? ? ? ?__syncthreads();
? ?}
? ?
? ?// write result for this block to global mem
? ?if(tid<32)warpReduce<blockSize>(sdata,tid);
? ?if(tid==0)d_out[blockIdx.x]=sdata[tid];
}
實驗結(jié)果
對于block的取值，我進行了微調(diào)，大概取2048會比512，1024，4096的效果要好。理論上來說，這個值取SM數(shù)量的倍數(shù)會比較合理。但是V100的SM是80，取一個完美的倍數(shù)還是比較困難。目前達到了768GB/s。
優(yōu)化技巧7：使用shuffle指令
現(xiàn)有問題
其實，對于Mharris的講義。reduce優(yōu)化就到此結(jié)束了。但是NV后來出了Shuffle指令，對于reduce優(yōu)化有著非常好的效果。目前絕大多數(shù)訪存類算子，像是softmax，batch_norm，reduce等，都是用Shuffle實現(xiàn)。所以，在這里談一下這么把shuffle指令用在reduce優(yōu)化上。
Shuffle指令是一組針對warp的指令。Shuffle指令最重要的特性就是warp內(nèi)的寄存器可以相互訪問。在沒有shuffle指令的時候，各個線程在進行通信時只能通過共享內(nèi)存來訪問彼此的寄存器。而采用了shuffle指令之后，warp內(nèi)的線程可以直接對其他線程的寄存器進行訪存。通過這種方式可以減少訪存的延時。除此之外，帶來的最大好處就是可編程性提高了，在某些場景下，就不用共享內(nèi)存了。畢竟，開發(fā)者要自己去控制 shared memory還是挺麻煩的一個事。
偽代碼如下：
template <unsigned int blockSize>
__device__ __forceinline__ float warpReduceSum(float sum){
? ?if(blockSize >= 32)sum += __shfl_down_sync(0xffffffff,sum,16);
? ?if(blockSize >= 16)sum += __shfl_down_sync(0xffffffff,sum,8);
? ?if(blockSize >= 8)sum += __shfl_down_sync(0xffffffff,sum,4);
? ?if(blockSize >= 4)sum += __shfl_down_sync(0xffffffff,sum,2);
? ?if(blockSize >= 2)sum += __shfl_down_sync(0xffffffff,sum,1);
? ?return sum;
}
template <unsigned int blockSize>
__global__ void reduce7(float *d_in,float *d_out, unsigned int n){
? ?float sum = 0;
? ?//each thread loads one element from global memory to shared mem
? ?unsigned int i=blockIdx.x*(blockDim.x*2)+threadIdx.x;
? ?unsigned int tid=threadIdx.x;
? ?unsigned int gridSize = blockSize * 2 * gridDim.x;
? ?while(i<n){
? ? ? ?sdata[tid] +=d_in[i]+d_in[i+blockSize];
? ? ? ?i+=gridSize;
? ?}
? ?// shared mem for partial sums(one per warp in the block
? ?static __shared__ float warpLevelSums[WARP_SIZE];
? ?const int laneId = threadIdx.x % WARP_SIZE;
? ?const int warpId = threadIdx.x / WARP_SIZE;
? ?sum = warpReduceSum<blockSize>(sum);
? ?if(laneId == 0)warpLevelSums[warpId]=sum;
? ?__syncthreads();
? ?sum = (threadIdx.x < blockDim.x / WARP_SIZE)? warpLevelSums[laneId]:0;
? ?// Final reduce using first warp
? ?if(warpId == 0)sum = warpReduceSum<blockSize/WARP_SIZE>(sum);
? ?// write result for this block to global mem
? ?if(tid==0)d_out[blockIdx.x]=sum;
}
實驗結(jié)果
總結(jié)與思考
通過上次一系列的優(yōu)化技巧，我們對reduce進行了不斷地優(yōu)化。最后效果是770.3GB/s。帶寬利用達到85%。PS：具體的reduce性能數(shù)據(jù)還需要大量測試，但由于我比較懶。所以測試工作就到此為此，大家有興趣可以自己再跑跑。代碼在這里，還有點亂，沒怎么整理，大家有疑問可以直接評論或者私信我。
reduce代碼
github.com/Liu-xiandong/How_to_optimize_in_GPU
然后不同的優(yōu)化技巧所帶來的性能表現(xiàn)如下：
可以看到，其實在reduce4的時候就已經(jīng)很難再提升了。而且這個數(shù)據(jù)跟NV在博客上的數(shù)據(jù)還是有比較大的出入，主要是對最后一維展開后，在V100上的帶寬利用率就很難再提升了，而G80還有顯著的提升。我覺得主要原因是硬件已經(jīng)更新了好幾代。
總而言之，我們通過這一系列的優(yōu)化已經(jīng)可以把reduce優(yōu)化到一個非常好的程度。我之前測過一次是160us，十分接近800GB/s。但不知道為啥就不能復(fù)現(xiàn)了。對于訪存型的算子，在V100上能做到接近800GB/s的帶寬就已經(jīng)接近極限了。而目前能夠做的優(yōu)化，也都列了出來。
《GPU優(yōu)化教程系列》是澎峰科技收集、整理、創(chuàng)作的一個公益系列課程
更多訊息可關(guān)注微信公眾號：澎峰科技PerfLXab
在B站上的《先進計算公益課》，歡迎收看。
https://space.bilibili.com/1444176265/channel/collectiondetail?sid=605205
