開發平台(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