稀疏矩陣與基礎儲存格式 (COO 與 CSR)
重點總覽 (Overview)
| 項目 | COO (Coordinate List) | CSR (Compressed Sparse Row) |
|---|---|---|
| 陣列 | rowIdx[], colIdx[], value[] (各 nnz) |
rowPtrs[] (numRows+1), colIdx[], value[] (各 nnz) |
| 平行化單位 | 一個 thread 對應一個 nonzero | 一個 thread 對應一個 row |
| 累加方式 | atomicAdd (多 thread 寫同一 y[row]) |
thread 區域 sum,最後寫入 (無 atomic) |
| 空間效率 (compaction) | 較差 (3 個 nnz 長度陣列) | 較佳 (rowPtrs 只需 numRows+1) |
| 彈性 (flexibility) | 高 (可任意重排、尾端 append) | 低 (插入需位移後續 row) |
| 可存取性 (accessibility) | 易取「某 nonzero 的 row/col」 | 易取「某 row 的所有 nonzero」 |
| 記憶體存取 | Coalesced | 非 coalesced (strided) |
| 負載平衡 | 平衡 (每 thread 1 個工作) | 嚴重 control divergence (row 長度不一) |
| 主要缺點 | 需要 atomic operations | 非 coalesced + control divergence |
SpMV 的核心難題不是運算量,而是 compaction (壓縮去零) 與 regularization (規則化) 之間的權衡:壓得越緊 → 資料越不規則 → 越容易引發非 coalesced 存取、control divergence、load imbalance。各格式的相對效能 高度取決於 nonzero 的分布 (data-dependent)。
背景與 SpMV (Background & SpMV)
- Sparse matrix (稀疏矩陣):絕大多數元素為 0。儲存/處理這些 0 會浪費 記憶體容量、頻寬、時間、能量。
- 來源:科學/工程/金融的線性方程組。每個 row = 一條方程式,每個 column = 一個變數的係數。方程式間 稀疏耦合 (每條方程式只牽涉少數變數)。
- 解線性系統
A·X + Y = 0:- 直接法 (求反矩陣,如 Gaussian elimination) 對大型稀疏矩陣不切實際 → 反矩陣會產生大量 fill-ins (額外 nonzero),比原矩陣還大。
- 迭代法 (iterative):若 A 為 positive-definite,可用 conjugate gradient。每次迭代猜 X、算
A·X + Y、看是否接近 0,再用梯度修正。與第 8 章 Stencil 的迭代解法相關。
- 迭代法最耗時的核心 = SpMV (Sparse Matrix-Vector multiplication and accumulation):計算
A·X + Y,其中 A 稀疏、X 與 Y 為 dense vector。因重要性而有標準函式庫介面SpMV。
SpMV 的算術強度極低 (OP/B ≈ 0.25):每個 nonzero 只做 2 FLOP (1 乘 1 加),卻至少要讀 value(4B)+colIdx(4B)=8B,且稀疏矩陣 沒有 data reuse。因此 SpMV 在 CPU/GPU 上都是 memory-bandwidth-bound,FLOPS 遠低於峰值。(詳細 roofline 討論見 sibling 14-Sparse-Matrix-Computation/02-Regularized-Formats-ELL-Hybrid-JDS)
五大設計考量 (The Five Design Considerations)
| 考量 | 含義 |
|---|---|
| Space efficiency (compaction) | 表示矩陣所需的記憶體容量 |
| Flexibility | 新增/移除 nonzero 的難易度 |
| Accessibility | 格式讓哪一類資料容易被存取 |
| Memory access efficiency | 是否支援高效 (coalesced) 的存取樣式 (regularization 一面) |
| Load balance | 是否能在 thread 間平衡工作 (regularization 另一面) |
範例稀疏矩陣 (Fig. 14.1):
x0 x1 x2 x3
eq0 [ 1 7 . . ]
eq1 [ 5 . 3 9 ]
eq2 [ . 2 8 . ]
eq3 [ . . . 6 ]
( . 代表 0 )
COO 格式與 SpMV (Coordinate List Format)
COO 把每個 nonzero 連同它的 row index 與 column index 各存一份,三個陣列同位置對應。
index : 0 1 2 3 4 5 6 7
value : [1 7 5 3 9 2 8 6]
colIdx: [0 1 0 2 3 1 2 3]
rowIdx: [0 0 1 1 1 2 2 3]
\__r0_/ \___r1___/ \_r2_/ r3
- COO 完全移除所有 0 元素,但引入
colIdx、rowIdx的 儲存負擔 (overhead)。 - 小矩陣時 overhead 可能 > 省下的空間;但 真正稀疏時 (如 1% nonzero),COO 總量 ≈ 原始的 3%。
平行 SpMV/COO:一個 thread 處理一個 nonzero
__global__ void spmv_coo_kernel(COOMatrix cooMatrix, float* x, float* y) {
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x; // 02 該 thread 的 nonzero 索引
if (i < cooMatrix.numNonzeros) { // 03 邊界檢查
unsigned int row = cooMatrix.rowIdx[i]; // 04
unsigned int col = cooMatrix.colIdx[i]; // 05
float value = cooMatrix.value[i]; // 06
atomicAdd(&y[row], x[col] * value); // 07 必須用 atomic!
}
}
Thread → 資料映射 (注意 row 0 由 t0、t1 兩個 thread 共同貢獻):
thread: t0 t1 t2 t3 t4 t5 t6 t7
value : 1 7 5 3 9 2 8 6
row : 0 0 1 1 1 2 2 3
\ / \ | / \ / |
atomicAdd: y[0] y[1] y[2] y[3]
^^^^^^^ t0,t1 同寫 y[0] → race → 需要 atomic
COO 在五大考量下:
| 考量 | COO 表現 |
|---|---|
| Space | 較差 (3 個 nnz 長度陣列),留待與其他格式比較 |
| Flexibility (優) | 可 任意重排 (只要三陣列同步重排,資訊不丟失);處理順序無關 (atomicAdd 結果與順序無關);新增 nonzero = 三陣列尾端 append |
| Accessibility | 易取「某 nonzero → 其 row/col」 → 支援跨 nonzero 平行;不易取「某 row/column 的所有 nonzero」 |
| Memory (優) | 連續 thread 存取三陣列的連續元素 → coalesced |
| Load balance (優) | 每 thread 恰好 1 個 nonzero → 無 control divergence (除邊界) |
因為「順序無關 + 可尾端 append」,COO 是 矩陣初次建構 或 計算中持續修改 時的熱門格式 (例如從未排序檔案讀入)。此彈性在 hybrid ELL-COO 中再次被利用 (見 sibling 筆記)。
SpMV/COO 的主要缺點是 atomic operations:因為同一 row 的多個 nonzero 被分配給不同 thread,必須以原子方式更新同一個 y[row]。要消除 atomic,就得讓 同一 row 的所有 nonzero 由同一 thread 處理 — 但 COO 不提供「以 row 為單位」的存取性。→ 這正是 CSR 要解決的問題。
CSR 格式與 SpMV (Compressed Sparse Row Format)
CSR 與 COO 一樣存 value 與 colIdx,但 nonzero 依 row 分組 (grouped by row),並用 rowPtrs[] 取代 rowIdx[]:rowPtrs[r] = 第 r 個 row 的 nonzero 在 value/colIdx 中的起始 offset。
index : 0 1 2 3 4 5 6 7
value : [1 7 | 5 3 9 | 2 8 | 6]
colIdx: [0 1 | 0 2 3 | 1 2 | 3]
rowPtrs:[0 2 5 7 8]
r0 r1 r2 r3 (end / 不存在的 row4)
rowPtrs[r]起始、rowPtrs[r+1]結束 (exclusive)。例如 row1 =value[2..4]。rowPtrs[numRows](此例 = 8) 多存一個「不存在 row 4」的起點 = 總 nonzero 數,方便界定最後一個 row 的結尾。- row 內若依 colIdx 遞增排序,
value等同「去除所有 0 後的 row-major layout」(排序非必要,kernel 仍正確)。
平行 SpMV/CSR:一個 thread 處理一個 row
__global__ void spmv_csr_kernel(CSRMatrix csrMatrix, float* x, float* y) {
unsigned int row = blockIdx.x * blockDim.x + threadIdx.x; // 02 該 thread 負責的 row
if (row < csrMatrix.numRows) { // 03 邊界檢查
float sum = 0.0f; // 04
for (unsigned int i = csrMatrix.rowPtrs[row]; // 05 起點
i < csrMatrix.rowPtrs[row + 1]; ++i) { // 06 終點 = 下一 row 起點
unsigned int col = csrMatrix.colIdx[i]; // 07
float value = csrMatrix.value[i]; // 08
sum += x[col] * value; // 09 區域累加
}
y[row] += sum; // 11 無 atomic!
}
}
Thread → row 映射 (每 thread 寫入 不同 y[row] → 不需 atomic):
t0 → row0: value[0..1] → y[0]
t1 → row1: value[2..4] → y[1]
t2 → row2: value[5..6] → y[2]
t3 → row3: value[7] → y[3] (各自獨立 → 無 race)
為何 CSR 非 coalesced?(strided access)
連續 thread 在同一迭代讀取的是「各 row 的起點」,彼此 相距很遠:
iter 0: t0→value[0] t1→value[2] t2→value[5] t3→value[7] (跳躍/strided)
iter 1: t0→value[1] t1→value[3] t2→value[6] t3→ (已結束)
iter 2: t0→ done t1→value[4] t2→ done t3→ done ← divergence
對比 COO 同一迭代是 t0→value[0], t1→value[1], ... 連續 → coalesced。
CSR 在五大考量下:
| 考量 | CSR 表現 (相對 COO) |
|---|---|
| Space (優) | 更省:只需 2 個 nnz 長度陣列;rowPtrs 僅 numRows+1,遠小於 COO 的 rowIdx |
| Flexibility (劣) | 新增 nonzero 須插入特定 row → 後續 row 的元素須整體位移、rowPtrs 須遞增 → 昂貴 |
| Accessibility (優) | 易取「某 row 的所有 nonzero」→ 跨 row 平行 → 避開 atomic;不易取某 column 的所有 nonzero |
| Memory (劣) | 連續 thread 存取相距甚遠 → 非 coalesced,浪費頻寬 |
| Load balance (劣) | 迴圈次數 = row 的 nonzero 數,分布隨機 → 多數/全部 warp 都有 control divergence |
「跨 row 平行」在 row 數達數千~數百萬、每 row 數十~數百 nonzero 時很理想 (thread 多、各 thread 工作量足)。但若 row 數太少,無法餵飽所有 GPU thread → 此時 COO 反而能榨出更多平行度 (nonzero 比 row 多)。
COO → CSR 轉換 是絕佳練習:用 histogram (數每 row 的 nonzero 數) + prefix sum / scan (算 rowPtrs) 兩個基礎平行原語即可完成。見 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram 與 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone。
考試/面試重點 (Exam / Test Patterns)
| 情境 / 關鍵字 | 答案 / 技巧 |
|---|---|
| 為何稀疏矩陣不直接求反矩陣解? | 反矩陣產生大量 fill-ins,比原矩陣更大 → 不切實際;改用 iterative (conjugate gradient) |
| SpMV 是什麼運算? | A·X + Y,A 稀疏、X/Y dense 的 matrix-vector multiply + accumulate |
| 為何 SpMV FLOPS 很低? | OP/B ≈ 0.25、無 data reuse → memory-bandwidth-bound |
| COO 平行化單位 / 為何要 atomic? | 一 thread = 一 nonzero;同 row 多 nonzero 寫同一 y[row] → 需 atomicAdd |
| CSR 如何避開 atomic? | 一 thread = 一 row,獨佔輸出 y[row] → 用區域 sum 累加,最後寫一次 |
rowPtrs 長度與末元素含義? |
長度 numRows+1;rowPtrs[numRows] = 總 nonzero 數 (界定最後 row 結尾) |
| 哪個格式 coalesced?哪個不是? | COO coalesced;CSR 非 coalesced (連續 thread 存取 strided) |
| 哪個格式有 control divergence? | CSR (row 長度不一);COO 每 thread 1 工作 → 無 divergence |
| 建構/動態修改矩陣選哪個? | COO (可任意重排、尾端 append) |
| 需逐 row 走訪矩陣選哪個? | CSR (易取某 row 全部 nonzero);COO 不適合 |
| row 數太少時? | 用 COO 榨取更多平行度 (nonzero > row) |
| COO→CSR 怎麼做? | histogram + prefix sum |
| 五大設計考量 | space、flexibility、accessibility、memory access efficiency、load balance |
Related Notes
- 14-Sparse-Matrix-Computation/02-Regularized-Formats-ELL-Hybrid-JDS
- 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
- 06-Performance-Considerations/01-Memory-Coalescing
- 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
- 11-Prefix-Sum-Scan/01-Scan-Foundations-and-Kogge-Stone
- 19-Parallel-Programming-And-Computational-Thinking/03-Problem-Decomposition-Output-vs-Input-Centric