Stanford CS149 PARALLEL COMPUTING Notes (Fall 2023) / CMU15-418

Stanford CS149 Parallel Computing: 斯坦福一门并行计算的课，最近不知道该做些什么，感觉要是继续颓下去就废了。打算学一学这门课，在这里记一下看 Slides 做的笔记

课程主页：Fall 2023

网上似乎没有 CS149 的视频，这门课前身是 CMU 的 15-418，可以看这门课的视频

Lecture 1: Why Parallelism? Why Efficiency?

2024/3/25

目标：并行处理加速程序计算

优化：减少线程之间通信 / 同步，线程过多将严重影响性能，平均分配任务

Fast != Efficient，考虑硬件消耗性价比

利用硬件中针对专门指令的处理单元

指令级并行 ILP

内存定义：一段连续的地址空间，通过地址定位其中的每一个字节，访问内存通常使得 CPU 等待影响性能

超标量 CPU：每个核心有多个逻辑单元（ALU），可以一个周期同时执行多条指令（若指令间相互独立）

Cache：将部分内存缓存到 CPU 自带的高速缓存中，若缓存命中则不需要访问内存，通常已 Cache Line 为单位进行缓存。当访问连续内存或多次访问同一内存时，Cache 命中率高

CPU 有多级缓存，每个 Processor 有单独的 L1, L2 Cache，多个 Processor 共享 L3 Cache

Cache 大致速率：L1 (4 cycles) < L2 (12 cycles) < L3 (40 cycles) < DRAM (250 cycles)

Lecture 2: A Modern Multi-Core Processor

2024/3/26

一个例子：给一个数组 \(x[i]\) 和简单函数 \(f(x)\)，计算 \(y[i] = f(x[i])\)

1. 将数组分成 \(N\) 个部分，每个部分分配给一个线程
2. 如果函数 \(f(x)\) 中的指令完全相同，只是传入的数据不同，可以利用 SIMD 指令集，将读取到的一条指令传给多个 ALU，在单个核心上同时计算多个数据，例如使用 AVX 指令集同时计算 8 个 32 位数据（SIMD 和循环展开是不一样的）

但如果代码中有分支，SIMD 就会失效，一次只能执行部分指令，所以 SIMD 要求程序一致性

到此有三种并行：Superscalar（硬件实现，多个取解码器、ALU）, SIMD（编译器、硬件实现）, Multi-Core

Cache LRU 替换策略，Cache 预读

提高单核吞吐量（Interleaved multi-threading）：当一个线程等待时切换到另一个线程执行。代价：上下文切换，适当的线程数量可以提高 CPU 利用率，但线程过多会导致频繁上下文切换

同步多线程（Simultaneous multi-threading , SMT）：每一个周期处理器选择多个线程的指令同时执行，如 Intel 的超线程技术

Lecture 3: Multi-core Arch Part II + ISPC Programming Abstractions

延迟（完成一整个任务的费时）、带宽、吞吐量（平均任务处理能力）的区别

在大量读取内存的情况下内存带宽成为瓶颈

流水线：将一个任务分成多个阶段，在前一个任务完成一个阶段后同时开始下一个任务的这个阶段，提高吞吐量。通常一个周期完成一条指令指的就是流水线下的吞吐量而非延迟

Intel SPMD Program Compiler (ISPC): 一个 C 语言变体的编译器，此语言针对 SPMD (Single Program Multiple Data) 进行扩展（利用 SIMD 指令）。通过使用 ISPC 可以自动利用 SIMD 指令集

2024/3/27

表层 Abstraction 与底层 Implementation

调用 ISPC 函数时生成多个程序实例，每个实例同时计算。ISPC 一些关键字：programCount 实例数量，programIndex 当前实例编号，uniform 一致变量（编译器自动优化，所有实例使用同一份变量），foreach 循环

使用泰勒展开计算 \(n\) 个 \(\sin(x)\)：

export void ispc_sinx(
  uniform int N,
  uniform int terms,
  uniform float* x,
  uniform float* result)
{
  // assumes N % programCount = 0
  for (uniform int i = 0; i < N; i += programCount) {
    int idx = i + programIndex;
    float value = x[idx];
    float numer = x[idx] * x[idx] * x[idx];
    uniform int denom = 6; // 3!
    uniform int sign = -1;
    for (uniform int j = 1; j <= terms; j++) {
      value += sign * numer / denom
      numer *= x[idx] * x[idx];
      denom *= (2*j+2) * (2*j+3);
      sign *= -1;
    }
    result[idx] = value;
  }
}

分组求和（不能直接 float sum 或者 uniform float sum）：

export uniform float sum_array(
  uniform int N,
  uniform float* x)
{
  uniform float sum;
  float partial = 0.0f;
  foreach (i = 0 ... N) {
    partial += x[i];
  }
  // reduce_add() is part of ISPC’s cross
  // program instance standard library
  sum = reduce_add(partial);
  return sum;
}

Lecture 4: Parallel Programming Basics

编写并行程序的基本步骤：Decomposition, Assignment, Orchestration, Mapping

加速效果 \(\textup{Speedup} = \frac{Time(Sequential)}{Time(Parallel)}\)

程序中无法并行部分占比 \(S\)，则 \(\textup{Speedup} \le \frac{1}{S}\)，\(P\) 个核心下 \(\textup{Speedup}\le \frac{1}{S + \frac{1 - S}{P}}\)

通常分解任务由程序员完成，分配任务由编译器完成

一个例子：使用 G-S 迭代法解 PDE（死去的线性代数与数值分析正在攻击我）。部分代码：

while (...) {
  for (int i = 1; i <= n; i++) {
    for (int j = 1; j <= n; j++) {
      prev = A[i, j];
      A[i, j] = 0.2 * (A[i, j] + A[i, j-1] + A[i-1, j] + A[i, j+1] + A[i+1, j]);
      diff += fabs(A[i, j] - prev);
    }
  }
  ...
}

这种情况下每一次迭代中 A[i, j] 依赖于 A[i - 1, j] 和 A[i, j - 1]。一种方法是按从左上往右下的顺序，这样虽然可以并行加速但比较麻烦，效果可能不明显。一种优化是使用 Red-Black Ordering，将矩阵交叉染色分成红黑两部分，分别计算。这种情况下根据局部性原理将连续的数据分配给一个线程，可以减少跨线程内存访问

老生常谈：多线程共享内存的同步问题、原子操作、Lock、Barrier

上述迭代算法中一个 Barrier 分析的例子：

while (!done) {
  float myDiff = 0.f;
  diff = 0.f;
  barrier(myBarrier, NUM_PROCESSORS);
  for (cells) {
    float prev = A[i,j];
    A[i,j] = ...;
    myDiff += abs(A[i,j] - prev);
  }
  lock(myLock);
  diff += myDiff;
  unlock(myLock);
  barrier(myBarrier, NUM_PROCESSORS);
  if (diff/(n*n) < TOLERANCE) done = true;
  barrier(myBarrier, NUM_PROCESSORS);
}

为什么有三个 Barrier？第一个：防止计算全局变量 diff 时有某个 Instance 调用开头的 diff = 0，第三个 Barrier 同理，第二个 Barrier 则是希望所有 Instance 得到相同结果，不会在计算完成时有某个 Instance 继续迭代

优化：上述 Barrier 基本都是由于 diff 变量，通过在每次迭代中使用不同的 diff 变量可以减少至一个 Barrier

（课件中是一个 diff[3]，但我感觉两个 diff 就够用了？暂时不是很清楚，视频中也没详细讲）

---

完成了 CS149 Programming Assignment 1

---

Lecture 5: Performance Optimization I: Work Distribution and Scheduling

2024/3/30

任务分配的不同方式：

• Static Assignment: 任务计算成本已知，效率最高
• Semi-Static Assignment: 任务计算成本不完全已知，但可以动态估计，运行时进行少量调整
• Dynamic Assignment: 任务计算成本不可预测，运行时随时动态调整，分割出的任务数量越多越容易平均分配，但也会增加额外开销

如何将任务平均分配给线程？

• 将大型任务分割成小任务：额外开销，且可能无法分割
• 先分配大型任务：需要运行时间可计算

动态分配：Task Queue

可以每个核独享一个任务队列减少核之间的竞争，当任务队列为空时可以从其他核的队列中获取任务

任务间有依赖：需要完成前置任务才能取出队列中的任务

Cilk Plus: C++ 的扩展语言，cilk_spawn + cilk_sync 实现任务并行（Fork-Join 范式）。底层实现：线程池，控制同时运行的线程数量，每个线程有自己的任务队列，队列为空时去其他线程队列获取任务

cilk_spawn 的两种实现（Cilk Plus 为后者）：生成线程后继续执行当前任务（Continuation First）、生成线程后挂起当前任务并切换到新线程执行（Child First）

线程间任务偷取的实现：任务队列为双端队列（Lock-Free Deque），线程自己生成 / 获取任务时从队头插入 / 取出，从其他线程偷取任务时从队尾取出

从其他线程偷取任务时可以一次偷取大量任务减少次数，在 Child First 的情况下双端队列可以保证线程间的局部性

cilk_sync 的实现：若没有线程间任务的偷取则无需特殊操作。否则将 cilk_sync 对应的任务添加同一个标识符，维护标识符对应的生成任务数量与完成任务数量，当两者相等即所有任务完成

Lecture 6: Performance Optimization II: Locality, Communication, and Contention

CPU 内部结构：每个核心一个 L1, L2 Cache，多个核心共享 L3 Cache。

NUMA 非一致性内存访问架构：内存被划分为 NUMA 节点，不同核心访问不同节点的延迟和带宽都可能不同

Message Passing Model：每个线程独享自己的内存空间，只能通过消息传递进行通信

以 G-S 迭代解 PDE 为例，每个数的计算依赖于周围 4 个数，因此中间部分矩阵对应的线程需要从其它线程额外获取上下两行（Ghost rows）。最后每个计算线程需要把结果同步给主线程：

if (tid != 0)
  send(&localA[1, 0], sizeof(float) * (N + 2), tid - 1, MSG_ID_ROW);
if (tid != get_num_threads() - 1)
  send(&localA[rows_per_thread, 0], sizeof(float) * (N + 2), tid + 1, MSG_ID_ROW);

if (tid != 0)
  recv(&localA[0, 0], sizeof(float) * (N + 2), tid - 1, MSG_ID_ROW);
if (tid != get_num_threads() - 1)
  recv(&localA[rows_per_thread + 1, 0], sizeof(float) * (N + 2), tid + 1, MSG_ID_ROW);

但如果使用 Synchronous (blocking) send / recv（send 会阻塞线程直到收到对方的 ack，recv 会阻塞线程直到接收完毕），上述代码就有可能死锁，需要修改消息传递的顺序避免死锁。同时也有异步版本的 Non-blocking asynchronous send / recv

通信无处不在：CPU - L1 - L2（或者其它核的 L2）- L3 - Memory（或者其它电脑的内存）- 硬盘 - 网络 ......，局部性原理

2024/3/31

运算强度 Arithmetic intensity = 运算量 / 通信量

Inherent communication 必要通信（算法原理所必须的通信，如 Ghost rows）/ Artifactual communication 人为通信（具体实现细节中的通信，如 Cache miss）

减少必要通信：已 G-S 迭代为例，\(N\times N\) 的网格，\(P\) 个处理器，若横向分成 \(P\) 块，则运算强度为 \(O(\frac{N}{P})\)；若横竖分为 \(\sqrt{P}\times \sqrt{P}\) 块，则运算强度为 \(O({\frac{N}{\sqrt{P}}})\)

减少人为通信：利用局部性原理减少 Cache miss

先尝试最简单的算法，再通过性能分析找出瓶颈进行优化：手动 / 利用工具。有时多核 Speedup 也能出现超线性（单核下问题规模大导致 Cache miss），因此针对不同情况要具体分析

Lecture 7: GPU architecture and CUDA Programming

图形处理单元 GPU：用于高性能并行计算，有单独的内存

CUDA 编程语言：类 C 语言，运行在 GPU 上。CUDA Thread ID 可以达到三维（CUDA Thread 仅抽象概念，实际底层与线程不同），一个 Block 内有多个 Thread

主要分为两部分，Host Code：和普通 C / C++ 一样在 CPU 上串行执行，Device Code：在 GPU 上并行执行 (SPMD)

有些类似不同层次的 Cache，CUDA 中有三种内存类型：Per-Thread Memory、Per-Block Memory、Device Global Memory

CUDA 同步相关：Barrier __syncthreads()、原子操作、Host / Device 同步（CUDA kernel 返回时的隐式 barrier）

Nvidia V100 SM (Streaming Multi-processor) Sub-core：一个 Block 中每 32 个 CUDA 线程组成一个 Warp（类似 ISPC 中的 Gang），每个 Sub-core 最多 16 个 Warp。如果 Warp 中的 CUDA 线程指令相同则以 SIMD 方式执行（硬件动态检查）

V100 结构：每个 Sub-core 中有一个取指单元、多个功能单元，16 个 Warp（每个 Warp 中有 32 个线程的寄存器）；4 个 Sub-Core 组成一个 SM unit，有共享内存和 L1 Cache；80 个 SM unit 共享 L2 Cache

Lecture 8: Data-Parallel Thinking

2024/4/1

Data-parallel model: Map, Fold, Scan

以求前缀和为例

• 直接法复杂度 \(O(N)\)，并行度 \(O(n)\)
• 倍增法总复杂度 \(O(N\log N)\)，并行度 \(O(\log N)\)
• 树状数组法复杂度 \(O(N)\)，并行度 \(O(\log N)\)

对不同规模数据使用不同规模算法，大型数据可以分治以保持局部性

Segmented scan -> Sparse matrix multiplication

Gather-Scatter, GroupBy, Filter, Sort

---

完成了 CS149 Programming Assignment 2

---

Lecture 9: Distributed Data-Parallel Computing Using Spark

2024/4/5

在分布式集群上的并行计算

Message passing model，分布式文件系统

MapReduce 模型

• Map：将数据映射到中间结果
• Reduce：将中间结果合并

调度器职责：利用数据局部性分配任务，处理节点故障、慢节点

MapReduce 局限：只支持简单模型，每次迭代都需要读写磁盘（目前大多数场景下内存足够）

Apache Spark：基于内存的分布式计算框架，适合中间数据多次使用的场景（迭代算法、数据查询）

保证容错性：复制多个计算任务 / 检查点 + 回滚 / 维护日志

RDD (Resilient Distributed Dataset)：弹性数据集，只读的有序集合，通过对持久性存储数据或现有 RDD 进行确定性转换而创建，对其操作返回数据

在 RDD 上链式操作会创建大量中间数据占用内存，优化：loop fusion（合并操作），tiling（分块操作）

Narrow dependency（一个 RDD 只被一个子 RDD 依赖，便于合并）/ Wide dependency（多个依赖，通信成本 / 容错成本）

RDD 实现：Lineage（记录操作历史，Log 开销低，可以从历史数据恢复）

Lecture 10: Efficiently Evaluating DNNs on GPUs

DNN (Deep Neural Networks)：深度神经网络

CNN (Convolutional Neural Networks)：卷积神经网络

Fully connected layer, Sparsely connected layer, Convolutional layer

卷积层改进：算法改进、代码改进

矩阵乘法优化：利用局部性原理、多线程、SIMD

Blocked dense matrix multiplication 分块矩阵乘法：减少 Cache miss

Hierarchical blocked matrix multiplication 分层分块矩阵乘法：减少 Cache miss，

多种不同的矩阵乘法优化 ......

底层对 DNN 的优化，CuDNN，Tensorflow，使用低精度 ......

Lecture 11: Cache Coherence

2024/4/6

Cache / 内存一致性问题

（Cache miss 时）Read-through：直接从内存读取，Read-allocate：将数据读入 Cache

（Cache hit）Write-back：只有在 Cache Line 被替换时才写回内存，Write-through：每次写操作都写回内存

（Cache miss）Write-allocate：将数据读入 Cache 后再写入，Write-no-allocate：直接写入内存

外部 IO：DMA (Direct Memory Access) 控制器直接访问内存

SWMR (Single Writer Multiple Reader) Invariant：一个写者多个读者，读者可以并行读取，写者需要独占

Data-Value Invariant (write serialization)

Snooping cache-coherence schemes：每个 Cache 都监听内存操作

Invalidation-based protocol：当一个 Cache 修改数据时，通知其它 Cache 使其失效

MSI protocol：Cache 状态 Modified, Shared, Invalid，处理器操作 PrRd, PrWr，总线事务 BusRd, BusRdX, BusUpd。对应有 MSI State Diagram

MESI protocol：Cache 状态 Modified, Exclusive, Shared, Invalid

监听广播的方案不可扩展，多核时会有大量广播，使用 Directory 优化

False sharing：多线程读取同一 Cache Line 中的不同变量，由于一致性协议导致额外开销

Lecture 12: Memory Consistency

Consistency：多线程操作一个数据，由硬件提供机制，软件解决

Coherence：一份数据有多份拷贝（如 Cache），由硬件解决

2024/4/7

Memory ordering：指令执行顺序

Sequential consistency：严格顺序执行

Relaxed consistency：允许乱序执行，提升性能

乱序执行导致指令的执行顺序改变，在单线程下不会有问题，但多线程中不同线程观察到的行为不同，会出现与预期行为不一致的问题

（感觉课件介绍太浅没看明白，正好一直没理解 C++ 的 Memory Ordering，抽时间再学了一下）

---

完成了 CS149 Programming Assignment 3

---

Lecture 13: Fine-Grained Synchronization and Lock-Free Programming

2024/4/15

死锁：线程之间限制导致无法继续执行。条件：互斥、持有并等待、不可抢占、循环等待

活锁：线程一直做无用功（如 retry）

饥饿：少量线程正常执行

Test-and-Set：原子操作，设置一个变量并返回原值

Compare-and-Swap / Exchange：原子操作，比较并交换

基于 TAS 的锁、基于 Test-and-Test-and-Set 的锁：延迟、通信开销、可扩展性的区别

Ticket Lock：每个线程取一个号，按顺序执行

基于基础的原子操作实现其它原子操作

细粒度锁：提高并行度与产生额外开销之间的权衡

2024/4/16

无锁数据结构：部分场景性能更优

Single reader single writer queue

Lock-free stack / linked list

Lecture 14: Domain Specific Programming Languages

2024/4/17

Domain Specific Language (DSL)：针对特定领域的编程语言

Halide：图像处理领域的 DSL

Lizst：求解网格上 PDE

Good DSL system：易用、高性能、好语法

Lecture 15: Transactional Memory 1

Declarative 声明式 / Imperative 命令式抽象：描述做什么 / 描述怎么做

Transactional Memory 事务性内存：原子并且隔离的内存操作序列（类似数据库事务）、原子性、隔离性（事务提交前不可见）、可串行性。语义抽象为 atomic block

事务性内存中的事务可以看成一串对部分内存地址的读写操作，抽象为读和写的集合

通过对事务内存读写集合的比较来判断是否真正有冲突（Read-Write 冲突、Write-Write 冲突）

2024/4/18

事务性内存优点：易用、高效、失败后自动回滚

事务（atomic block）与锁的不同：抽象级别、锁的功能更多

TM 实现：数据版本控制策略（Undo log / Write buffer）、冲突检测（乐观 / 悲观）

Lecture 16: Transactional Memory 2

（这后面的内容感觉基本都是介绍性质，就简单的看一遍过了）

STM（Software Transactional Memory）实现：Intel STM

HTM（Hardware Transactional Memory）实现

异构计算：使用混合计算资源，高效。如 CPU + GPU，具体到 CPU 中也有多种计算单元

不同架构硬件的例子

Lecture 17: Hardware Specialization

2024/4/19

追求能耗比，Choosing the right tool for the job

Spatial Language：描述硬件结构

Streaming Execution Model：数据流执行模型

Kernel vs Stream

HBM (High Bandwidth Memory)：高带宽内存

Lecture 18: Accessing Memory + Course Wrap Up

DRAM access latency is not fixed

最佳情况：Read（CAS）

最坏情况：Precharge -> Activate -> Read（PRE + RAS + CAS）

---

完成了 CS149 Programming Assignment 4

---

完成了 CS149 Programming Assignment 5