作者hardman1110 (笨小孩)
看板C_and_CPP
標題[問題] CUDA shared-memory
時間Tue Oct 3 10:12:05 2017
開發平台(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://webptt.com/m.aspx?n=bbs/C_and_CPP/M.1506996728.A.C64.html
1F:→ johnjohnlin: do something 裡面通常還要有一個 syncthread 10/03 10:30
2F:→ hardman1110: 原因是什? 前面的同步不算嗎? 困惑中= = 10/03 10:40
3F:→ johnjohnlin: 你迴圈繞回去的時候會寫到 shared memory 10/03 10:44
4F:→ hardman1110: Do something 之後就不會更改值了 10/03 11:27
5F:→ hardman1110: 所以我才在一開始同步 10/03 11:28
6F:→ hardman1110: 就算繞回去應該再同步一次不是嗎? 10/03 11:32
7F:推 a1u1usul3: do something的時候有的thread提早做完先去改值了,有 10/03 12:16
8F:→ a1u1usul3: 的thread還沒做完需要用舊的值,但被改了 10/03 12:16
9F:→ hardman1110: 所以我只要在使用前一刻同步就好囉? 10/03 12:37
10F:→ hardman1110: 還有在assign值前同步 10/03 13:00
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:24:40
11F:→ hardman1110: 已嘗試在assign前後都同步,但結果還是會錯(暈 10/03 13:26
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:31:48
12F:→ johnjohnlin: 那 do something 裡面是不是有 break 之類的 10/03 14:00
13F:→ johnjohnlin: BTW, blockDim.x = 135 是個很糟糕的選擇,盡量避免 10/03 14:01
14F:→ a1u1usul3: code還是貼在codepad吧 10/03 14:09
15F:→ a1u1usul3: 然後你的code會不會邏輯上就錯了 10/03 14:19
16F:→ a1u1usul3: 不用sharedMeomry的時候,每個thread從自己的TmpPos拿 10/03 14:19
17F:→ a1u1usul3: 32個元素進private memory,而TmpPos每個thread都不同 10/03 14:20
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 14:21:21
18F:→ a1u1usul3: 結果用SharedMemory的時候32人從自己的TmpPos拿一個 10/03 14:22
19F:→ a1u1usul3: 元素進SharedMemory 10/03 14:23
20F:→ hardman1110: a1大 已補上github好讀版連結 10/03 14:24
21F:→ hardman1110: 我這邊純粹想讓多個thread 同時assign值 甭跑回圈 10/03 14:25
22F:→ a1u1usul3: SharedMemory的功能是讓多個thread共用的資料不用重複 10/03 14:27
23F:→ a1u1usul3: 你資料的並沒有共用不是嗎@@? 10/03 14:27
24F:→ hardman1110: 我想通了~抱歉 確實把y當執行緒切 每個thread y不同 10/03 14:28
25F:→ hardman1110: 純共用的話 感覺用register 就好 陣列大小不大 10/03 14:30
26F:→ a1u1usul3: thread內共用->register threads間共用->sharedMemory 10/03 14:31
27F:→ a1u1usul3: Blocks間共用->GlobalMemory 10/03 14:32
28F:→ a1u1usul3: 好像還有很潮的shuffle,threads間共用的樣子 10/03 14:36
29F:→ hardman1110: 要在加速的話 好像還可以用surface memory來讀寫? 10/03 14:55
30F:→ hardman1110: 感謝各位大大指點 10/03 14:56