# 平行程式 HW3 環境: 依作業提供之apollo server。 ## 3-1 ### Implementation #### Which algorithm do you choose in hw3-1? 使用FLOYD-WARSHALL演算法，Floyd-Warshall演算法是一種用於計算加權圖中所有頂點對之間最短路徑的演算法。它的主要特點是能夠處理帶有負權邊的圖，但要求圖中沒有負權迴路。以下是Floyd-Warshall演算法的基本步驟： 1. **初始化**：建立一個距離矩陣，初始時將每對頂點之間的距離設為無限大（如果兩個頂點之間沒有直接邊），並將每個頂點到自身的距離設為0。如果有直接邊，則將距離設為該邊的權重。 2. **動態規劃**：對於每一個中間頂點 $( k )$，更新距離矩陣，使得對於每一對頂點 $( (i, j))$，如果通過頂點 $( k )$ 的路徑比直接從 $( i )$ 到 $( j )$ 的路徑更短，則更新距離矩陣中的值。具體公式為： $\text{dist}[i][j] = \min(\text{dist}[i][j], \text{dist}[i][k] + \text{dist}[k][j])$ 3. **重複更新**：重複上述步驟，對所有頂點作為中間點進行更新，直到所有頂點對之間的最短距離都被計算出來。 這樣，最終的距離矩陣將包含所有頂點對之間的最短路徑距離。Floyd-Warshall演算法的時間複雜度為 $(O(n^3))$，其中 $(n)$ 是圖中的頂點數。 Sudo code ```c FLOYD-WARSHALL( W ): n = W.rows D(0) = W for k = 1 to n let D(k) = = [ d[i][j](k) ] be a new n x n matrix for i = 1 to n for j = 1 to n d[i][j](k) = min(d[i][j](k-1), d[i][k](k) + d[k][j](k-1)) return D(n) ``` #### 實作 * 主函數： * 讀取輸入檔案並初始化距離矩陣。 * 獲取CPU的執行緒數量並初始化障礙變數。 * 創建多個執行緒，每個執行緒執行block_FW函數。 * 等待所有執行緒完成計算。 * 將結果寫入輸出檔案。 ```clike= int main(int argc, char* argv[]) { if (pthread_barrier_init(&barrier, NULL, NUM_THREADS) != 0) { printf("Barrier initialization failed\n"); return 1; } //..... for(int i = 0 ; i < NUM_THREADS ; i++){ ids[i] = i; pthread_create(&threads[i], nullptr, block_FW,(void*)&ids[i]); } // gether threads for(int i = 0 ; i < NUM_THREADS ; i++){ pthread_join(threads[i], nullptr); } output(argv[2]); return 0; } void* block_FW(void* args) { int pid = *(int*)args; //for(int k = pid ; k < n ; k+=NUM_THREADS){ // 不能用Round，會卡pthread barrier for(int k = 0 ; k < n ; k++){ //for(int i = 0 ; i < n ; i++){ for(int i = pid ; i < n ; i+=NUM_THREADS){ for(int j = 0 ; j < n ; j++){ if (Dist[i][k] + Dist[k][j] < Dist[i][j]) { Dist[i][j] = Dist[i][k] + Dist[k][j];} } } //因為下一個round的結果depands on 上一個round pthread_barrier_wait(&barrier); } pthread_exit(nullptr); return NULL; } ``` #### 心得 因為下一個round的結果depands on 上一個round的結果，所以需在每個round結束時放barrier。最外層的迴圈(k)代表round次數，需注意的是不能把round次數分散在各pthreads，因為若有thread提早跳出for loop，便可能導致其他threads卡死在`pthread_barrier_wait`。 ---- ## 3-2 此部分包含2種實作的實驗，第一種是依照作業說明的每次launch phase1~3，重複次數等於$\frac{n}{B}$，其中B代表block數量。第二種是展開整個matrix，launch次數$n$次，但每個kernel內容較簡單。 本report包含實作並比較兩種極端做法的performance差異：(1).Kernel launch較少，但每個kernel複雜、(2)Kernel launch較多，但每個Kernel簡單。 ### 作法(一): Cuda實做Blocked FLOYD-WARSHALL 屬於Kernel launch較少，但每個kernel複雜的作法。 原理:每個iteration會launch 3個phase的cuda函數，每個pahse的函數各自處理$B\times B$大小的sub-matrix。 ![image](https://hackmd.io/_uploads/ByeNBueGJg.png) ```cpp= #include <stdio.h> #include <stdlib.h> #include <cuda_runtime.h> #include <iostream> //====================== const int INF = ((1 << 30) - 1); //const int V = 50010; int* Dist; int n, m; int ceil(int a, int b) { return (a + b - 1) / b; } __device__ void calKernel(int* dist, int B, int Round, int block_start_x, int block_start_y, int block_width, int block_height , int n) { int b_i = blockIdx.x * blockDim.x + threadIdx.x + block_start_x; int b_j = blockIdx.y * blockDim.y + threadIdx.y + block_start_y; if (b_i < block_start_x + block_height && b_j < block_start_y + block_width) { for (int k = Round * B; k < (Round + 1) * B && k < n; ++k) { int block_internal_start_x = b_i * B; int block_internal_end_x = (b_i + 1) * B; int block_internal_start_y = b_j * B; int block_internal_end_y = (b_j + 1) * B; if (block_internal_end_x > n) block_internal_end_x = n; if (block_internal_end_y > n) block_internal_end_y = n; for (int i = block_internal_start_x; i < block_internal_end_x; ++i) { for (int j = block_internal_start_y; j < block_internal_end_y; ++j) { if (dist[i * n + k] + dist[k * n + j] < dist[i * n + j]) { dist[i * n + j] = dist[i * n + k] + dist[k * n + j]; } } } } } } __global__ void blockedFloydWarshallPhase1(int* dist, int round ,int r , int B , int n){ calKernel(dist,B,r,r,r,1,1,n); } __global__ void blockedFloydWarshallPhase2(int* dist, int round ,int r , int B , int n){ calKernel(dist,B, r, r, 0, r, 1,n); calKernel(dist,B, r, r, r + 1, round - r - 1, 1,n); calKernel(dist,B, r, 0, r, 1, r,n); calKernel(dist,B, r, r + 1, r, 1, round - r - 1,n); } __global__ void blockedFloydWarshallPhase3(int* dist, int round ,int r , int B , int n){ calKernel(dist,B, r, 0, 0, r, r,n); calKernel(dist,B, r, 0, r + 1, round - r - 1, r,n); calKernel(dist,B, r, r + 1, 0, r, round - r - 1,n); calKernel(dist,B, r, r + 1, r + 1, round - r - 1, round - r - 1,n); } #define B 512 int main(int argc, char* argv[]) { input(argv[1]); //============================= // cuda memory //============================= int* device_D = NULL; cudaHostRegister(&Dist[0], sizeof(int) * n*n, cudaHostRegisterDefault); cudaMalloc((void **)&device_D, sizeof(int) * n*n); cudaMemcpy(device_D , &Dist[0], sizeof(int) * n*n , cudaMemcpyHostToDevice); //============================= // Cuda launch //============================= dim3 threadsPerBlock(16, 16); dim3 numBlocks((n + threadsPerBlock.x - 1) / threadsPerBlock.x, (n + threadsPerBlock.y - 1) / threadsPerBlock.y); int round = ceil(n, B); for (int r = 0; r < round; ++r) { // Phase 1 blockedFloydWarshallPhase1<<<numBlocks, threadsPerBlock>>>(device_D , round,r,B,n ); // Phase 2 blockedFloydWarshallPhase2<<<numBlocks, threadsPerBlock>>>(device_D, round,r,B,n ); // Phase 3 blockedFloydWarshallPhase3<<<numBlocks, threadsPerBlock>>>(device_D, round,r,B,n ); } cudaMemcpy(Dist, device_D, sizeof(int)*n*n, cudaMemcpyDeviceToHost); cudaHostUnregister(&Dist[0]); cudaFree(device_D); output(argv[2]); delete[] Dist; return 0; } ``` #### 潛在缺點 1. Uncoalesced index 2. Not using shadered memory 以下圖來說明當k=0時，一個block內程式主要的access pattern: `dist[i * n + k] + dist[k * n + j] < dist[i * n + j]` * i * n + k:黑框區域，其column-major的特性不利於memory access * k * n + j:紅框區域，雖然是row-major，但每個k-iteration會經歷i次非連續的記憶體access。 * i * n + j:籃框區域。 如此雜亂的access pattern組合，若不先把dist的內容放到shared-memory，則會大大的被Memory access時間拖慢。 ![image](https://hackmd.io/_uploads/B10ncdlz1e.png) #### Unroll對Performance的影響 在主要的核心方法上加上`unroll`指示如下： ```cpp= __device__ void calKernel(int* dist, int B, int Round, int block_start_x, int block_start_y, int block_width, int block_height , int n) { int b_i = blockIdx.x * blockDim.x + threadIdx.x + block_start_x; int b_j = blockIdx.y * blockDim.y + threadIdx.y + block_start_y; if (b_i < block_start_x + block_height && b_j < block_start_y + block_width) { #pragma unroll 32 for (int k = Round * B; k < (Round + 1) * B && k < n; ++k) { int block_internal_start_x = b_i * B; int block_internal_end_x = (b_i + 1) * B; int block_internal_start_y = b_j * B; int block_internal_end_y = (b_j + 1) * B; if (block_internal_end_x > n) block_internal_end_x = n; if (block_internal_end_y > n) block_internal_end_y = n; #pragma unroll 32 for (int i = block_internal_start_x; i < block_internal_end_x; ++i) { #pragma unroll 32 for (int j = block_internal_start_y; j < block_internal_end_y; ++j) { if (dist[i * n + k] + dist[k * n + j] < dist[i * n + j]) { dist[i * n + j] = dist[i * n + k] + dist[k * n + j]; } } } } } } ``` | 方法 | c10.1 | c20.1 | | -------- | -------- | -------- | | 無unroll | 15.86 | timeout | | unroll | 15.56 | timeout | 並無提升的原因可能有: * 迴圈迭代次數：如果迴圈迭代次數不是32的倍數，編譯器可能無法完全展開迴圈，導致代碼效率低下。 * **記憶體訪問模式**：如果展開的迴圈導致非合併的記憶體訪問(Uncoalesced index)，性能可能會下降。確保記憶體訪問是合併的對性能至關重要。 * 指令快取：展開迴圈會增加代碼大小，這可能導致指令快取未命中，對性能產生負面影響。 * 編譯器優化：編譯器可能已經在有效地優化迴圈。手動展開迴圈可能會干擾這些優化。 可發現**Uncoalesced index問題**對於後續優化的影響極大。 --- #### B的設定對Performance的影響 B的設定控制了每個sub-matrix要切成$B\times B$大小。 | B | c10.1 | c20.1 | | -------- | -------- | -------- | | 512 | 15.86 | timeout | | 256 | 15.51 | timeout | | 64 | 2.22 | timeout | | 32 | 0.82 | 23.88 | | 16 | 0.42 | **timeout** | | 8 | 0.32 | 19.97 | | 4 | 0.27 | 11.90 | | **2** | **0.22** | **8.64** | | 1 | 0.22 | 14.05 | 經過實驗，當**B=2**是最好的，可能原因是較小的B能提升記憶體access pattern的緊密程度(跳躍幅度較小)。 --- ### 作法(二) Cuda實做FLOYD-WARSHALL 屬於 Kernel launch較多，但每個Kernel簡單。 原理: 回歸到原本的FLOYD-WARSHALL演算法取代多個Phase，以提升程式的可讀性，並簡化問題。 ```cpp= __global__ void floydWarshallKernel(int* dist, int k, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < n && j < n) { int ik = i * n + k; int kj = k * n + j; int ij = i * n + j; if (dist[ik] + dist[kj] < dist[ij]) { dist[ij] = dist[ik] + dist[kj]; } } dim3 threadsPerBlock(16, 16); dim3 numBlocks((n + threadsPerBlock.x - 1) / threadsPerBlock.x, (n + threadsPerBlock.y - 1) / threadsPerBlock.y); for (int k = 0; k < n; ++k) { floydWarshallKernel<<<numBlocks, threadsPerBlock>>>(device_D, k, n); cudaDeviceSynchronize(); } ``` 在雙方皆無任何優化的情況下，此做法能跟Blocked FLOYD-WARSHALL有相近的Performance且無須調整參數B。 | 方法 |c10.1 | c20.1 |p11k1| | -------- | -------- | -------- |-------- | | 作法1:Blocked FLOYD-WARSHALL with B=32|0.82 | 23.88 |Timeout | | 作法1:Blocked FLOYD-WARSHALL with B=2 |0.22 | 8.64 |Timeout | | 作法2:FLOYD-WARSHALL | 0.22 |10.45 |Timeout| | 方法 | 評分系統總體時間 | | -------- | -------- | | Blocked FLOYD-WARSHALL with B=2 | 25.13 | | FLOYD-WARSHALL | 29.35 | #### 嘗試用Shared memory ```cpp= __device__ void calKernel(int* dist, int Round, int block_start_x, int block_start_y, int block_width, int block_height, int n) { extern __shared__ int sharedMem[]; int b_i = blockIdx.x * blockDim.x + threadIdx.x + block_start_x; int b_j = blockIdx.y * blockDim.y + threadIdx.y + block_start_y; int tid = threadIdx.y * blockDim.x + threadIdx.x; if (b_i < block_start_x + block_height && b_j < block_start_y + block_width) { for (int k = Round * B; k < (Round + 1) * B && k < n; ++k) { int block_internal_start_x = b_i * B; int block_internal_end_x = (b_i + 1) * B; int block_internal_start_y = b_j * B; int block_internal_end_y = (b_j + 1) * B; if (block_internal_end_x > n) block_internal_end_x = n; if (block_internal_end_y > n) block_internal_end_y = n; // Load data into shared memory if (threadIdx.x < B && threadIdx.y < B) { sharedMem[threadIdx.y * B + threadIdx.x] = dist[block_internal_start_x * n + block_internal_start_y + threadIdx.y * n + threadIdx.x]; } __syncthreads(); for (int i = block_internal_start_x; i < block_internal_end_x; ++i) { for (int j = block_internal_start_y; j < block_internal_end_y; ++j) { int shared_i = i - block_internal_start_x; int shared_j = j - block_internal_start_y; if (sharedMem[shared_i * B + k] + sharedMem[k * B + shared_j] < sharedMem[shared_i * B + shared_j]) { sharedMem[shared_i * B + shared_j] = sharedMem[shared_i * B + k] + sharedMem[k * B + shared_j]; } } } __syncthreads(); // Write data back to global memory if (threadIdx.x < B && threadIdx.y < B) { dist[block_internal_start_x * n + block_internal_start_y + threadIdx.y * n + threadIdx.x] = sharedMem[threadIdx.y * B + threadIdx.x]; } } } } ``` | 方法 | c10.1 | c20.1 | | -------- | -------- | -------- | | 用Shared memory | timeout | timeout | | 無優化 | 0.22 | 8.64 | ---- ## Profling 分別對上述的Blocked FLOYD-WARSHALL與FLOYD-WARSHALL 做Profling。以data c10.1與c20.1來看。 ### Nsight Data * Algorithm : BLOCK * Data: C10.1 ![image](https://hackmd.io/_uploads/rJspItVMJe.png) --- * Algorithm : BLOCK * Data: C20.1 ![image](https://hackmd.io/_uploads/HJHNsFEGJg.png) ### Performance分析 Performance分析分為CPU端與GPU端的面向。 #### CPU 以CPU端來看，有近9成的時間在launch kernel的for-loop中。 | Time | Total Time | Num Calls | Avg | Med | Min | Max | StdDev | Name | | --- | --- | --- | --- | --- | --- | --- | --- | --- | | 85.90% | 7.581 s | 7500 | 1.011 ms | 390.149 μs | 2.769 μs | 4.204 ms | 1.172 ms | cudaLaunchKernel | | 12.70% | 1.123 s | 2 | 561.465 ms | 561.465 ms | 8.229 ms | 1.115 s | 782.393 ms | cudaMemcpy | | 1.30% | 114.300 ms | 1 | 114.300 ms | 114.300 ms | 114.300 ms | 114.300 ms | 0 ns | cudaHostRegister | | 0.00% | 2.881 ms | 1 | 2.881 ms | 2.881 ms | 2.881 ms | 2.881 ms | 0 ns | cudaHostUnregister | | 0.00% | 197.125 μs | 1 | 197.125 μs | 197.125 μs | 197.125 μs | 197.125 μs | 0 ns | cudaFree | | 0.00% | 161.312 μs | 1 | 161.312 μs | 161.312 μs | 161.312 μs | 161.312 μs | 0 ns | cudaMalloc | | 0.00% | 878 ns | 1 | 878 ns | 878 ns | 878 ns | 878 ns | 0 ns | cuModuleGetLoadingMode | --- #### GPU GPU有8成的時間在執行Phase3，而Phase2與Phase1各占約10%。因為Phase3執行phase1與phase2之外的較多Pixel導致。 | Time | Total Time | Instances | Avg | Med | Min | Max | StdDev | Name | | --- | --- | --- | --- | --- | --- | --- | --- | --- | | 80.60% | 7.001 s | 2500 | 2.800 ms | 2.773 ms | 2.516 ms | 4.234 ms | 207.766 μs | blockedFloydWarshallPhase3(int *, int, int, int) | | 11.30% | 979.608 ms | 2500 | 391.843 μs | 391.270 μs | 387.718 μs | 455.750 μs | 4.960 μs | blockedFloydWarshallPhase2(int *, int, int, int) | | 8.10% | 707.639 ms | 2500 | 283.055 μs | 282.628 μs | 279.843 μs | 327.524 μs | 3.423 μs | blockedFloydWarshallPhase1(int *, int, int, int) | 雖然乍看之下會下意識覺得要優化Phase3的工作量，但若從所有的工作量佔比來看的話，耗時最久的是Kernel launch的overhead。即使Blocked FLOYD-WARSHALL能有效減少Kernel launch數量， | Time | Total Time | Instances | Avg | Med | Min | Max | StdDev | Category | Operation | | --- | --- | --- | --- | --- | --- | --- | --- | --- | --- | | 43.30% | 7.581 s | 7500 | 1.011 ms | 390.149 μs | 2.769 μs | 4.204 ms | 1.172 ms | CUDA_API | cudaLaunchKernel | | 39.90% | 7.001 s | 2500 | 2.800 ms | 2.773 ms | 2.516 ms | 4.234 ms | 207.766 μs | CUDA_KERNEL | blockedFloydWarshallPhase3(int *, int, int, int) | | 6.40% | 1.123 s | 2 | 561.465 ms | 561.465 ms | 8.229 ms | 1.115 s | 782.393 ms | CUDA_API | cudaMemcpy | | 5.60% | 979.608 ms | 2500 | 391.843 μs | 391.270 μs | 387.718 μs | 455.750 μs | 4.960 μs | CUDA_KERNEL | blockedFloydWarshallPhase2(int *, int, int, int) | | 4.00% | 707.639 ms | 2500 | 283.055 μs | 282.628 μs | 279.843 μs | 327.524 μs | 3.423 μs | CUDA_KERNEL | blockedFloydWarshallPhase1(int *, int, int, int) | | 0.70% | 114.300 ms | 1 | 114.300 ms | 114.300 ms | 114.300 ms | 114.300 ms | 0 ns | CUDA_API | cudaHostRegister | | 0.00% | 8.186 ms | 1 | 8.186 ms | 8.186 ms | 8.186 ms | 8.186 ms | 0 ns | MEMORY_OPER | [CUDA memcpy Host-to-Device] | | 0.00% | 7.635 ms | 1 | 7.635 ms | 7.635 ms | 7.635 ms | 7.635 ms | 0 ns | MEMORY_OPER | [CUDA memcpy Device-to-Host] | | 0.00% | 2.881 ms | 1 | 2.881 ms | 2.881 ms | 2.881 ms | 2.881 ms | 0 ns | CUDA_API | cudaHostUnregister | | 0.00% | 197.125 μs | 1 | 197.125 μs | 197.125 μs | 197.125 μs | 197.125 μs | 0 ns | CUDA_API | cudaFree | | 0.00% | 161.312 μs | 1 | 161.312 μs | 161.312 μs | 161.312 μs | 161.312 μs | 0 ns | CUDA_API | cudaMalloc | | 0.00% | 878 ns | 1 | 878 ns | 878 ns | 878 ns | 878 ns | 0 ns | CUDA_API | cuModuleGetLoadingMode | Kernel launch 的 overhead 會佔很大比例的原因有幾個： 1. **CPU-GPU 通訊延遲**：每次 kernel launch 都需要 CPU 和 GPU 之間的通訊，這會產生一定的延遲。 2. **上下文切換**：每次 kernel launch 都會涉及上下文切換，這會增加額外的延遲。 3. **同步開銷**：在某些情況下，kernel launch 需要進行同步操作，這會進一步增加 overhead。 4. **小 kernel 的影響**：對於執行時間較短的小 kernel，launch overhead 會顯得特別明顯，因為 kernel 執行時間相對較短。 #### GPU Memory 用量 Blocked FLOYD-WARSHALL與FLOYD-WARSHALL一樣在測資C20.1共使用了約95.37MB存放資料。 | Total | Count | Avg | Med | Min | Max | StdDev | Operation | | --- | --- | --- | --- | --- | --- | --- | --- | | 95.37 MiB | 1 | 95.37 MiB | 95.37 MiB | 95.37 MiB | 95.37 MiB | 0 B | [CUDA memcpy Device-to-Host] | | 95.37 MiB | 1 | 95.37 MiB | 95.37 MiB | 95.37 MiB | 95.37 MiB | 0 B | [CUDA memcpy Host-to-Device] | --- ## HW3-3 Multi-GPU 我只有成功做出資料切割的部分。 ### 資料切割 範例: 將原本5*5的matrix切成上下兩部分。 ![image](https://hackmd.io/_uploads/SJb54qGrJx.png) 以下程式能正確將資料分散在2個gpu上。 ```cpp! int main(int argc, char* argv[]) { input(argv[1]); int* device_D[2]; int n_per_device = n / 2 ; for (int i = 0; i < 2; ++i) { int start_index = i * n * n_per_device; if(i%2==1 && n%2==1){ n_per_device++; } cudaSetDevice(i); cudaMalloc((void**)&device_D[i], sizeof(int) * n * n_per_device); cudaMemcpy(device_D[i], &Dist[start_index], sizeof(int) * n * n_per_device, cudaMemcpyHostToDevice); } dim3 threadsPerBlock(32, 32); dim3 numBlocks( (n_per_device + threadsPerBlock.x - 1) / threadsPerBlock.x, (n + threadsPerBlock.y - 1) / threadsPerBlock.y); #pragma omp parallel num_threads(2) { int gpu_id = omp_get_thread_num(); cudaSetDevice(gpu_id); runFloydWarshall(device_D[gpu_id], n, numBlocks, threadsPerBlock); } n_per_device = n / 2 ; for (int i = 0; i < 2; ++i) { int start_index = i * n * n_per_device; if(i%2==1 && n%2==1){ n_per_device++; } cudaSetDevice(i); cudaMemcpy(&Dist[start_index], device_D[i], sizeof(int) * n * n_per_device, cudaMemcpyDeviceToHost); cudaFree(device_D[i]); } output(argv[2]); delete[] Dist; return 0; } ``` 我沒有成功做出來，但以下闡述相關構想。 * 構想: 將Kernel launch分散給各GPU device。 ```cpp for (int i = 0; i < numDevices; ++i) { cudaSetDevice(i); for (int k = i ; k < n; k+=numDevices) { floydWarshallKernel<<<numBlocks, threadsPerBlock>>>(device_D[i], k, n); cudaDeviceSynchronize(); } cudaDeviceSynchronize(); } ``` 但由於這樣每隔會少跟另一半的k做比較，因此需再用個function比較最短的。 ```cpp __global__ void cmp(int* dist1 , int* dist2 , int n){ int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; printf("cmp i %d j %d " , i , j); if (i < n && j < n) { int ij = i * n + j; printf("dist1[ij] %d dist2[ij] %d \n" , dist1[ij], dist2[ij]); if(dist1[ij] > dist2[ij]){ dist1[ij] = dist2[ij]; } } } .... cudaSetDevice(0); cmp<<<numBlocks , threadsPerBlock>>>(device_D[0] , device_D[1] , n); cudaDeviceSynchronize(); ``` --- * 構想2: 由上面Nsight-system的數據可知道Phase3佔了最大一部份，因此這邊把Phase3額外拆給另一個GPU做。 ```cpp= for (int r = 0; r < round; ++r) { { cudaSetDevice(0); blockedFloydWarshallPhase1<<<numBlocks, threadsPerBlock>>>(device_D[0], round, r, n); blockedFloydWarshallPhase2<<<numBlocks, threadsPerBlock>>>(device_D[0], round, r, n); cudaDeviceSynchronize(); cudaSetDevice(1); blockedFloydWarshallPhase3<<<numBlocks, threadsPerBlock>>>(device_D[0], round, r, n); cudaDeviceSynchronize(); } } ``` ### Weak scalability Judge執行時間(s)關係如下： |方法 | c01.1 | c04.1 | | -------- | -------- | -------- | | Single-GPU | 0.27 | 8.79 | | 構想1 | 0.37 | 5.33 | | 構想2 | 0.32 | 15.01 | --- ## Experiment on AMD GPU 將hw3-2編譯成.hip檔案 ``` hipify-clang hw3-2.cu ``` 測資: c20.1 ### 測量方法# 1 在頭尾加上計時程式: ``` #include <chrono> int main(){ auto start = std::chrono::high_resolution_clock::now(); .... auto end = std::chrono::high_resolution_clock::now(); std::chrono::duration<double> duration = end - start; printf("Your code runs for : %f secondes \n" , duration.count()); } ``` ### 測量方法# 2 使用Nsight System。 由於是執行AMD的程式，CUDA event無法捕捉，使NSight system的資訊只剩下CPU端的執行時間。 ![image](https://hackmd.io/_uploads/rJqSJo2Mkg.png) | 方法 | c20.1執行時間(s) |c01.1執行時間(s) | | -------- | -------- | -------- | | Nvidia | 8.64 |0.27| | AMD | 1.31 |0.02| ## 心得 雖然將程式轉換為AMD版本看似速度很快，但其輸出結果卻是錯誤的（Nvidia版本的輸出是正確的）。我認為在現階段，直接將Nvidia程式轉換為AMD程式仍存在一定的風險。 --- ## Experience & conclusion 這頁筆記主要涵蓋以下幾個重點： - **Floyd-Warshall演算法**：詳細介紹了Floyd-Warshall演算法的步驟和實作，包括初始化距離矩陣、動態規劃更新距離矩陣、以及重複更新直到計算出所有頂點對之間的最短路徑。 - **CUDA實作**：展示了如何使用CUDA進行Blocked Floyd-Warshall演算法的實作，並探討了不同B值對性能的影響，以及使用shared memory和unroll指示的優化效果。 - **多GPU構想**：提出了將Kernel launch分散給多個GPU device的構想，並進行了相關實驗和性能分析。 - **AMD GPU實驗**：將CUDA程式轉換為HIP程式在AMD GPU上執行，並比較了Nvidia和AMD GPU的性能差異。