这次回顾ECE408 Lecture 19，这次介绍了并行稀疏方法。

课程主页：

搬运视频：

目标

  • 学习在并行稀疏方法中压缩输入数据以减少内存带宽消耗的关键技术；
    • 更好地利用片上存储器；
    • 更少的字节传输到片上存储器；
    • 更好地利用全局内存；
    • 挑战：保持规律性；

稀疏矩阵

  • 许多现实世界的系统本质上都是稀疏的；
    • 描述为稀疏矩阵的线性系统；
  • 求解稀疏线性系统
    • 基于稀疏矩阵向量乘法的迭代共轭梯度求解器是一种常用方法；
  • PDE系统的解可以表述为表示为稀疏矩阵向量乘法的线性运算；

分析和人工智能中的稀疏矩阵

稀疏矩阵-向量乘法(SpMV)

挑战

  • 与密集矩阵向量乘法相比，SpMV；
    • 不规则/非结构化；
    • 从优化技巧、先前讨论的想法中获益甚微；
  • 实现最佳性能的关键；
    • 减少稀疏性（通过去除零）；
    • 最大化规律性（通过减少分歧和负载不平衡）；
    • 最大化DRAM突发利用率（布局安排）；

简单的并行SpMV

每个线程处理一行。

压缩稀疏行(CSR)格式

记录：

  • 数据值；
  • 列索引；
  • 行指针；
    • 例如下图中表示col_index的第$[0, 2)$个元素是第1行，第$[2, 2)$个元素是第2行，第$[2, 5)$个元素是第3行；第$[5, 7)$个元素是第4行；

CSR数据布局

CSR Kernel设计

并行SpMV/CSR Kernel(CUDA)

__global__ void SpMV_CSR(int num_rows, float *data, int *col_index, int *row_ptr, float *x, float *y) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < num_rows) {
        float dot = 0;
        int row_start = row_ptr[row];
        int row_end = row_ptr[row+1];
        for (int elem = row_start; elem < row_end; elem++)
        	dot += data[elem] * x[col_index[elem]];
        y[row] = dot;
    }
}

CSR内核负载不平衡

线程执行不同次数的迭代：

CSR内核内存发散（未合并访问）

  • 相邻线程访问不相邻的内存位置；
    • 灰色元素在第0次迭代中被所有线程访问；

使用ELL(PACK)格式规则化SpMV

  • 将所有行填充到相同的长度；
    • 如果某几行比其他行长得多，效率低下；
  • 转置（Column Major）以提高DRAM效率；
  • 数据和col_index都被填充/转置；

ELL Kernel设计

并行SpMV/ELL kernel

 __global__ void SpMV_ELL(int num_rows, float *data, int *col_index, int num_elem, float *x, float *y) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < num_rows) {
        float dot = 0;
        for (int i = 0; i < num_elem; i++)
            dot += data[row+i*num_rows]*x[col_index[row+i*num_rows]];
        y[row] = dot;
    }
}

使用ELL进行内存合并

Coordinate(COO)格式

明确列出每个非零元素的列和行索引：

COO允许重新排序元素

实现SpMV/COO的顺序循环

for (int i = 0; i < num_elem; row++)
	y[row_index[i]] += data[i] * x[col_index[i]];

COO内核设计

访问输入矩阵和向量

所有线程都可以访问矩阵数据和col_index以访问向量：

累加到输出向量

每个线程使用其元素的row_index累加到输出，需要原子操作：

Hybrid格式

  • ELL处理典型条目；
  • COO处理异常条目；
    • 以分段减少的方式实施；
  • 在实践中经常在顺序主机代码中实现；