作者v00623 (阿哩他命EX PLUS)
看板C_and_CPP
标题[问题] CUDA调整block大小 使用GPGPU-Sim
时间Mon Sep 19 20:22:32 2016
抱歉 原本的数据有错 已修改
各位前辈好
最近在使用GPGPU-Sim 3.2.2来模拟CUDA程式，想了解block大小对於程式的影响
我挑选模拟器附的Benchmark BFS来改
原本程式中的block数为256，而每个block中有256个thread
而我修改成2048个block，每个block有32个thread
原本预想thread总数相同，模拟结果应该OK，顶多执行速度变慢
但是模拟器吐出的讯息显示修改後的模拟cycle以及指令数量与原本相差许多
原本:
gpu_tot_sim_cycle = 773568
gpu_tot_sim_insn = 15889228
gpgpu_simulation_time = 0 days, 0 hrs, 4 min, 23 sec (263 sec)
修改後:
gpu_tot_sim_cycle = 697192
gpu_tot_sim_insn = 14920542
gpgpu_simulation_time = 0 days, 0 hrs, 2 min, 47 sec (167 sec)
模拟的数值差了一些 这是有问题的吗?
而且比对输出的result也不同
请问问题出在哪? 是因为block数量不同 而blockID影响结果?
我对於block数量与程式的关系还不熟悉
应该怎麽调整block以及thread的数量，而不影响程式结果呢
另外想请教GPGPU-Sim使用上遇到问题有哪里可以发问吗
目前只有看到GOOGLE Group，不过似乎不太热络?
--
※ 发信站: 批踢踢实业坊(ptt.cc), 来自: 140.118.155.204
※ 文章网址: https://webptt.com/cn.aspx?n=bbs/C_and_CPP/M.1474287759.A.F6E.html
1F：→ johnjohnlin: block size 32 会造成 occupancy 低落，导致效能不佳 09/19 22:02
2F：→ johnjohnlin: 但是改了 block size 之後差那麽多我觉得应该是 bug 09/19 22:03
恩恩 我知道occupancy会不好 官方建议block size128~256，不过我是想量测每个大小的效能影响。
所以我改的地方是对的? 我是修改<<<block, thread, shared memory>>>(参数...);这边
不过差这麽多让我觉得怪怪的 执行时间也少了一半..
※ 编辑: v00623 (114.25.14.78), 09/19/2016 22:33:24
※ 编辑: v00623 (140.118.155.204), 09/20/2016 10:43:10
3F：→ freef1y3: 如果要确认是 benchmark 还是 gpgpu-sim 的问题 09/20 10:44
4F：→ freef1y3: 我会建议用真正的 GPU 跑一次看看 09/20 10:44
5F：→ freef1y3: block 和 thread 数量这有时候是跟 benchmark 有关的 09/20 10:45
6F：→ freef1y3: 甚至有的 CUDA 程式必须 block 大小是 2 的次方才能跑 09/20 10:46
後来用模拟器多跑几种block大小 下图为结果
http://imgur.com/a/mMjkr
只看IPC的话 原本256threads的相对比较不好
有点不知道如何解释这个状况
猜想原本benchmark是否考虑到john大讲的occupancy之类的层面
※ 编辑: v00623 (140.118.155.204), 09/20/2016 11:01:50
7F：→ freef1y3: 多个小 block 比少数大 block 好 09/20 14:16
8F：→ freef1y3: 我会猜是因为 __syncthreads() 造成的 overhead 09/20 14:17
9F：→ freef1y3: 若是 occupancy 的关系，应该多个小 block 会比较差 09/20 14:18
10F：→ freef1y3: 但是看这 IPC 的差异，执行时间似乎不该差这麽多 09/20 14:19
11F：→ freef1y3: 除非执行的 Instruction 总数也会随着 block size 改变 09/20 14:20
12F：→ freef1y3: 若你的学校有授权 visual studio，且你有实体显示卡 09/20 14:23
13F：→ freef1y3: 可用 nvidia visual profiler 测测看 09/20 14:23
14F：→ freef1y3: 我没用过 gpgpu-sim，所以这方面的问题就无法回答了 09/20 14:25
程式中没有使用到__syncthreads()
https://github.com/gpgpu-sim/ispass2009-benchmarks/tree/master/BFS
後来我发现他kernel的执行次数有差 所以inst数量才会不同
http://imgur.com/a/nwIjN
不过我有个疑问 就算kernel跑的次数相同 那inst数量一定会一样吗? (这方面还不太熟悉
上图第二第四列都是跑8次kernel 不过inst还是有差一些
虽然我觉得应该还要看程式怎麽写...
※ 编辑: v00623 (140.118.155.204), 09/21/2016 10:50:20