作者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/cn.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