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