--- # System prepended metadata title: 平行程式設計 HW3 楊淨 109062530 tags: [資訊工程相關, 平行程式設計] --- ###### tags `平行程式設計` `資訊工程相關` # 平行程式設計 HW3 楊淨 109062530 ## Implementation 1. Which algorithm do you choose in hw3-1? 我使用一般的Floyd-Warshall，搭配openmp 對for 平行化 ![](https://i.imgur.com/MbGVXkt.png) 2. How do you divide your data in hw3-2, hw3-3? 本次作業就是直接做題目給的範例切法 ![](https://i.imgur.com/PJWERXQ.png) 這邊加速的問題在於一個block跟shared memory可以開到多大 一個BLOCK就是 `dim3 bk(B, B);` 這樣B最大測試後只能開到32 最一開始直接用範例給的round來設計，也就是說這樣要跑到⌈n/B⌉次 `int round = ceil(n, B);` 每一次就是運算一個block量的東西，share memory 就是開到`sdata[B][B]` 但這樣做會造成shared memory使用效率不足，所以到了P22之後的所有大測資我都超時。 後來參考很多網路上的討論與論文文件後，發現了一個能將shared memory使用效率更高的切法。 這邊是以shared memory最高大約能開到多少去設定Big Block ![](https://i.imgur.com/IkKCg6K.png) 這樣一個round就會變成`int round = ceil(n, BB);` 整體要的loop量就加速2倍了，為了達到這樣子能對齊每個sm 與 block 最一開始讀資料時要先將邊長填充，這樣才好實作shared mem 的copy ![](https://i.imgur.com/2ZvxEcp.png) 有了被填充後的邊長之後圖大概是長這樣， 橘色部位是多出來被填充以方便對齊SM與GLOBLE DATA， 資料有72 * 72， SHARE MEM 是64 * 64 一個BLOCKG是32 * 32 PaddingN 是 72+56=128 ![](https://i.imgur.com/Tk2Aezu.png) 再來就是一樣根據作業3提供的那張圖去算，將綠色區域想像成下圖的灰色區域 只是不一樣的是這次一次可以算4倍的範圍，只要設定好 offset就可以進行偏移運算。 ![](https://i.imgur.com/PJWERXQ.png) ![](https://i.imgur.com/mtOnubG.png) 以PHASE1為例子，紅色就是一個THREAD會運算到的點 這樣就能充分運用 SHARED MEMORY ![](https://i.imgur.com/mCTabe2.png) ![](https://i.imgur.com/2LM7NiZ.png) 3. What’s your configuration in hw3-2, hw3-3? And why? (e.g. blocking factor, #blocks, #threads) ![](https://i.imgur.com/mtOnubG.png) 如同我在第2節提到的，因為threads上限是32*32，所以一個BLOCK就直接設定成32 * 32 再來就是 SHARED MEMORY 這邊是使用到64 * 64 INT ![](https://i.imgur.com/PJWERXQ.png) 如圖第一個PHASE的BLOCK 只要一個就好了 第二個PHASE的BLOCK 他是2個邊長，我這邊實作上只開一個邊長，但利用OFFSET去做垂直跟水平的偏移。 第三個PHASE的BLOCK 就是直接取PaddingN的ceiling平方去做剩餘的運算。 4. How do you implement the communication in hw3-3? 5. Briefly describe your implementations in diagrams, figures or sentences. 如同前面說的 資料切成這樣 在PHASE1執行這一步，紅色是指1個THREAD做的 綠色是sdata，更新全部的值 ![](https://i.imgur.com/WkaNfdV.png) ![](https://i.imgur.com/Tk2Aezu.png) 全部的thread一起做就會變成一次將左上角的這個大PIVOT完成 ![](https://i.imgur.com/kL2iZMv.png) 在PHASE2執行這步，phase2同時定位垂直與水平， 以下圖第2個BLOCK blockIdx.x == 1為例， 他會將大block的東西都計算了，包含大於總資料n的區域， 不過那些計算我們最終資料是用不到的，我們只要非橘色區域。 ![](https://i.imgur.com/YWjDQMj.png) ![](https://i.imgur.com/GS0ASk0.png) 在PHASE3執行這步，phase2的時候已經將垂直與水平邊界都算好了， 圖以 blockIdx.y == 1 blockIdx.x == 1 時為例 ![](https://i.imgur.com/pkb5mZu.png) 紫色的格子是這個block計算的東西，紅色是已經算完的區域 ![](https://i.imgur.com/FEkrbC9.png) 這樣一來只要round跑完，中心就會由左上算到右下。 --- ## Profiling Results (hw3-2) Provide the profiling results of following metrics on the biggest kernel of your program using NVIDIA profiling tools. 這一部分我拿P21K1的資料來當測試，我認為這個資料量剛剛好 occupancy ![](https://i.imgur.com/1ogr9iL.png) sm efficiency ![](https://i.imgur.com/gyYkHqK.png) shared memory load/store throughput ![](https://i.imgur.com/r8ykkdq.png) global load/store throughput ![](https://i.imgur.com/PoHeb1V.png) ![](https://i.imgur.com/i0R2HAm.png) ## Experiment & Analysis 1. System Spec 這一部分我是使用實驗室的機器APOLLO(3-1)與HADES(3-2 3-3) ![](https://i.imgur.com/ugbwgv8.png) Hades: ![](https://i.imgur.com/SkLJ5d7.png) ![](https://i.imgur.com/ur85hQ4.png) ![](https://i.imgur.com/oA14N7a.png) ![](https://i.imgur.com/6OlJbl1.png) 2. ### Blocking Factor (hw3-2) Observe what happened with different blocking factors, and plot the trend in terms of Integer GOPS and global/shared memory bandwidth. (You can get the information from profiling tools or manual) (You might want to check nvprof and Metrics Reference) 因為p21k1 的資料量太大，一直超時不給跑 這部分我拿 p11k1 進行 --metrics inst_integer 計算再除以執行時間 首先 B=32 * 32 一個kernel 加起來共 13304960768 約133億次*172 /耗時1.2秒 =19092 GOPS ![](https://i.imgur.com/Qr8Br1S.png) ![](https://i.imgur.com/Gdy1VG1.png) B=16 * 16 ![](https://i.imgur.com/0Q3IPOB.png) ![](https://i.imgur.com/0JHxgCT.png) B=8 * 8 跑到超時 無法計算 ![](https://i.imgur.com/QmRrZ61.png) 使用的blockfactor: B * B | 8 * 8 = 64 | 16 * 16 = 256 | 32 * 32 = 1024 | | -------- | -------- | -------- | | X | 12040 GOPS | 19092 GOPS | ![](https://i.imgur.com/SfCdNDf.png) ### global/shared memory bandwidth B1024 ![](https://i.imgur.com/BFrZjuT.png) B256 ![](https://i.imgur.com/Y2iLsh1.png) B64 ![](https://i.imgur.com/wAJeXfS.png) ![](https://i.imgur.com/EaEx4x9.png) ![](https://i.imgur.com/xDFxivu.png) 3. Optimization (hw3-2) Any optimizations after you port the algorithm on GPU, describe them with sentences and charts. Here are some techniques you can implement: 這一部分我以P11K1當範例，比較了幾個版本: * CPU版本的 * 最初完成的CUDA版本(SM與BLOCK一樣大) * 最終版本(4倍SM以在一個THREAD中進行偏移計算) ![](https://i.imgur.com/0Ktg3q8.png) ![](https://i.imgur.com/0k1N79H.png) 4. Weak scalability (hw3-3) Observe weak scalability of the multi-GPU implementations 5. Time Distribution (hw3-2) Analyze the time spent in: computing communication memory copy (H2D, D2H) I/O of your program w.r.t. input size. 以測資p11k1 p16k1 p21k1 p25k1 為例 | | p11k1 | p16k1 | p21k1 | p25k1 | | -------- | -------- | -------- | -------- |-------- | | V | 11000 | 16000 | 20959 | 25000 | | E | 505586 | 5444739 | 987205 | 5780158 | ![](https://i.imgur.com/3g4cyBe.png) ![](https://i.imgur.com/4BPtLJF.png) ## Experience & conclusion What have you learned from this homework? 本次作業學到了不少與CUDA操作有關的用法，比起之前的LAB，這次最大的收穫在於實作SHARED MEMORY拓展後，更清楚如何運用SHARED MEMORY來加速運算，還有BLOCK大小的各項影響，原本認識的平行化都只是簡單的將FOR迴圈想辦法拆分，但這次作業為了將東西分散到GPU上的BLOCK，整個思維需要跳開好幾階，一開始一直想不到怎麼加速才好。 Feedback (optional) 希望作業報告能給再多一些實作的範例或步驟，有些東西沒碰過實在不太清楚要如何實作，學生可能會多花不少時間去猜要怎麼做才對。