關於cuda實現歸約求和算法的問題


      歸約算法的基本思想是,對一個輸入數組執行某種計算,然后產生一個更小的結果數組。當大量的數進行加和計算時,可以利用歸約算法,多線程同時進行求和計算,使得時間復雜度下降。

算法思想如下:

      設數據總數為N,輸入數組為a[N]。首先啟用N/2個線程,對於第i個線程,計算 a[i]=a[i]+a[i+N/2],此時,每個線程都將兩個值合並為一個值,得到部分和數組,即a[0]到a[N/2-1];然后,對部分和數組執行操作,第i個線程計算a[i]=a[i]+a[i+N/4],部分和數組長度又減半;接着,以此類推,在此操作執行log2 N次后,a[0]即為數組總和。

       這可通過下面的並行算法實現(記為代碼段00):

輸入:a[N],N=2k                
輸出:數組值的總和存儲在a[0]中。
BEGIN
  1. p=N/2
  2. While p>0 do
  3. For i=1 to p do in parallel
  4. a[i]=a[i]+a[i+p]
  5. End parallel
  6. p=[p/2]
  7. End while
END

        while循環重復log2 N次,時間復雜度為O(log2 N),與串行時時間復雜度O(N)相比,減小了計算時間。 而此時,Watch out!!

1.不要讓“拖后腿的短板”影響程序的正確性

      由於在設備端上各線程的計算速度不同,在執行中,速度快的線程可能會讀取速度慢的線程還未寫好的那塊內存空間。這很可能是很危險的操作。此時,我們可以選擇用共享內存,利用同一線程塊的線程的同步來抑制這種情況。

      相應的代碼段(代碼段01)如下:

__shared__ int cache[N];
int tid = threadIdx.x;
cache[tid]=a[tid];//a[]為函數的參數,傳入輸入數組 __syncthreads();//線程同步
int i=N/2;
while(i!=0)
{
 if (tid<i){
    cache[tid]+=cache[tid+i];
}
__syncthreads();
i=i/2;
}

      Ta-da!記得用共享內存與其同步機制哦,如題所言,不要讓拖后腿的慢線程影響你程序的正確性~    

2.__syncthreads()只適用於同一線程塊里;not 線程塊間

      當數組長度N超過一個線程塊內的最大線程數(threadsPerBlock)時,要同時啟用多個線程塊。相應地將輸入數組分成多個組,每組為threadsPerBlock個,每個線程塊做上述代碼段中的操作,求得部分和(即每組數據的總和)。其中,每個線程塊只將threadsPerBlock個(而非N個)數據拷入到共享內存數據,所以數組索引改為以下形式(代碼段02):

if(tid+blockIdx.x*blockDim.x<N)
{
   cache[tid]=a[tid+blockIdx.x*blockDim.x];
}

     求得了部分和(即每組數據的總和),存放在每個線程塊的cache[0]里,通過以下語句(代碼段03)將每個線程塊內計算出的和放入一個數組中:

if(tid==0)
{
partial_a[blockIdx.x]=cache[0];
}

     接着,對此部分和數組也執行代碼段01中的操作,進行同樣的歸約求和,來求得數據總和。

     然而,卻不可直接在代碼段03下進行數組partial_a[]求和操作,因為每個線程塊的執行速度又不同,直接在后面進行歸約求和操作,容易發生競態條件。因為線程塊間沒有同步機制(這一點值得注意)。此時,可以再調用一個核函數totalAdd()來求總和。

partial_add<<<BlockPerGrid,threadsPerBlock>>>(partial_a); //歸約求和里數組和又存放在本數組中,因此兩核函數參數可以相同。都為a或partial_a。 
total_add<<<1,BlockPerGrid>>>(partial_a); //取partial_a名字是為形象起見,可把本頁partial_a都寫為a

3.小心訪問到不該訪問的內存區域

     以上的討論限定N=2k(在代碼段00中),若N不是2k,或者N/threadsPerBlock不是整數,此時需要多開一線程塊,而這線程塊的線程沒有全部用上。cache[threadsPerBlock]后部分有可能沒有被賦值為我們的輸入數據。則代碼段01中兩數相加cache[tid]+=cache[tid+i]會使程序變得不正確。給對相應的代碼加上限制條件,避免相加這些危險數據。代碼如下:

while(i!=0)
{
     if(tid<i)
     {
           if(tid+i+blockIdx.x*blockDim.x<N)
           {
                 cache[tid]+=cache[tid+i];
           }
           else
          {  
                 cache[tid];
          }
      }
      __syncthreads();
      i=i/2;
}

     當然,如果你足夠叛逆,不喜歡用前一半數加上后一半數的算法,亦可采用其他算法。比如,相鄰兩數相加,其和存在前一半數中。更改代碼段00:

//a[i]=a[i]+a[i+p]
 a[i]=a[2*i-1]+a[2*i]

      不過,其他地方也需要作出相應修改。

     Ta-da!至此,就可以完成任意大量數據的歸約求和算法咯。


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM