PTT數位生活區 / C_and_CPP (C/C++)

[問題] CUDA調整block大小 使用GPGPU-Sim

看板C_and_CPP (C/C++)作者 (阿哩他命EX PLUS)時間9年前 (2016/09/19 20:22), 9年前編輯推噓0(0推 0噓 14→)
留言14則, 2人參與, 最新討論串1/1
抱歉 原本的數據有錯 已修改 各位前輩好 最近在使用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://www.ptt.cc/bbs/C_and_CPP/M.1474287759.A.F6E.html
→
09/19 22:02, , 1F
block size 32 會造成 occupancy 低落，導致效能不佳
09/19 22:02, 1F
→
09/19 22:03, , 2F
但是改了 block size 之後差那麼多我覺得應該是 bug
09/19 22:03, 2F
恩恩 我知道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
→
09/20 10:44, , 3F
如果要確認是 benchmark 還是 gpgpu-sim 的問題
09/20 10:44, 3F
→
09/20 10:44, , 4F
我會建議用真正的 GPU 跑一次看看
09/20 10:44, 4F
→
09/20 10:45, , 5F
block 和 thread 數量這有時候是跟 benchmark 有關的
09/20 10:45, 5F
→
09/20 10:46, , 6F
甚至有的 CUDA 程式必須 block 大小是 2 的次方才能跑
09/20 10:46, 6F
後來用模擬器多跑幾種block大小 下圖為結果 http://imgur.com/a/mMjkr 只看IPC的話 原本256threads的相對比較不好 有點不知道如何解釋這個狀況 猜想原本benchmark是否考慮到john大講的occupancy之類的層面 ※ 編輯: v00623 (140.118.155.204), 09/20/2016 11:01:50
→
09/20 14:16, , 7F
多個小 block 比少數大 block 好
09/20 14:16, 7F
→
09/20 14:17, , 8F
我會猜是因為 __syncthreads() 造成的 overhead
09/20 14:17, 8F
→
09/20 14:18, , 9F
若是 occupancy 的關係，應該多個小 block 會比較差
09/20 14:18, 9F
→
09/20 14:19, , 10F
但是看這 IPC 的差異，執行時間似乎不該差這麼多
09/20 14:19, 10F
→
09/20 14:20, , 11F
除非執行的 Instruction 總數也會隨著 block size 改變
09/20 14:20, 11F
→
09/20 14:23, , 12F
若你的學校有授權 visual studio，且你有實體顯示卡
09/20 14:23, 12F
→
09/20 14:23, , 13F
可用 nvidia visual profiler 測測看
09/20 14:23, 13F
→
09/20 14:25, , 14F
我沒用過 gpgpu-sim，所以這方面的問題就無法回答了
09/20 14:25, 14F
程式中沒有使用到__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
‣ 更多 v00623 的文章
文章代碼(AID): #1NtzYFzk (C_and_CPP)
C_and_CPP 近期熱門文章
更多 近期熱門文章 >>
PTT數位生活區 即時熱門文章
更多 即時熱門文章 >>
‣ 更多 v00623 的文章
文章代碼(AID): #1NtzYFzk (C_and_CPP)