开发平台(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