---
tags: [2025_Fall]
---

# 平行程式設計筆記

## Datapath Structure

1. Instruction Fetch: 用 Program Counter 從 Instruction Memory 抓指令
2. Instruction Decode: Decoder 解碼出 OP code, register selections
3. Execute: ALU 做運算，從 register 抓資料
4. Memory Access: Load/Store 從 Data Memory 抓資料或儲存結果。
5. Write Back: 把結果寫回 register 或 Data Memory

## Instruction Level Parallelism (ILP)

Instruction 介紹(跳過)

有些 instruction 之間沒有相依性，可以同時執行，這樣就可以提升效能。對程式設計師來說，ILP 通常是不可見的。

**為何不能實現 N 倍加速？**

* 處理器在同一指令流中找到獨立指令並平行執行的能力 (ILP) 已達到**極限**
* 傳統透過提高時脈頻率和探索 ILP 來提升性能，但受到 **Power Wall** 和製造技術限制
* 因此現在傾向**增加處理單元** (multicore)，但這要求軟體必須寫成平行程式才能獲得性能提升

## Vector Program

比如 Intel AVX2, AVX-512: 用更大的 register 做 vectorization，存好幾個數字，一次進行多個運算，算是一種 **SIMD**，使用 OpenMP 的 `#pragma omp simd` 來告訴 compiler 做 vectorization。

## SIMD on CPUs and GPUs

* CPU: 所有 ALU 要對固定長的 vector 同時做同樣的運算
* GPU: SIMT (single instruction, multiple threads)，執行同樣的 instruction，但每個 thread 可以是不同路徑(if-else)

比如 CPU/GPU 中多個 ALU/thread 同時執行有 if-else 的 code，會因為有些執行 if、有些執行 else，導致 **Divergent execution (分歧執行)**。為確保正確性，所有執行緒必須執行**兩條路徑**，並使用 **Masking (遮罩)** 來停用不需要寫入結果的 ALU。這會造成資源利用效率低下，最壞情況下吞吐量會大幅下降。

## Memory Constraints

記憶體的 **latency (延遲)** 和 **bandwidth (頻寬)** 還是個瓶頸。

**Stalls (停滯)**: 處理器因依賴於前一個 instruction (如 memory access) 而無法執行下一條 instruction 時發生停滯。一個 memory access 可能要幾百、幾千個 cycle。

因此有 Cache，暫存 memory 中常用的資料，L1 cache hit 的話只需要幾個 cycle。

## Hiding memory latencies

**Prefetch (預取)**: 提前把可能會用到的資料從 memory 抓到 cache，主要是 **CPU** 硬體操控，透過分析 memory access pattern 來預測並隱藏 latency。

用 **multi-threading**，有的 thread stall 時就可以執行其他 thread。這是在硬體層面 schedule logical threads (如 SMT/Hyper-Threading)，需要 CPU core 儲存多個 execution contexts。與 OS 作的 process context switch 不同，更快一點。

**Multi-threading/SMT 的優缺點**:

* **優點**: 更有效利用 ALU、hide memory latency、當單一 thread 的 ILP 不足時填補 functional units
* **缺點**: 需要額外儲存 execution contexts、單一 thread 執行時間變長、高度依賴 memory bandwidth、cache 被多個 thread 瓜分而更容易 miss

## Data Parallel Execution Models

### SPMD (Single Program, Multiple Data)

以多個 process 執行同樣程式，但處理**不同的資料集**，之間可能用 shared memory 等方式溝通。
算是 MPMD (Multiple Program, Multiple Data) 的特例。

### Execution Model Comparison

| 模型 | 特點 | 優點 | 缺點 |
|------|------|------|------|
| **MPMD/SPMD** | 多個獨立 threads (Multicore CPUs) | 更通用，支援 thread-level parallelism，可混合 sequential 和 parallel code | 對 data parallelism 效率較差 |
| **SIMD/Vector** | 一個 thread 具有寬執行 datapath (許多 ALU) (x86 SSE/AVX) | 易於為 data parallelism 編程 | Vector programming 較困難 |
| **SIMT** | 多個 lockstep threads (可遵循不同 control flow) (GPUs) | Data parallelism 很簡單 | Divergent execution 會嚴重損害 performance |

## Decomposition of a program

把問題分成小任務，不一定要是靜態的切，程式執行過程也可能產生新的任務。重要的是建立足夠的任務來**保持所有 execution units 都 busy**。

### 簡單的範例

要把一張 N×N 圖片的綠色通道調亮，並計算紅色通道的平均值：

* **Step 1 (調亮)**: 可以輕易平行作，時間約 $N^2/P$
* **Step 2 (計算平均值)**:
* **方法一 (序列合併)**: Step 2 需序列執行 ($N^2$)。總時間約 $N^2/P + N^2$。加速比 $\le 2$
* **方法二 (平行優化)**: Step 2 先平行計算**局部和** ($N^2/P$)，再花 $P$ 的時間**序列合併**。總時間約 $2N^2/P + P$。當 $N >> P$ 時，**加速比接近 $P$**

## Task scheduling and distribution

會想盡量讓所有 processors 的負載平均，讓它們都 busy。使用 assign tasks to threads 的方式。

### Static assignment

先決定好如何分配，好處是簡單。可能基於預測的執行時間。

**適用性**: 當工作成本 (執行時間) 和總工作量是可預測的時候。

### Semi-static assignment

執行時週期性的重新分配任務，在重新分配間是 static 的。
例如粒子模擬，因為會動，根據區域切給不同 processor 的方法需要動態調整。

### Dynamic assignment

比如 thread pool，thread 執行完一個任務後再從任務佇列拿下一個任務來執行。適用於執行時間不可預測的情況。

* **Granularity (任務粒度)**:
* **太細**: 有利於良好的負載平衡，但可能導致高同步成本 (頻繁進入 critical section 存取 shared queue)
* **太粗**: 可減少同步 overhead，但可能犧牲負載平衡
* **理想**: 任務數量應比 processor 多以確保負載平衡，但任務數量應盡可能少以減少管理 overhead 和進入 critical section 的頻率
* **A big task**: 如果突然冒出一個大任務，如果放在最後執行會拖很久，可以考慮獨立用部分 threads。
* **Distributed Task Queues (分佈式任務佇列)**: 每個 worker thread 都有自己的任務佇列，只有當本地佇列為空時，才從其他佇列中 **steal** 任務。這有助於減少同步 overhead，因為不需要在單一 shared queue 上同步。

## Example: 2D grid solver

### 問題敘述

假設 (N+2)×(N+2) 方格中，周圍一圈不會變，而中央 N×N 的點需要在每次運算設為：四周與自己的值的平均（是否是取舊值不重要，因為是要取得最終的平衡點）。

如果用從上而下、從左而右掃描計算的話，可以獨立算每個右上到左下的對角線，但開始與結尾的平行化效果不好、會需要常常同步（下一層很需要上一層的值）。

### 平行化方法

比較好的方法是 **Red-Black Coloring (紅黑著色)**，交錯放置，先更新**所有紅點**（這些點只依賴黑點），完成後**同步**；再更新**所有黑點**（這些點只依賴紅點），再同步。

考量 **Communication overhead** 時，要設計 assignment 給 processes 的方法：

* **Blocked assignment (連續幾列)**: 每個 process 負責連續幾列，與其他 process 的溝通只需要頭尾兩列
* **Interleaved assignment (交錯分配，如 12341234)**: 需要更多的溝通

因此前者較好。

### Message Passing Model

在計算完後，假設所有 processes 都先 send 再 recv(receive) 的話，因為 MPI 的 send, recv 是 **blocking (同步)** 的，呼叫會等待完成，會導致 **deadlock**。

**避免 deadlock 的策略**:

* **同步傳輸**: 可以改為：**奇數 process 先 send 再 recv，偶數 process 先 recv 再 send** (或交替進行)，確保有人負責接收
* **非同步/無阻塞傳輸**: 使用 `MPI_Isend` 和 `MPI_Irecv`，這些呼叫會**立即返回**，允許 processor 在訊息實際傳輸時執行其他工作。但程式設計師必須確保在操作完成前 (通常使用 `MPI_Wait()` 檢查狀態)，**不要修改用於傳輸的 buffer**

另外，關於是否完成，所有 process 傳送 diff 給 root，root 決定是否結束再傳給所有 process。

## Communication Cost Model

總通訊時間 = **Overhead (開銷)** + **Occupancy (佔用時間)** + **Network Delay (網路延遲)**

**Occupancy** 由系統中最慢的組件決定，決定了穩態下的訊息傳輸速率。

### Inherent Communication (內在溝通)

演算法本身**根本上必需**的通訊。可以透過**改變演算法**或**優化核心間的資料分配**來減少。

**例子**: Grid Solver 中，採用 **Blocked assignment** (區塊分配) 而非 Interleaved assignment，可大幅降低溝通與計算的比率。算術強度 $\propto N/P$。

> 核心間溝通也是屬於 Inherent Communication。

**Communication-to-Computation Ratio**:

$$\frac{\text{Amount of communication (e.g. bytes)}}{\text{Amount of computation (e.g. instructions)}}$$

Arithmetic intensity = 1 / Communication-to-Computation Ratio，越高越好

### Artifactual Communication (人為溝通)

由於**系統實作細節**造成的額外通訊：

* 採用 Write through cache 時頻繁寫入會需要一直儲存到記憶體，改用 Write back 就不用
* Cache Line Granularity，如 cache line 大小為 64 bytes，但實際需要用 65 bytes，會導致多傳輸 63 bytes 的無用資料

#### Cache Miss 類型

* **Cold miss (Compulsory miss)**: 第一次存取，sequential 程式無法避免，但平行程式可以 hide latency
* **Capacity miss**: Working set 大於 cache
* **Conflict miss**: Set-associative cache 的映射衝突
* **Communication miss (Coherent miss)**: 比如 P1 使用 `a[0:3]`，P2 使用 `a[4:7]`，兩者都在同一 cache line，P1 寫完後 P2 讀就會 miss。

#### Locality Optimization

透過 **Loop Fusion (迴圈融合)**，讓多個運算步驟重複使用已載入的資料，從而**提高 Arithmetic Intensity**，減少對 memory 的存取次數。

#### Contention

資源競爭，比如 critical section (dynamic assignment 存取 shared task queue)、GPU 存取共享記憶體的同一個 memory bank。

可透過增加 ginarity 或是 distributed task queue 來減少。

## CUDA Programming Model

2007 Tesla 架構後才有 CUDA，之前只能用來處理 Vertex/Fragment shaders 等 graphics 任務。

* **Host (CPU)**: 執行 sequential code，呼叫 **Kernel** 函數
* **Device (GPU)**: 執行 parallel kernel 函數

### Execution Hierarchy

Kernel 函數被組織成三個層次：

* **Grid**: 整個 kernel，執行在整個 GPU 上
* **Block**: 一組 threads，執行在單一 **SM (Streaming Multiprocessor)** 上，且**不會遷移**。視資源使用情況而定，有可能一個 SM 上同時執行多個 blocks。
* **Thread**: 最小執行單位，執行在 **CUDA Core (SP, Streaming Processor)** 上

### Warp

Threads 被分組為 **Warps** (通常 32 個 threads)。Warp 是 **SIMT (Single Instruction, Multiple Threads)** 架構下的基本執行單位。

如果 block 中的 thread 數量是 **warp 大小的倍數**，資源使用效率會更高。

### CUDA Function Qualifiers

* **`__global__`**: Host 呼叫，在 Device 執行 (kernel 函數)
* **`__device__`**: 僅在 Device 上呼叫和執行 (device 函數)
* **`__host__`**: 在 Host 上執行 (預設)

### Occupancy (佔用率)

$$\frac{\text{Active Warps}}{\text{Maximum Allowable Active Warps}}$$

衡量 GPU 計算資源的利用率。**高 occupancy** 表示有足夠多的 active threads 來**隱藏 memory latency**。

每個 block 使用的資源過多會限制 occupancy：

* **Register 使用量**: $\text{Active threads} \le \frac{\text{Total Registers per SM}}{\text{Registers per Thread}}$
* **Shared memory 使用量**: $\text{Active threads} \le \frac{\text{Total Shared memory per SM}}{\text{Shared memory per Thread}}$，shared memory 還是各 block 共用，只是通常會隨著 thread 數增加而增加。
* **Block size**: $\text{Active threads} = \text{Threads per block}\times \text{Blocks per SM} \le \text{Max Threads per SM}$

**低 occupancy** 可能導致 **GPU 資源利用不足**。

### CUDA Memory Hierarchy

GPU memory 分為三種 address spaces：

* **Private**: 每個 thread 的 registers/local variables
* **Shared**: Block 內所有 threads 共享，速度比 global memory 快約 **100 倍**
* **Global**: 所有 threads 可存取，生命週期與應用程式相同

### Memory Coalescing (記憶體合併)

為了提高 global memory 存取效率，Warp 中 threads 的請求會被**分組 (grouped)**，並以 128 bytes 等對齊單位進行存取。

**最佳情況**: Warp 中相鄰 threads 存取連續的 memory addresses。

### Shared Memory Bank Conflicts

Shared memory 劃分為多個 Banks。如果同一個 Warp 中的多個 threads 同時存取**相同的 Bank** (但不同地址)，會發生衝突，導致請求被序列化處理。

**例外**: 多個 threads 讀取同一地址不會有衝突 (broadcast)。

SM 與 L1 cache 是用同個空間，可以用 `cudaFuncSetCacheConÞg` 來設定 L1 與 shared memory 的比例。

### Synchronization

* **`__syncthreads()`**: 在 **block 內**作為 **barrier**，確保 block 內所有 threads 到達此點
* **Atomic operations**: 如 `atomicAdd`，用於保護 critical section，確保能存取變數的 threads 中只有一個 thread 執行操作
* Host 與 Device 之間的同步: `cudaDeviceSynchronize()`、kernel 結束時。

## Unified Memory (UM)

一種更簡單的 memory 管理方式，CPU 和 GPU 可以存取相同的變數，使用 `cudaMallocManaged`。

CPU 或 GPU 存取時觸發 **page fault**，自動將 memory pages 遷移到存取它的 processor 上

**適用**: 存取模式**稀疏或不可預測**的情況。

### Asynchronous Prefetch

使用 `cudaMemPrefetchAsync` 可以在資料實際被需要之前，**非同步地**將 UM 資料塊遷移到 CPU 或 GPU memory 中，以**減少 page fault 和遷移 overhead**。

### Asynchronous Transfers

允許在 **GPU 執行 kernel 時，同時在 PCIe bus 上傳輸資料**。透過使用 **CUDA Streams** 和 `cudaMemcpyAsync` 實現，可有效利用通訊通道 (pipelining)。

## CUDA Streams

GPU 上的**操作序列**。不同 Streams 之間的操作可以**並行執行 (Concurrent)**。

**Default Stream** 比較特殊，會阻塞所有其他 Streams。

### Host/Device Synchronization

* **`cudaDeviceSynchronize()`**: CPU 等待**所有** GPU 操作完成
* **`cudaStreamSynchronize(stream)`**: 只等待**特定 Stream** 的操作完成

## GPU Dynamic Scheduling

**Hardware scheduler** 負責將 **Thread Blocks 動態地映射到 SMs**。系統假設 blocks 間沒有依賴性，可以**以任何順序執行**。

**NVIDIA 4090 SM**: 具有四個 **Warp schedulers**，可同時運行來自四個 Warps 的 instructions。每個 SM 可同時維護多達 **48 個 Warp 的 execution contexts**。

### Block Dependency 問題

雖然 blocks 可按**任何順序**執行，但若不同 blocks 存取 global memory 的**相同變數** (例如在 Histogram 範例中使用 `atomicAdd`)，則必須確保執行順序不影響結果。

**程式設計師必須非常小心**，因為系統不保證 blocks 的執行順序。

## Barnes-Hut Algorithm (BH)

用於加速 N-body 模擬 (如星系演化) 的演算法。它將計算複雜度從 $O(N^2)$ 降低到 **$O(N \log N)$**，通過將遠處的粒子群組視為單一質量中心來近似計算引力。

### BH 挑戰與分解

* 每個粒子的**工作量是非均勻的**
* 粒子會移動，導致工作成本和通訊模式**隨時間變化**
* BH 使用 **Quad-tree (四元樹)** 來表示粒子空間分佈

### BH 負載平衡策略

採用 **Semi-static assignment (半靜態分配)**。根據每個粒子**先前計算的交互成本**，**週期性地重新分配**工作給 processors。這利用了粒子空間分佈緩慢演變的特性。

## Performance Evaluation

### Basic Metrics

* **Speedup**: $\text{Speedup} = \frac{\text{Execution time of sequential program}}{\text{Execution time on } P \text{ processors}}$
* **Efficiency**: Performance per unit resource (e.g., per watt, per dollar, per chip area).

> **Pitfall (常見陷阱)**:
> 算 Speedup 時，Baseline (分母) 必須是 **Fastest sequential algorithm on one core**，而不是把原本的 Parallel program 跑在一個 core 上。
> 因為 Parallel program 會有 synchronization/communication overhead，跑在單核上通常比純 sequential 版本慢，這樣算出來的 Speedup 會灌水 (unfair)。

### Scaling Constraints

設計系統或評估效能時，受限於問題規模與硬體資源的關係：

1. **Problem Constrained Scaling (PC)**
* **定義**: 固定 Problem size，目標是縮短執行時間 (Solve faster)。
* **場景**: 大部分課堂作業、批次處理。
* **問題**: 如果機器變大 ($P$ 增加) 但問題太小，Parallelism overhead 會蓋過好處。

2. **Time Constrained Scaling (TC)**
* **定義**: 固定執行時間，目標是做更多的事 (Do more work)。
* **場景**:
* **Real-time Gaming**: 每一幀必須在 16ms 內算完，顯卡變強通常是為了讓畫質更好 (更多 work)，而不是讓 FPS 衝到幾千。
* **High-frequency Trading**: 必須在固定 latency 內完成決策。

3. **Memory Constrained Scaling (MC)**
* **定義**: 固定每個 Processor 的 Memory usage，目標是跑更大的 Problem size (Run largest problem possible)。
* **場景**: 科學運算 (N-body simulation)、大型 ML 模型訓練。
* 重點是把記憶體塞滿，避免 Thrashing (Swap to disk)。

## Roofline Model

用來視覺化與分析系統效能瓶頸的模型。

* **X 軸**: **Arithmetic Intensity** (FLOPS/Byte)
* 每從 Memory 搬運 1 Byte 資料進來，能做幾次運算。
* **Y 軸**: **Performance** (GFLOPS)
* **公式**: $P = \min(\pi, \beta \times I)$
* $\pi$: Peak Performance (硬體運算極限)
* $\beta$: Peak Bandwidth (記憶體頻寬極限)
* $I$: Arithmetic Intensity

### 兩個主要區域

1. **Memory Bound (Diagonal region)**
* 圖表的左半邊斜線區。
* $I \le \pi / \beta$
* 效能被 **Memory Bandwidth** 卡住。
* **Optimization**:
* Software Prefetching (隱藏 latency)
* 改善 Data Locality (如 loop fusion，增加 Cache hit rate，提升 effective $I$)
* NUMA optimization
2. **Compute Bound (Horizontal region)**
* 圖表的右半邊平坦區。
* $I > \pi / \beta$
* 效能被 **Compute capability** 卡住 (已達硬體極限)。
* **Optimization**: Parallelism
* ILP (Instruction Level Parallelism)
* SIMD/Vectorization
* FMA (Fused Multiply-Add)

* **Ridge Point**: 斜線轉平線的轉折點，代表達到 Peak performance 所需的最低 Arithmetic Intensity。

## Interconnection Networks

### 簡介

互連網路 (Interconnection Networks)，也稱為**片上網路 (Networks-on-Chips, NoCs)**，是連接處理器節點的基礎設施。

**為什麼 NoC 很重要？**

* **系統可擴展性 (Scalability)**: 匯流排 (Bus) 式互連難以良好擴展（單一匯流排難以共享），而 NoC 易於擴展到大型系統，並可輕鬆添加新組件
* **系統性能**: 實現組件之間的**快速通訊**和**可管理延遲**
* **能源效率**: 相比匯流排，NoC 通常具有更低的能耗

### 核心術語

| 術語 | 定義 |
| :--- | :--- |
| **Node (節點)** | 連接到互連網路的組件，例如處理器、快取、記憶體等 |
| **Links (鏈路)** | 用於傳輸訊號的連接結構（例如電線） |
| **Switch/Router (交換器/路由器)** | 連接固定數量的輸入與輸出鏈路 |
| **Network interface** | 連接節點到網路的介面 |
| **Topology (拓撲)** | 交換器/路由器透過鏈路連接的方式；影響路由、吞吐量、延遲、複雜性和成本 |
| **Routing (路由)** | 訊息從源頭到目的地傳輸的方式；可以是**靜態路由**（預定義策略）或**動態路由**（基於流量負載調整） |
| **Flow control (流量控制)** | 管理網路流量和緩衝區空間的演算法 |

### 拓撲的屬性

| 屬性 | 定義 |
| :--- | :--- |
| **Routing distance (路由距離/跳數)** | 兩節點之間沿著路徑的鏈路數量 |
| **Diameter (直徑)** | 最大的路由距離 |
| **Average distance (平均距離)** | 所有有效路徑的平均路由距離 |
| **Bisection bandwidth (對分頻寬)** | 將網路切成兩半，計算所有被切斷鏈路的總頻寬，是衡量網路的常見指標 |
| **Blocking vs. Non-blocking (阻塞 vs. 非阻塞)** | 如果任意一對節點之間的通訊**總是可能**，則網路為非阻塞 (Non-blocking) |

### 網路類型

* **Direct Network (直接網路)**: 每個節點都有自己的交換器/路由器，節點**位於網路內部**
* **Indirect Network (間接網路)**: 節點**位於網路外部**

### 常見拓撲比較

| 拓撲 (Topology) | 類型 | 成本 (Cost) | 延遲 (Latency) | 特點與應用 |
| :--- | :--- | :--- | :--- | :--- |
| **Bus (匯流排)** | 直接 | 低 (小系統) | N/A | 爭用 (Contention) 高、頻寬有限、擴展性差 |
| **Crossbar (交叉開關)** | 間接 | $O(N^2)$ | $O(1)$ | 高頻寬、低延遲；難以擴展，成本高 |
| **Ring (環狀)** | 直接 | $O(N)$ | 高 | 簡單；對分頻寬不變，擴展性問題 |
| **Mesh (網格)** | 直接 | $O(N)$ | $O(\sqrt{N})$ | **優異的可擴展性**，易於晶片佈局 |
| **Torus (環面)** | 直接 | $O(N)$ | N/A | 結合 Mesh 和 Ring，有更多鏈路和路由選項；佈局複雜 |
| **Tree (樹狀)** | 間接 | N/A | $O(\log N)$ | 分層拓撲，用 **Fat tree** 緩解根節點頻寬瓶頸 |
| **Multi-Stage Network** | 間接 | $O(N\log N)$ | $O(\log N)$ | 阻塞網路 |

### 通訊粒度

* **Message (訊息)**: 節點間通訊的資訊，可由多個 Packet 組成
* **Packet (封包)**: 網路傳輸的單位，可由多個 Flit 組成
* **Flit (Flow Control Digit)**: 封包分解成的更小單位，是**鏈路傳輸的基本單元**，也是**緩衝和路由的最小粒度**
* 封包包含：**Header flit** (含路由資訊)、**Payload/Body flit** (主體)、**Tail flit** (含錯誤校正碼)

### 交換方法

* **Circuit Switching (電路交換)**: 在傳輸前建立發送方和接收方之間的**完整路徑**
* **優點**: 傳輸期間**無競爭**，高頻寬傳輸，不需緩衝
* **缺點**: 需要設置和釋放鏈路的開銷，可能導致**鏈路利用率低**
* **Packet Switching (封包交換)**: 針對**每個封包**做路由決策
* **優點**: 無需設置開銷，鏈路閒置時可被立即分配

### 流量控制方法

流量控制決定了如何管理網路中的流量和緩衝空間。

#### Flit (FLow control digIT)

NoC 使用的網路架構中，當一個封包要在網路中傳輸時，會被切成多個固定數量的 Flits：

* **Header flit**: 包含目的地位址和路由資訊，決定封包要往哪裡走
* **Body/Payload flits**: 實際要傳輸的資料內容
* **Tail flit**: 標記封包結束，可能包含錯誤檢查碼 (checksum)

這樣的設計讓 pipeline 的方式傳輸變得可能，進一步壓低延遲。

#### Store and Forward (儲存與轉發)

* **流程控制單元**: 整個**封包**
* **機制**: 封包必須在路由器中**完整複製**，然後才能轉發到下一節點
* **缺點**: **高延遲**。延遲隨著跳數 (number of hops) 線性增加

#### Cut-Through (直通式)

* **機制**: 交換器/路由器在**接收到封包 Header 後**，立即開始將資料轉發到下一鏈路
* **優點**: 顯著**降低延遲**
* **缺點**: 仍需要**緩衝整個封包**。如果輸出鏈路阻塞 (blocked)，整個傳輸會**退化成 Store-and-Forward**

#### Wormhole (蟲孔式)

* **機制**: Body flits 跟隨 Header flit。如果 Header 停止，其餘 flits 也停止，形成**完全管線化 (pipelined)** 傳輸
* **優點**: **所需緩衝區數量較少**
* **總延遲**: 主要取決於單個 Flit 的傳輸時間和 Flit 總數
* **問題**: **Head of Line Blocking (隊頭阻塞)**: 當一個封包的 Header 被阻塞時，它佔用了路由器中的緩衝區，阻止了其他**可以通過**的封包。也有可能導致 deadlock。

#### Virtual Channel (虛擬通道, VC) 流量控制

* **目的**: 允許**不同的封包共享一個物理通道**，主要用於解決**隊頭阻塞 (Head of Line Blocking)**
* **機制**: 路由器的輸入緩衝區被劃分為多個 **虛擬通道 (VCs)**。如果一個 VC 中的 Flits 被阻塞，其他 VC 仍然可以使用共享鏈路傳輸 Flits
* **裁決**: 使用 **Virtual Channel Arbiter** 來決定哪個 VC 可以使用共享鏈路
* **優勢**:
* **避免死結 (Deadlock avoidance)**: 可用於打破資源的**循環依賴**
* **流量優先級 (Prioritization)**: 某些 VC 可以設定更高的優先級，提供服務品質 (QoS) 保證

一個 Link 一次還是只能走一個 Flit，VC 只是解決了 Head-of-Line Blocking、deadlock。

### 現代路由器管線

現代 NoC 路由器由多個組件構成，包括輸入緩衝區（劃分為 VCs）、**Crossbar Switch**、**VC Allocator** 和 **Switch Allocator**。

標準的**五級路由器管線 (Five-stage router pipeline)** 階段：

1. **BW** (Buffer Write, 緩衝區寫入)
2. **RC** (Route Calculation, 路由計算)
3. **VA** (VC Allocation, VC 分配)
4. **SA** (Switch Allocation, 交換器分配)
5. **ST** (Switch Traversal, 交換器穿越)