[問題] cuda tree reduction
小弟最近剛開始學習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
推
05/07 22:37, , 1F
05/07 22:37, 1F
→
05/07 22:38, , 2F
05/07 22:38, 2F
→
05/07 23:10, , 3F
05/07 23:10, 3F
→
05/07 23:10, , 4F
05/07 23:10, 4F
→
05/07 23:10, , 5F
05/07 23:10, 5F
→
05/07 23:11, , 6F
05/07 23:11, 6F
推
05/07 23:28, , 7F
05/07 23:28, 7F
→
05/08 00:23, , 8F
05/08 00:23, 8F
→
05/08 01:02, , 9F
05/08 01:02, 9F
→
05/08 01:11, , 10F
05/08 01:11, 10F
→
05/08 01:13, , 11F
05/08 01:13, 11F
感謝兩位的指教，以下是我目前改出來的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)
推
05/08 11:13, , 12F
05/08 11:13, 12F
推
05/08 11:33, , 13F
05/08 11:33, 13F
→
05/08 16:03, , 14F
05/08 16:03, 14F
→
05/08 16:07, , 15F
05/08 16:07, 15F
→
05/08 16:33, , 16F
05/08 16:33, 16F
→
05/08 16:34, , 17F
05/08 16:34, 17F