[問題] cuda tree reduction

[godman362 (青)]

小弟最近剛開始學習cuda

對於習慣寫單一thread程式的我有點轉不過來，還請各位指點一下

我目前做一個練習，宣告一個512*512的陣列

目的要將這512*512用tree reduction加總成一個數

我在host端呼叫第一個kernel function：


add<<<512, 512>>>(da, db, SIZE)

da：512*512長度陣列位址

db：512長度陣列位址

dc：512*512(size)

然後我第一個kernel function是這樣寫的：



__global__ void add(int *a, int *b, int N)
{
  int t = threadIdx.x;
  int bk = blockIdx.x;

  __shared__ int w;
  w = 0;

  w += a[t];
  __syncthreads();

  b[bk] = w;
}


第一個kernel我目的是要將一個block裡面的thread讀到的值加總

也就是說，我第一個kernel會把512*512個數，加總到剩下512個數

接著在呼叫第二個kernel做tree reduction

不過我在這個kernel呼叫就產生問題了

我一開始把da陣列裡面的元素全部設為「10」

也就是說，我的每個w加總出來的結果應該要是「5120」

但是我的w卻只有10

還請各位幫忙指正一下錯誤，謝謝

--
                               【一路說到掛】                     ︻           
              空谷殘聲  簫中劍  蕭無人  簫中劍  空谷殘聲  簫中劍  現           
                                                                  正           
                               瘋狂兌現俠道精神                   熱           
                                                                  映           
                          黃文擇拒絕再配音的武痴傳人              ︼           

--
※ 發信站: 批踢踢實業坊(ptt.cc)
◆ From: 122.118.9.133
推 LPH66:+= 不是 atomic  平行做的話會有同步問題
→ LPH66:你該做的是開 512 個 thread,每個 thread 把 512 個數加起來
→ godman362:但是thread和thread之間不能直接溝通啊？
→ godman362:所以我應該要用__shared__ int w[512]這樣子嗎？
→ godman362:然後把元素copy到w[512]之後做sum
→ godman362:不知道L大是不是這個意思？
推 LPH66:他們因為 __shared__ 的 w 而間接溝通了
→ janyfor:按照你的想法每個執行緒要跑"512次的累加迴圈"
→ janyfor:        寫                                 才有累加
→ janyfor:512 個 threads "平行化累加" 得到1個數值
→ janyfor:因有 512 blocks, 所以最後會得到 512 個數值
感謝兩位的指教，以下是我目前改出來的code

__global__ void add(int *a, int *b, int N)
{
  int t = threadIdx.x;
  int bk = blockIdx.x;

  __shared__ int w[512];
  w[t] = a[t];

  for (int i = 256; i > 0; i >>= 1) {
    if (t < i) {
      w[t] += w[t+i];
    }
    __syncthreads();
  }

  b[bk] = w[0];
}

我將global memory讀到shared memory，然後在每個block做tree reduction

再將每個w[0]丟給b陣列對應的位置

再呼叫第二個kernel做tree reduction就可以了

謝謝兩位指點
※ 編輯: godman362       來自: 122.118.3.222        (05/08 07:45)
推 andyjy12:去看一下SDK的 reduction怎麼做比較有效率吧~
推 littleshan:你整個搞錯 tree reduction 的意思了
→ godman362:我搞錯tree reduction的意思...請問我哪邊搞錯了？
→ godman362:另外謝謝a大指點，我會去參考一下SDK的部份
→ godman362:剛剛看了SDK裡面的ppt，有提到wrap這個部份
→ godman362:不知道是不是a大說的可以增加效率的部份？