PTT數位生活區 / C_and_CPP (C/C++)

[問題] CUDA shared-memory

看板C_and_CPP (C/C++)作者 (笨小孩)時間8年前 (2017/10/03 10:12), 8年前編輯推噓1(1推 0噓 29→)
留言30則, 3人參與, 8年前最新討論串1/1
開發平台(Platform): (Ex: Win10, Linux, ...) WIN10 編譯器(Ex: GCC, clang, VC++...)+目標環境(跟開發平台不同的話需列出) VC2017 額外使用到的函數庫(Library Used): (Ex: OpenGL, ...) CUDA 9.0 問題(Question)： 想透過 shared memory 來加速kernal的效能 利用treadid 平行assign資料 也有用__syncthreads 來同步 但資料還是跟用迴圈跑的不一樣 (結果有錯) 想請問大大們我的使用方式有錯嗎? 還有vc上可以單步執行來看CUDA變數嗎? 餵入的資料(Input)： 一維陣列的輸入與輸出指標 預期的正確結果(Expected Output)： USE_SHARED_MEM = 0 與 = 1 data值要一樣 錯誤結果(Wrong Output)： github: https://github.com/ChiFang/question/blob/master/CUDA_SharedMem.cu USE_SHARED_MEM = 1 會導致最後結果錯誤，表示data值不一樣 (後面程式完全一模一樣) 程式碼(Code)：(請善用置底文網頁, 記得排版) #define USE_SHARED_MEM 1 __global__ void kernal_test(const int a_RangeUpScale, const int *a_CostData, int *a_Input) { // Get the work index of the current element to be processed int y = blockIdx.x*blockDim.x + threadIdx.x; //執行緒在陣列中 對應的位置 #if USE_SHARED_MEM == 1 __shared__ int Buff[32]; #else int Buff[32]; #endif // Do the operation for (int x = 1; (x < g_ImgWidth_CUDA); x++) { int TmpPos = y*Area + (x-1)*a_RangeUpScale; #if USE_SHARED_MEM == 1 // Synchronize to make sure the sub-matrices are loaded before starting the computation __syncthreads(); if (threadIdx.x < 32) { Buff[threadIdx.x] = a_CostSmooth[TmpPos + threadIdx.x]; } // Synchronize to make sure the sub-matrices are loaded before starting the computation __syncthreads(); #else for (int cnt = 0; cnt < 32 ;cnt++) { Buff[cnt] = a_CostSmooth[TmpPos + cnt]; } #endif // use Buff to do something } } 補充說明(Supplement)： grid size = 8 block size = 135 所以thread id 一定會大於32 -- ※ 發信站: 批踢踢實業坊(ptt.cc), 來自: 114.34.230.27 ※ 文章網址: https://www.ptt.cc/bbs/C_and_CPP/M.1506996728.A.C64.html
→
10/03 10:30, 8年前 , 1F
do something 裡面通常還要有一個 syncthread
10/03 10:30, 1F
→
10/03 10:40, 8年前 , 2F
原因是什? 前面的同步不算嗎? 困惑中= =
10/03 10:40, 2F
→
10/03 10:44, 8年前 , 3F
你迴圈繞回去的時候會寫到 shared memory
10/03 10:44, 3F
→
10/03 11:27, 8年前 , 4F
Do something 之後就不會更改值了
10/03 11:27, 4F
→
10/03 11:28, 8年前 , 5F
所以我才在一開始同步
10/03 11:28, 5F
→
10/03 11:32, 8年前 , 6F
就算繞回去應該再同步一次不是嗎？
10/03 11:32, 6F
推
10/03 12:16, 8年前 , 7F
do something的時候有的thread提早做完先去改值了，有
10/03 12:16, 7F
→
10/03 12:16, 8年前 , 8F
的thread還沒做完需要用舊的值，但被改了
10/03 12:16, 8F
→
10/03 12:37, 8年前 , 9F
所以我只要在使用前一刻同步就好囉？
10/03 12:37, 9F
→
10/03 13:00, 8年前 , 10F
還有在assign值前同步
10/03 13:00, 10F
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:24:40
→
10/03 13:26, 8年前 , 11F
已嘗試在assign前後都同步，但結果還是會錯(暈
10/03 13:26, 11F
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:31:48
→
10/03 14:00, 8年前 , 12F
那 do something 裡面是不是有 break 之類的
10/03 14:00, 12F
→
10/03 14:01, 8年前 , 13F
BTW, blockDim.x = 135 是個很糟糕的選擇，盡量避免
10/03 14:01, 13F
→
10/03 14:09, 8年前 , 14F
code還是貼在codepad吧
10/03 14:09, 14F
→
10/03 14:19, 8年前 , 15F
然後你的code會不會邏輯上就錯了
10/03 14:19, 15F
→
10/03 14:19, 8年前 , 16F
不用sharedMeomry的時候，每個thread從自己的TmpPos拿
10/03 14:19, 16F
→
10/03 14:20, 8年前 , 17F
32個元素進private memory，而TmpPos每個thread都不同
10/03 14:20, 17F
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 14:21:21
→
10/03 14:22, 8年前 , 18F
結果用SharedMemory的時候32人從自己的TmpPos拿一個
10/03 14:22, 18F
→
10/03 14:23, 8年前 , 19F
元素進SharedMemory
10/03 14:23, 19F
→
10/03 14:24, 8年前 , 20F
a1大 已補上github好讀版連結
10/03 14:24, 20F
→
10/03 14:25, 8年前 , 21F
我這邊純粹想讓多個thread 同時assign值 甭跑回圈
10/03 14:25, 21F
→
10/03 14:27, 8年前 , 22F
SharedMemory的功能是讓多個thread共用的資料不用重複
10/03 14:27, 22F
→
10/03 14:27, 8年前 , 23F
你資料的並沒有共用不是嗎@@?
10/03 14:27, 23F
→
10/03 14:28, 8年前 , 24F
我想通了~抱歉 確實把y當執行緒切 每個thread y不同
10/03 14:28, 24F
→
10/03 14:30, 8年前 , 25F
純共用的話 感覺用register 就好 陣列大小不大
10/03 14:30, 25F
→
10/03 14:31, 8年前 , 26F
thread內共用->register threads間共用->sharedMemory
10/03 14:31, 26F
→
10/03 14:32, 8年前 , 27F
Blocks間共用->GlobalMemory
10/03 14:32, 27F
→
10/03 14:36, 8年前 , 28F
好像還有很潮的shuffle，threads間共用的樣子
10/03 14:36, 28F
→
10/03 14:55, 8年前 , 29F
要在加速的話 好像還可以用surface memory來讀寫?
10/03 14:55, 29F
→
10/03 14:56, 8年前 , 30F
感謝各位大大指點
10/03 14:56, 30F
‣ 更多 hardman1110 的文章
文章代碼(AID): #1Pql7una (C_and_CPP)
C_and_CPP 近期熱門文章
更多 近期熱門文章 >>
PTT數位生活區 即時熱門文章
更多 即時熱門文章 >>
‣ 更多 hardman1110 的文章
文章代碼(AID): #1Pql7una (C_and_CPP)