[問題] CUDA 新手, threads 之間相加的問題

看板C_and_CPP作者 (@@)時間12年前 (2013/05/07 23:05), 編輯推噓1(1010)
留言11則, 5人參與, 最新討論串1/1
開發平台(Platform): (Ex: VC++, GCC, Linux, ...) Unix 額外使用到的函數庫(Library Used): (Ex: OpenGL, ...) CUDA 問題(Question): global 和 shared 的變數 累加後的答案一樣 餵入的資料(Input): N/A 預期的正確結果(Expected Output): a_sh[threadIdx.x] = 1 *a = 8 錯誤結果(Wrong Output): a_sh[threadIdx.x] = 1 *a = 1 程式碼(Code):(請善用置底文網頁, 記得排版) #define N_Block 2 #define N_Thread 4 __global__ void test(int* a) { // *a = 0; __shared__ int a_sh[N_Thread]; a_sh[threadIdx.x] = 0; // __syncthreads(); *a += 1; a_sh[threadIdx.x] += 1; } int main(int argc,char **argv) { int *d_a; cudaMalloc(&d_a, sizeof(int)); cudaMemcpy(預設 *d_a = 0); test<<<N_Block, N_Thread>>>(d_a); return 0; } 補充說明(Supplement): 我應該是有觀念錯誤 但怎麼樣都想不通 *a 是存在 global 裡面 且是唯一的 每個 blocks 裡面的 threads 都會執行 *a += 1 應該是執行了 4x2 = 8 次吧 怎麼答案還是 1 呢? 感謝 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 24.11.249.170 ※ 編輯: MikePhysics 來自: 24.11.249.170 (05/08 07:08)

05/08 08:38, , 1F
因為第一行就*a = 0;
05/08 08:38, 1F

05/08 08:44, , 2F
但之後每個 threads 都會把 *a 加上了 1
05/08 08:44, 2F

05/08 09:26, , 3F
google "Race condition"
05/08 09:26, 3F

05/08 09:42, , 4F
謝謝您 其實我知道 reduction 的演算法
05/08 09:42, 4F

05/08 09:43, , 5F
但還是不了解 為什麼這個問題的答案會是 1
05/08 09:43, 5F

05/08 09:44, , 6F
只知道 race condition 會讓效率變差 還有讀錯資料
05/08 09:44, 6F

05/08 09:46, , 7F
但在這裡 不是不管 每個 threads 的速度 都會加 1 嗎?
05/08 09:46, 7F

05/08 11:47, , 8F
syncthread是block level的~最後一個block會設成0 再加1
05/08 11:47, 8F

05/08 11:48, , 9F
並不是所有block都執行完 *a=0, 才執行後面的*a+=1
05/08 11:48, 9F
感謝說明 我旦改了一下程式碼 利用 cudaMemcpy 傳初始值給 kernel 答案是一樣的 好像和 一開始的 *a = 0 無關 ※ 編輯: MikePhysics 來自: 24.11.249.170 (05/08 12:01)

05/08 12:37, , 10F
每個thread是同時執行的所以就是0+1而已,你這樣要用atomic
05/08 12:37, 10F

05/08 12:39, , 11F
這種資料有相依性的不適合用GPU作,GPU是作SIMD
05/08 12:39, 11F
文章代碼(AID): #1HYOZKV_ (C_and_CPP)