稀疏矩陣與基礎儲存格式 (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
Important

SpMV 的核心難題不是運算量,而是 compaction (壓縮去零)regularization (規則化) 之間的權衡:壓得越緊 → 資料越不規則 → 越容易引發非 coalesced 存取、control divergence、load imbalance。各格式的相對效能 高度取決於 nonzero 的分布 (data-dependent)


背景與 SpMV (Background & SpMV)

Tip

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

平行 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 (除邊界)
Tip

因為「順序無關 + 可尾端 append」,COO 是 矩陣初次建構計算中持續修改 時的熱門格式 (例如從未排序檔案讀入)。此彈性在 hybrid ELL-COO 中再次被利用 (見 sibling 筆記)。

Warning

SpMV/COO 的主要缺點是 atomic operations:因為同一 row 的多個 nonzero 被分配給不同 thread,必須以原子方式更新同一個 y[row]。要消除 atomic,就得讓 同一 row 的所有 nonzero 由同一 thread 處理 — 但 COO 不提供「以 row 為單位」的存取性。→ 這正是 CSR 要解決的問題。


CSR 格式與 SpMV (Compressed Sparse Row Format)

CSR 與 COO 一樣存 valuecolIdx,但 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)

平行 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
Warning

「跨 row 平行」在 row 數達數千~數百萬、每 row 數十~數百 nonzero 時很理想 (thread 多、各 thread 工作量足)。但若 row 數太少,無法餵飽所有 GPU thread → 此時 COO 反而能榨出更多平行度 (nonzero 比 row 多)。

Tip

COO → CSR 轉換 是絕佳練習:用 histogram (數每 row 的 nonzero 數) + prefix sum / scan (算 rowPtrs) 兩個基礎平行原語即可完成。見 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram11-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+1rowPtrs[numRows] = 總 nonzero 數 (界定最後 row 結尾)
哪個格式 coalesced?哪個不是? COO coalescedCSR 非 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