作者v00623 (阿哩他命EX PLUS)
看板C_and_CPP
标题[问题] CUDA Memcpy相关问题
时间Wed Jul 12 17:25:08 2017
开发平台(Platform): (Ex: Win10, Linux, ...)
Linux
GPGPU-Sim
问题(Question)：
目前正在测试dependent kernel
kernel A 负责两个阵列相加存到result array
kernel B 把kernel A的reult + 1
再将结果与CPU运算结果比对
一般的写法是：
1.Memcpy两个阵列的data到device
2.launch kernel A
3.将result复制回host
4.把result传到device
5.launch kernel B
6.将运算结果复制回host
7.比对
但我把第三四步骤省略 比对结果也是正确的
请问是为什麽呢？
是因为kernel B需要的data与kernel A的result array记忆体位置相同？
或是有其他原因？
程式码如下
https://gist.github.com/anonymous/5d8b6c58ce7ecd0407f6595d41fd8a2c
後来我改变他的data_size 例如200(原本是100)
结果前100个正确 後100个错误 这又是为什麽？
--
※ 发信站: 批踢踢实业坊(ptt.cc), 来自: 140.118.155.192
※ 文章网址: https://webptt.com/cn.aspx?n=bbs/C_and_CPP/M.1499851515.A.90A.html
1F：→ Schottky: 第三和第四步骤 (36,37 两行) 本来就是多余的 07/12 17:41
2F：→ wheatdog: 3和4感觉没有必要，可以直接把gpu的pointer传给下个ker 07/12 17:44
3F：→ wheatdog: nel 07/12 17:44
4F：→ Schottky: 错误的原因... 你问我我问谁... 07/12 17:45
5F：→ Schottky: VectorAdd2 到底实际作用是什麽只有你知道 07/12 17:45
6F：→ Schottky: 但我觉得 17 行那个 idx 值可能跟你想的不一样 07/12 17:46
7F：→ Schottky: 试试 int idx = blockIdx.x*blockDim.x + threadIdx.x; 07/12 17:46
为什麽3,4是多余的？ 是遇到要使用相同的data就可以不用这两个步骤的意思？
VectorAdd2作用是把VectorAdd1的结果再加一
整个程式目的是要测试不做3,4会不会影响结果
改成blockIdx.x*blockDim.x + threadIdx.x确实是正确了
不过我block只有一个
那麽blockIdx.x为0 跟原本只用threadIdx.x有差别吗？
还是我哪里误会了？
8F：→ Schottky: 你的 block size 是不是高於 100 却不到 200 07/12 18:46
你是指VectorAdd2<<<1, size>>>(....)
^^^这个吗
我block size设跟data size相同 原本100 结果正确 换成200後错一半
9F：→ Schottky: 这麽说好了，你把 blockDim.x 存进 result array 07/12 19:05
10F：→ Schottky: 再 printf 出来看看和你想的一样不一样... 07/12 19:05
blockIdx.x 印出来是0
blockDim.x 印出来是200 也就是data_size
threadIdx.x 印出来是0~199
跟我想的一样....
现在有点搞混了
我把idx那行改回我最原本那样 结果是正确的
我再看清楚我到底改了哪里
--------------------------------------
我比较想了解的地方是 有相依的kernel间的data 可以不用回传?
请问原理是什麽?
11F：→ Schottky: 这个问题嘛，只是你脑筋打结而已... 07/12 19:57
12F：→ Schottky: device memory 又还没有 free 掉，本来就不会平白消失 07/12 19:57
13F：→ Schottky: 跟前後 kernel 有没有相依一点关系也没有 07/12 19:58
14F：→ Schottky: 是说你害我很想买新显示卡..... 07/12 20:01
0.0 哈哈 犒赏自己 买下去!
不过回到我问的 记忆体没有free掉
所以意思是他存取同个地方罗
那麽 会存取同位址 是为什麽?
这个可以解释一下吗
15F：→ Schottky: 位址不就存在 dev_A, dev_B ... 里面，位址值有变吗? 07/12 21:20
16F：→ Schottky: 没变的话为什麽不是存取同位址 07/12 21:20
你的意思是 因为都是dev_C 记忆体位置就是那里
所以我launch VectorAdd2 後面带参数dev_C
那麽他就会去读这个记忆体位址 所以就会拿到正确的资料
是这样吗?
17F：→ Schottky: 而且 device 和 host 之间的 memcpy 非常之慢 07/12 21:24
18F：→ Schottky: 资料可以直接留给下一个 kernel 何必再传回来传过去 07/12 21:27
很慢我知道 只是我不知道原来可以这样做
19F：→ Schottky: 是的，从 allocate 到 free 之间那块记忆体都固定给你用 07/13 08:45
20F：→ Schottky: 所以你负责记住那个位址(放在dev_C)并且用完要负责free 07/13 08:45
21F：→ Schottky: 跟平常写 C 语言使用 host memory 的状况是一样的 07/13 08:46
22F：→ Schottky: kernel 凭着 dev_C 的位址就可以存取到同一块记忆体 07/13 08:48
好的 这部分我了解了 感谢你
※ 编辑: v00623 (140.118.155.204), 07/13/2017 10:16:46
23F：推 jun0325: 加个cudaThreadSynchronize();确保GPU和CPU之间的 07/16 14:04
24F：→ jun0325: data consistency。你变成data size变成200，有可能是GPU 07/16 14:05
25F：→ jun0325: 算太久，你又没保证data consistency，所以CPU读到不正确 07/16 14:06
26F：→ jun0325: 的资料 07/16 14:06