[問題] CUDA shared-memory

[hardman1110 (笨小孩)]

開發平台(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
→ johnjohnlin: do something 裡面通常還要有一個 syncthread
→ hardman1110: 原因是什?  前面的同步不算嗎? 困惑中= =
→ johnjohnlin: 你迴圈繞回去的時候會寫到 shared memory
→ hardman1110: Do something 之後就不會更改值了
→ hardman1110: 所以我才在一開始同步
→ hardman1110: 就算繞回去應該再同步一次不是嗎？
推 a1u1usul3: do something的時候有的thread提早做完先去改值了，有
→ a1u1usul3: 的thread還沒做完需要用舊的值，但被改了
→ hardman1110: 所以我只要在使用前一刻同步就好囉？
→ hardman1110: 還有在assign值前同步
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:24:40
→ hardman1110: 已嘗試在assign前後都同步，但結果還是會錯(暈
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 13:31:48
→ johnjohnlin: 那 do something 裡面是不是有 break 之類的
→ johnjohnlin: BTW, blockDim.x = 135 是個很糟糕的選擇，盡量避免
→ a1u1usul3: code還是貼在codepad吧
→ a1u1usul3: 然後你的code會不會邏輯上就錯了
→ a1u1usul3: 不用sharedMeomry的時候，每個thread從自己的TmpPos拿
→ a1u1usul3: 32個元素進private memory，而TmpPos每個thread都不同
※ 編輯: hardman1110 (114.34.230.27), 10/03/2017 14:21:21
→ a1u1usul3: 結果用SharedMemory的時候32人從自己的TmpPos拿一個
→ a1u1usul3: 元素進SharedMemory
→ hardman1110: a1大 已補上github好讀版連結
→ hardman1110: 我這邊純粹想讓多個thread 同時assign值 甭跑回圈
→ a1u1usul3: SharedMemory的功能是讓多個thread共用的資料不用重複
→ a1u1usul3: 你資料的並沒有共用不是嗎@@?
→ hardman1110: 我想通了~抱歉 確實把y當執行緒切 每個thread y不同
→ hardman1110: 純共用的話 感覺用register 就好 陣列大小不大
→ a1u1usul3: thread內共用->register threads間共用->sharedMemory
→ a1u1usul3: Blocks間共用->GlobalMemory
→ a1u1usul3: 好像還有很潮的shuffle，threads間共用的樣子
→ hardman1110: 要在加速的話  好像還可以用surface memory來讀寫?
→ hardman1110: 感謝各位大大指點