2017-03-23 13:53
1052人閱讀收藏
舉報cuda（32）
從硬體角度分析，支援cuda的nvidia 顯示卡，都是由多個multiprocessors 組成。每個 multiprocessor 裡包含了8個stream processors，其組成是四個四個一組，也就是兩組4d的處理器。每個 multiprocessor 還具有 很多個（比如8192個）暫存器，一定的（比如16kb） share memory，以及 texture cache 和 constant cache。
在 cuda 中，大部份基本的運算動作，都可以由 stream processor 進行。每個 stream processor 都包含乙個 fma（fused-multiply-add）單元，可以進行乙個乘法和乙個加法。比較複雜的運算則會需要比較長的時間。
在執行 cuda 程式的時候，每個 stream processor 就是對應乙個 thread。每個 multiprocessor 則對應乙個 block。但是我們乙個block往往有很大量的執行緒，之前我們用到了256個和1024個，遠超乙個 multiprocessor 所有的8個 stream processor 。
實際上，雖然乙個 multiprocessor 只有八個 stream processor，但是由於 stream processor 進行各種運算都有 latency，更不用提記憶體訪問的 latency，因此 cuda 在執行程式的時候，是以warp 為單位。
比如乙個 warp 裡面有 32 個 threads，分成兩組 16 threads 的 half-warp。由於 stream processor 的運算至少有 4 cycles 的 latency，因此對乙個 4d 的stream processors 來說，一次至少執行 16 個 threads（即 half-warp）才能有效隱藏各種運算的 latency（ 如果你開始運算，再開乙個執行緒，開始運算，再開乙個執行緒，開始運算，再開乙個執行緒開始運算，這時候第乙個執行緒就ok了，第乙個執行緒再開始運算 ， 看起來就沒有延遲了， 每個處理單元上最少開4個可以達到隱藏延遲的目的，也就是4*4=16個執行緒）。也因此，執行緒數達到隱藏各種latency的程度後，之後數量的提公升就沒有太大的作用了。
問：同乙個sm中某乙個時刻是不是只能執行同乙個block裡面的wrap呢？如果是的話，同乙個時刻sm是否能夠同時執行同乙個block裡面的多個wrap呢？如果是的話，由於同乙個block共享32個bank，這樣的話不同wrap，有沒有可能產生bank confick呢，例如乙個wrap0-31中的a[0]與第二個32-63wrap中的a[32]訪問的是同乙個bank？謝謝~
答：（1）否。乙個sm可以執行來自多個blocks的warps的。而不是只能乙個。這點你可以參考手冊裡的計算能力區別**那裡，可以清晰的看到不同計算能力的卡，能同時上多少個blocks。而它們的數字都不是1.（2）否。不同的blocks間的warps不會導致bank conflict. 目前只有乙個warp內部的執行緒間才有可能會導致bank conflict.而不同warp不能同時導致bank conflict的原因是它們不能同時訪問同乙個sm裡的shared memory的。同乙個sm裡的shared memory只能同時服務來自1個warp的請求。實際上，我們考慮shared memory是否有bank衝突的時候，只需要看某個warp的某個瞬間即可。 而無需考慮本次訪問和它自身的之前、之後的訪問以及和來自其他warp的訪問之間的關係。 硬體本身的特性決定了，這些不能同時進行。不能同時進行自然也無bank conflict的。 你可以寫個簡單的程式，模擬你的懷疑，讓執行緒31和執行緒32（這歸屬於2個不同的warp了）能同時下標對映到同乙個bank上， 然後跑下profiler, 看下這兩個events的計數資訊: shared_ld_bank_conflict和shared_st_bank_conflict你會發現這兩個數值為0.
問： 好的，謝謝，關於第乙個問題還有點疑問：乙個sm可以執行來自多個blocks的warps的，這個是同乙個時刻執行來自多個blocks的warps？還是乙個block結束後退出，然後接著執行下個？還點疑問，shared_ld_bank_conflict與shared_st_bank_conflict中ld與st代表什麼呢？（這應該還是算第二次提問的延續吧，第三次我去論壇
答：可以同時執行來自多個blocks的多個warps的。不是必須需要等待乙個block完成才能進行下乙個的。 但是資源的分配是以block為單位的，哪怕該block最後還殘存1個warp在執行， 也需要等待此warp結束，才能block集體結束。 這會導致效能分析工具報告的achieved occupancy下降的。如果你寫**之類的話，可以在kernel結束前新增乙個__syncthreads()，這會有效的提公升achieved occupancy （但這對實際效能無任何用途）（只是數字好看了）（適合匯報使用）對kepler到pascal有效。
CUDA之Thread Wrap執行詳解
從硬體角度分析，支援cuda的nvidia 顯示卡，都是由多個multiprocessors 組成。每個 multiprocessor 裡包含了8個stream processors，其組成是四個四個一組，也就是兩組4d的處理器。每個 multiprocessor 還具有 很多個 比如8192個 暫...
CUDA學習筆記之 CUDA儲存器模型
gpu片內 register，shared memory 板載視訊記憶體 local memory,constant memory,texture memory,texture memory,global memory host 記憶體 host memory,pinned memory.cuda儲...
CUDA學習筆記之 CUDA儲存器模型
標籤 cuda 儲存binding cache 程式設計api 2010 01 03 20 32 5577人閱讀收藏 舉報 cuda 6 cuda儲存器模型 gpu片內 register，shared memory 板載視訊記憶體 local memory,constant memory,textu...