2個矩陣乘法的例子：

1. 分塊：每個block負責目標矩陣中的一塊，好處：（讀顯存的數據量：計算次數=1：小塊的邊長）；如果直接每個thread負責1個目標值，讀顯存數據量：計算次數=1：1，而且讀column顯存時可能無法連續讀顯存；（？？？好像不對）

2. 分塊矩陣把數據讀到了shared-memory，充分利用高速緩存進行矩陣乘法；減少了訪問顯存的量，使得計算：訪存是O(N^3) : O(N^3/K) ; 只要shared-memory能放的下，K越大越好；

?

1. A*X+Y的例子：https://devblogs.nvidia.com/easy-introduction-cuda-c-and-c/

注意：應對向量長度不是block大小的整數倍，1. block個數向上取整 2. kernel里用if判斷是否有活兒干

2. 卡時間：https://devblogs.nvidia.com/how-implement-performance-metrics-cuda-cc/

同步操作(cudaMemcpy)和異步操作(kernel)解釋的很好；（cudaMemcpy是同步的（阻塞的），即上面的kernel執行完畢，cudaMemcpy才能開始，cudaMemcpy執行完畢，后面的kernel才能開始；kernel是異步的，即發起調用后，控制立即返回CPU端繼續往下執行）

CPU端計時：調cudaDeviceSynchronize()等kernel執行完再打點。缺點：等待期間CPU啥也干不了，浪費了；

GPU端計時：event, 會記下這個stream執行到這里的時間點；

global-mem帶寬計算：1. 理論帶寬（同時讀寫）是查硬件手冊計算得到的，MHZ(每秒多少次傳輸，memoryClockRate）* 內存接口寬度（384bit，memoryBusWidth) * 2(DDR RAM可以讀寫同時）2. 實際帶寬，跑程序卡時間得到（也要把讀的和寫的都算進去）；（2012年的GPU，理論雙向帶寬148GB/s)

GFLOPS是針對float數據計算的，double數據的性能一般是GFLOPS的一半。A*X+Y看成是2個Float計算；這個例子讀寫12個byte才做2個float計算，顯然是memory-bound的，不是compute-bound。復雜程序性能優化建議用工具來profiling，看看bottleneck在哪兒。

3. 查GPU屬性，錯誤處理：https://mp.csdn.net/postedit/88734169

通過調API來獲取GPU參數來計算memory帶寬的例子；

nvcc編譯選項可以指定編譯成在X.X版本上跑；

每個block多少threads，很重要：太少，則每個SM上的block數（硬件規定有上限的）填不滿SM；太多，不能超過block上限thread總數（報錯）；

同步操作，會返回錯誤碼；kernel啟動等異步操作，在出錯的時候系統會往系統的變量里寫入錯誤碼，host可以同步一下（cudaDeviceSynchronize），然后調GetLastError得到錯誤碼（于此同時錯誤碼被清除了）；kernel之后立馬調的GetLastError只能返回launch相關的錯誤，不同步完的話無法返回執行期間的錯誤。cudaDeviceSynchronize拖慢速度所以盡量只在debug的時候用，release以后別用；

?

?

1. 優化數據傳輸：https://devblogs.nvidia.com/how-optimize-data-transfers-cuda-cc/

host<-->device的PCI帶寬8GB/s; device memory<-->GPU的帶寬144GB/s; 相差一個數量級；

優化原則：1. 盡量少在host和device之間傳輸數據；2.使用page-locked(pinned)主存（雙刃劍）；3.多次小傳輸打包成一次大傳輸，可減少每次傳輸的額外開銷；4.該數據傳輸可以和kernel執行或者其他數據傳輸并行起來（stream)

卡時間的方法：1.?cudaEventElapsedTime()； 2. nvprof工具查看數據傳輸耗時

對某個任務來說，GPU和CPU誰快，應該把GPU和device的數據傳輸時間也計算進去，否則不準。

host<-->device之間copy數據，CUDA無論如何會使用host的page-locked(pinned)主存，如果用戶開辟的不是這個則CUDA自動開辟一份先做一次copy，如果用戶開辟的是這個則省了這一步，會有百分之小幾十的速度提升。

開辟pinned內存，把小傳輸打包到這個大pinned內存里，再一次性傳輸；2D/3D的數據調用cudaMemcpy2D()/cudaMemcpy3D()會更快

2.把默認stream和"stream們各個操作提交順序造成的性能差異"講的很透徹的文章：https://devblogs.nvidia.com/how-overlap-data-transfers-cuda-cc/???

默認stream(NULL)就是他要執行必須之前提交的所有stream的操作都完成，他執行的時候后面所有stream提交的操作都不能開始；

對于host而言，cudaMemcpy是阻塞的，kernel提交是非阻塞的（可以和接下來的CPU計算并行），對device而言，kernel執行也是阻塞的，他執行完，后面的操作才能開始（同stream內的）；

想copy和計算并行，必須用cudaMemcpyAsync版本，且pinned host內存；

提交完一個stream的3個操作再提交下一個stream的3個操作那個例子，講的很清楚。兩種情況：copy和kernel各有1個engine對應1個隊列，H2D和D2H和kernel各有1個engine對應1個隊列；按提交順序queue到不同的engine隊列里（必須按先入先出來執行），且滿足同一個stream內的順序；Hyper-Q之后這些trick沒用了。

3. global memory:??https://devblogs.nvidia.com/how-access-global-memory-efficiently-cuda-c-kernels/

用__device__定義變量，或者用cudaMalloc()開辟動態數組；

每個硬件warp上是32個threads；warp是SIMD執行（齊步走）；早期1代顯卡要求訪問global memory要對齊且按thread順序訪問，不這樣的話就有1個數量級的帶寬損失（分成多次訪存了）；2代顯卡之后有cache，湊齊了才去訪存，所以沒有什么帶寬損失

row-major的矩陣，如果這批threads要訪問矩陣某一列，涉及到跨stride訪問，造成帶寬嚴重下降（GPU沒辦法將這些訪問請求合并到一起）；data-locality很重要！解決方法：1. 按column-major存放矩陣；2. 使用shared-memory(沒有stride降帶寬這說）；

4. shared memory:?https://devblogs.nvidia.com/using-shared-memory-cuda-cc/

bank處理好了，shared memory延遲比global memory好100倍；

同一份shared memory可以被同一個block里所有線程訪問；

block里的線程們不一定并行執行（不同warp時），先寫后讀的例子寫的快的線程讀寫的慢的線程的目標數據會race, 要用__syncthreads()來barrier; 注意：有分支的時候要確保所有threads都調了，否則會死鎖！

分配語法：1. 編譯之前知道大小，直接在kernel里定義數組即可；2. 不知道大小，則在kernel<<<>>>里告訴大小，在kernel里聲明extern __shared int s[]即可，當多個數組用的話手動設置那些數組的指針即可；

劃分成32個bank（同warp大小）; 每個bank帶寬是4B/cycle; 3.0之后可以配置bank為8B; bank之間的地址是連續的，threads訪問連續shared-memory可以跨越多個bank，從而達到最佳帶寬。多個threads同時訪問同一個bank，會被拆成多次請求（訪問同一地址例外，廣播和多播確保其一次完成）

64KB/SM的shared memory，可以被user配置成多種L1-cache和shared-memory比例；如果系統需要更多shared-memory則會強制調整。

5. 矩陣轉置例子：https://devblogs.nvidia.com/efficient-matrix-transpose-cuda-cc/

用矩陣copy來做對比，來看出來各種不同的轉置實現有多少潛力。每個sub-matrix由每個線程搬4輪，可以將kernel里計算index的開銷"均攤薄"

全用global-memory來轉置，比copy慢了一兩倍。原因：寫入的時候threads沒有“連續”訪存（典型的跨stride訪問）

使用32*32的shared-memory來中轉，讀的時候正常讀入（讀一行寫一行），寫的時候讀（讀shared-mem)一列寫(寫global-mem)一行, 使得對global-mem的訪問都是“連續”的了；速度還是比copy慢了百分之大幾十

為了調查是不是__syncthreads()造成的拖累，對矩陣copy也用shared-mem來實現并用上__syncthreads(), 發現并沒有慢，說明不是這個的原因；

真正原因：threads讀shared-memory（一列）的時候，訪問的是同一個bank; 解決：shared-mem開辟[32][32+1]，從而讓同一列的相鄰元素都分布在不同的bank; 完美達到copy例子的帶寬！

總結：兩個有效提速：global-memory的"連續"訪問，shared-memory的bank散開訪問，有效

?

6. CUDA7 Thread default stream:?https://devblogs.nvidia.com/gpu-pro-tip-cuda-7-streams-simplify-concurrency/

兩個例子：1. Thread default stream上執行kernel，不會等之前提交的所有stream上操作完成，也不會阻塞后續提交的所有stream上操作開始；2. 每個thread在自己的default stream上提交任務，可以實現多線程的stream并行執行的效果；

?

7. Unified Memory:?https://devblogs.nvidia.com/unified-memory-cuda-beginners/

老卡：cudaMallocManaged()調用的時候就在GPU上分配存儲了；CPU這邊訪問內存時，會觸發page fault，從GPU上一頁一頁搬到host上，CPU改了內存后，GPU上啟動kernel時會先自動把全部這部分內存搬到GPU里，再運行kernel； nvprof顯示的kernel耗時不包括之前的CPU->GPU內存搬家耗時；

新卡：lazy策略，cudaMallocManaged()調用時也許并不在GPU上分配存儲；用的時候才分配；GPU kernel啟動后用到的時候觸發page fault一頁一頁吧數據從CPU搬到GPU; nvprof顯示的kernel耗時就高了；

該怎么做：CPU的活兒讓GPU來干，從而減少CPU干完的東西搬到GPU的開銷；盡早prefetch即將訪問的數據；

老卡上CPU和GPU不能同時訪問同一地址（會報錯）；新卡支持page fault，因此可以同時訪問同一地址（用戶自己確保數據一致性）; 新卡支持Unified Memory上的原子操作；
另一篇好文章：https://devblogs.nvidia.com/beyond-gpu-memory-limits-unified-memory-pascal/

?

8. Cooperative Groups:?https://devblogs.nvidia.com/cooperative-groups/

A.折半reduce的例子（每一步都要sync)（很像MPI_reduce)

B.使用上面的例子(reduce_sum)來完成一個數組的加和：每個thread加和一部分數組(stride-style)得到1個數，block內部用折半reduce來把threads個數加成1個數，用原子加操作把每個block的數加到一起；

C.thread group：可以是所在的block，也可使是自定義拆分的；

D. 把B例子的block的thread-group替換成手工切分的32threads的thread-group, 改動很小；group內部的sync必須該group的所有threads都參與！所以有分支的時候，盡量拆分group，避免后續代碼的sync造成死鎖；

E. 把線程個數放到模板參數里，可以在編譯時告訴編譯器thread-group大小，從而讓編譯器去有機會優化代碼（比如循環展開，32的時候利用warp來去掉同步）

F. 使用shuffle命令來加速warp（直接使用同warp內其他thread的寄存器數據，不用shared memory了）；B例子使用shfl_down和模板參數分thread-group來實現；https://blog.csdn.net/Bruce_0712/article/details/64926471

G.?coalesced_threads可以拿到本warp內當前活躍的線程，編成一個group; 可以用于同步，可以用于選舉rank0來做事；(里面的例子用到了這個：https://devblogs.nvidia.com/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/?）（shfl可以用來做warp內廣播）

?

9. copy_if()的實現：https://devblogs.nvidia.com/cuda-pro-tip-optimized-filtering-warp-aggregated-atomics/

用原子操作遞增offset；很慢；

開一個shared memory變量（__shared__)，用原子操作遞增該變量，每個線程能有自己的offset，每個block派一個thread去原子更新這段的offset；（代碼很巧妙）；稍快一點兒；

用以下代碼實現，比thrust還快半個數量級：（如果用shuffle指令集也可實現，只是代碼多些）

__device__ int atomicAggInc(int *ctr) {auto g = coalesced_threads();int warp_res;if(g.thread_rank() == 0)warp_res = atomicAdd(ctr, g.size());return g.shfl(warp_res, 0) + g.thread_rank(); }

調用端：

if(src[i] > 0)dst[atomicAggInc(nres)] = src[i];

總結

以上是生活随笔為你收集整理的GPGPU的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。