优化策略

warp divergence

ref 实现中的 warp divergence 主要是因为将不同行归入一个 warp 计算, 而不同行的 NNZ 可能有很大差异, 产生 warp divergence. 因此, 只要避免将不同行划入同一 warp 即可. 因此, 令 block.x = 1, 使每个 thread block 至多处理一行数据.

如图中 ref 与 phase_1 对比, 提升有限.

shared memory

在 SpMM 中, 稀疏矩阵的一个元素代表了对于稠密矩阵的一行的访问. 因此可以将稀疏矩阵的一部分缓存在 shared memory 中, 以减少重复从 global memory 中读取稀疏矩阵.

Task 0

Performance

Environment

P100 GPU 最大支持每个 SM 64KB shared memory，但每个 thread block 最多只支持 48KB

Method

使用 实验三 - 高性能计算导论实验文档 (tsinghua.edu.cn) 中的分块方法. 一个 thread block 处理一个矩阵分块. 每个 thread block 所需使用的数据全部拷贝到 shared memory 中. 在 threadIdx 的基础上偏移 i_start, j_start 或 center_block_start 即可将 shared memory 中的坐标映射到 global memory 中的不同矩阵分块.

分析 test_gmem.cu 的性能变化来源

__global__ void stride_copy(float *dst, float *src) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  dst[i * STRIDE] = src[i * STRIDE];
}

Size: 512 x 512 x 512

Naive

No OPT Performance

Performance

Performance

Performance

Implementation

void a_plus_b_intrinsic(float* a, float* b, float* c, int n) {
  for (int i = 0; i < n; i += 8) {
    _mm256_store_ps(
        c + i, _mm256_add_ps(_mm256_load_ps(a + i), _mm256_load_ps(b + i)));
  }
}

Ring Allreduce 算法

首先将每个结点的数据分为 comm_sz 个数据块, 每个数据块大小为 count = n / comm_sz 个 float.

第一阶段共 comm_sz - 1 步. 在第 k 步, 第 my_rank 个进程会将自己的第 (my_rank - k) % comm_sz 对应数据块发送给第 succ = my_rank + 1 个进程并累加. 注意到对于 my_rank 进程, 第 k 步的 Send 与 Recv 使用的数据块不同, 但第 k + 1 步的 Send 依赖于第 k 步的 Recv 得到的数据块. 因此 Send 可以是非阻塞的, 但 Recv 必须是阻塞的, 以确保在第 k + 1 步 Send 前, 第 k 步 Recv 已完成.