正則化儲存格式 (ELL、Hybrid ELL-COO 與 JDS)

重點總覽 (Overview)

本篇承接 14-Sparse-Matrix-Computation/01-Sparse-Matrices-and-Basic-Formats-COO-CSR。CSR 雖省空間又免 atomic,但兩大缺陷:記憶體存取無法 coalesce、control divergence 嚴重。本篇三種「正則化」格式逐一補救。

ELL 格式:Padding + Transposition 達成 Coalescing (Improving Memory Coalescing with the ELL Format)

名稱源自 ELLPACK(解 elliptic boundary value problems 的套件)。從 CSR 出發兩步驟:

1. Padding:找出 nonzero 最多的 row(下例 row 1 有 3 個),其餘 row 在尾端**補 padding(值 0)**到相同長度 → 變成矩形矩陣。colIdx 一併補。
2. Transposition:把矩形矩陣以 column-major(逐欄連續) 排放 = C 語言 row-major 下的轉置。rowPtrs 不再需要(row r 的首元素就在 value[r])。

原矩陣 (4x4)            CSR 分組          padding 成矩形(*=補零)
   c0 c1 c2 c3        row0: 1 7        row0: 1 7 *
r0  1  7  .  .        row1: 5 3 9      row1: 5 3 9   <- 最長 row, max=3
r1  5  .  3  9   -->  row2: 2 8    --> row2: 2 8 *
r2  .  2  8  .        row3: 6          row3: 6 * *
r3  .  .  .  6

       transpose 成 column-major(逐欄攤平)
                 ┌── col0 ──┐┌─ col1 ─┐┌─ col2 ─┐
       value  = [ 1  5  2  6 , 7  3  8  *, *  9  *  * ]
       colIdx = [ 0  0  1  3 , 1  2  2  *, *  3  *  * ]
       index i:  0  1  2  3   4  5  6  7   8  9 10 11
                 └thr0..thr3┘ (第 t 次迭代 = 第 t 欄)

__global__ void spmv_ell_kernel(ELLMatrix ellMatrix, float *x, float *y) {
    unsigned int row = blockIdx.x*blockDim.x + threadIdx.x;       // 02 一 thread 一 row
    if (row < ellMatrix.numRows) {                               // 03 邊界
        float sum = 0.0f;                                        // 04
        for (unsigned int t = 0; t < ellMatrix.nnzPerRow[row]; ++t) { // 05
            unsigned int i = t*ellMatrix.numRows + row;          // 06 column-major 索引
            unsigned int col = ellMatrix.colIdx[i];              // 07
            float value      = ellMatrix.value[i];              // 08
            sum += x[col]*value;                                 // 09
        }
        y[row] += sum;                                          // 11 免 atomic(一 row 一 thread)
    }
}

Coalescing 為何成立:索引 i = t*numRows + row,row 直接綁 threadIdx.x → 同一迭代(同 t),相鄰 thread 的 i 相鄰 → 連續位址 → coalesced。對比 CSR(row-major)第一次迭代 thread 0/1/2/3 讀 value[0],[2],[5],[7](相隔遠、不 coalesced)。

ELL 迭代 t=0:  thr0 thr1 thr2 thr3
               v[0] v[1] v[2] v[3]   <- 相鄰 => COALESCED ✓
CSR 迭代 0  :  thr0 thr1 thr2 thr3
               v[0] v[2] v[5] v[7]   <- 跳躍 => NOT coalesced ✗

Accessibility:ELL 同時具備 CSR(給 row → 取該 row nonzeros)與 COO(給 nonzero → 反推 row/col)的存取性。反推公式:

正向(算 i):   i   = t*numRows + row
反推(算 row): row = i % numRows     (因 row < numRows,故 row%numRows == row)

Hybrid ELL-COO 格式:抑制過量 Padding (Regulating Padding with the Hybrid ELL-COO Format)

問題根源:ELL 的 padding 爆量與 divergence,都來自少數超長 row。對策:轉 ELL 前,把超長 row 尾端的多餘 nonzero 抽出來放進獨立的 COO,剩餘部分做 ELL,最後兩部分結果相加(hybrid method)。

僅用 ELL(row1、row6 超長)→ max=5,大量 padding(22 個)
                 ┌──────── 抽走 row 尾端多餘元素 ────────┐
ELL 部分(max 降到 2,padding 22→3)        COO 部分(收容多餘元素)
   row : [ a b ]                            value : [ ... 被抽走的 nnz ... ]
   row1: [ a b ]  (其餘 3 個 -> COO)         rowIdx: [  1  1  1  6  6 ]
   ...                                       colIdx: [ ...           ]
   row6: [ a b ]  (其餘 2 個 -> COO)
   => 所有 thread 只需迭代 2 次 (divergence 大減)

兩段各自跑既有 kernel:SpMV/ELL(規則的多數元素)+ SpMV/COO(少量長尾元素,用 atomic 累加,見 sibling 筆記)。

JDS 格式:排序消除 Control Divergence (Reducing Control Divergence with the JDS Format)

Jagged Diagonal Storage:不靠 padding,而是依 row 長度排序(例如長→短),排序後矩陣呈鋸齒三角狀,故名。

步驟:(1) 依 row 分組(同 CSR);(2) 依 row 的 nnz 排序,同時用 row 陣列記錄原始 row 索引;(3) column-major 存放;(4) 用 iterPtr 標記每次迭代(每欄)nonzeros 的起點。

排序前 nnz       排序後(長→短)      column-major + iterPtr
 r0: 2           s0(原r1): x x x      iter0(col0): r1 r0 r2 r3   <- 4 個
 r1: 3   sort    s1(原r0): x x        iter1(col1): r1 r0 r2      <- 3 個
 r2: 2  ----->   s2(原r2): x x        iter2(col2): r1            <- 1 個
 r3: 1           s3(原r3): x
 row 陣列 = [1,0,2,3]   value = [r1c0 r0c0 r2c0 r3c0 | r1c1 r0c1 r2c1 | r1c2]
 (保留原始順序)         iterPtr = [0,             4,           7,        8]

 物理視圖(第一次迭代):thr0 thr1 thr2 thr3 讀 value[0..3] 相鄰 => COALESCED ✓

每個 thread 認領一 row,用 iterPtr 找各迭代起點走 dot product(kernel 留作習題)。

為何排序不影響解:重排 row = 重排線性方程式;只要 y 隨之重排即等價於換方程順序,解不變。最後再用 row 陣列把解 x 排回原順序即可。

JDS 變體(分段 ELL):排序後把列切成數段(section),因已排序,同段內各列長度接近;對每段各做一份 ELL,只需 pad 到該段最長列 → padding 遠少於整體單一 ELL。此變體用 section pointer 取代 iterPtr。

本章總結:Data-Dependent 效能與 OP/B (Summary: Data-Dependent Performance & OP/B)

• 正則化的兩大手法:hybrid method(混用兩格式)與 sorting/partitioning(JDS)。兩者都可能重新引入零元素;hybrid 用來緩解「padding 爆量」的病態案例。
• Data-dependent performance:SpMV 的執行效率與頻寬效率取決於輸入矩陣 nonzero 的分佈——這與先前多數 kernel 不同,卻是真實大型應用的常態,也是 SpMV 作為 parallel pattern 的重要性所在。

為何 SpMV 的 FLOPS 很低——關鍵在 arithmetic intensity OP/B:

每個 nonzero 的工作:
  記憶體流量 = value(4B) + colIdx(4B)         = 8 bytes   (稀疏矩陣無 data reuse!)
  運算量     = 1 個 multiply + 1 個 add        = 2 ops
  => OP/B = 2 ops / 8 bytes = 0.25 op/byte

五格式總對照(承 sibling 的 COO/CSR):

考試/面試重點 (Exam / Test Patterns)

Related Notes

• 14-Sparse-Matrix-Computation/01-Sparse-Matrices-and-Basic-Formats-COO-CSR
• 06-Performance-Considerations/01-Memory-Coalescing
• 04-Compute-Architecture-And-Scheduling/02-Warps-SIMD-and-Control-Divergence
• 09-Parallel-Histogram/01-Atomic-Operations-and-Basic-Histogram
• 22-Advanced-Practices-And-Future-Evolution/03-Memory-Bandwidth-and-Compute-Throughput
• 06-Performance-Considerations/04-Optimization-Checklist-and-Bottlenecks