[Lecture Notes] CMU 15-618 Parallel Computer Architecture & Programming

Lec 2. A modern multi-core processor

4 key concepts: 两个与parallel execution有关, 两个与challenges of accessing memory 有关

Parallel execution

Superscalar processor: Instruction level parallelism (ILP)
- ILP读未来的指令（每个周期读两条指令），有两个fetch/decode单元和两个exec单元，能够同时执行两条指令
Multi-core: 多个processing cores
- 多核之前，处理器提升重点在更大缓存，更多分支预测predictor等；同时更多晶体管（才能放得下更多缓存和更多predictor和乱序执行逻辑）促生更小的晶体管，促进更高的计算机主频
- 2004年多核出现之后，人们在一个chip上放多个processor，用更多晶体管放更多核心。
SIMD processing (aka Vector processing): 多个ALU(同一个core内)
- 仍然只需要一个fetch/decode单元，多个ALU。
- conditional execution: 如果想simd的程序块有if else，要通过mask处理
- 手写avx代码（cpu指令）是explicit SIMD; 而GPU是implicit SIMD，因为compiler生成的并不是并行指令（是普通的scalar instructions），只有GPU硬件运行才是SIMD的

Accessing memory

cache: reduce latency
prefetching reduces stalls: hides latency
Multi-threading, interleave processing of multiple threads
- 跟prefetching一样，也是hide latency，不能reduce latency
- 指的是：开多个线程，在一个线程卡住的时候执行别的线程
- 在下图中，创建thread1的时候不仅仅创建thread1，还会告诉电脑创建了thread 2 3 4，硬件检测线程是否发生了stall（被等待内存操作卡住），如果发生了stall会很快切换到别的线程，想juggling一样。硬件决定如何juggle这些线程。
- 这样memory latency仍然存在，但是被hide了。memory latency在后台发生，前台CPU一直在执行有用的工作。
- 这种操作会导致单个线程的执行时间变长（因为thread1从runnable到重新开始执行有一段空挡（这段空隙在执行thread 2 3 4）。
- 需要更多硬件资源，存储线程的register等状态信息，这样切换线程才会快。且需要较大的memory bandwidth。

GPU设计成处理大量数据（远大于核内缓存的数据量）。

与CPU内存相比，GPU显存带宽更高，但延迟更高。

Lec 3. Programming Models

Abstraction vs Implementation

abstraction和implementation不是一个东西！

ISPC (Intel SPMD Program Compiler)

SPMD: single program multiple data

一个ISPC计算sin(x)的例子

Interleaved

export void sinx(
  uniform int N, uniform int terms,
  uniform float* x, uniform float* result) {
  // assume N % programCount = 0
  for (uniform int i=0; i<N; i+=programCount) {
    int idx = i + programIndex;
    // 不重要
  }
}
<!--code0-->

在这个示例中blocked比interleaved更好。

因为每个iteration工作量完全相同，SIMD指令load连续内存（直接_mm_load_ps1）比load不连续内存(这种操作叫gather，只在AVX2及以后才支持）更快。

我们可以使用foreach来代替。

foreach表示循环中的每一次iteration都是独立的，由ISPC决定如何分配

1
2
3

foreach(i = 0 ... N) {
	// index ...
}

ISPC: abstraction VS. Implementation
- Programming model: SPMD
  - 程序员想的是，我们有programCount个逻辑指令流，写代码也按照这样的abstration去写
- Implementation: SIMD
  - ISPC输出SSE4或AVX这种指令来实现逻辑操作

Four Programming Models

Shared Address Space model
- 共享内存，不同线程通过读写同一块内存来通信
- 很多人认为这是最方便易用的model，因为这和sequential programming最相近
- 每个处理器都能访问任意内存地址
- Uniform memory access time; Symmetric shared-memory multi-processor (SMP), : 指每个处理器都能访问内存且访问内存所需时间一致。Scalability不好。
- Non-uniform memory access (NUMA): 每个处理器都能访问完整内存，但所需时间不一致。这样Scalability比较好。但是performance tuning可能需要更多精力。
Message Passing model
- 线程之间不共享内存（address space不共享），只能通过发送和接收messages通信
- 相比shared memory的优点：不需要别的硬件，可以纯软件做，实现起来简单。常用的库是MPI (message passing interface)
- 现代很常用的操作是，在一个节点（节点内是多核的）内用shared address，在不同节点间用message passing

这里很搞：abstraction can target different types of machines.

分清abstraction和implementation的区别！

比如说：

message passing 这个 abstraction 可以使用硬件上的 shared address space 来 implement。

发送/接收消息就是读写message library的buffer。

shared address space 这个 abstraction 在不支持硬件 shared address space 的机器上也可以用软件来 implement (但是低效)

所有涉及共享变量的page都标记成invalid，然后使用page-fault handler来处理网络请求

The data-parallel model
- 上面的shared address space和message passing更general
- data-parallel更specialized, rigid。
- 在不同的数据（数据块）上执行相同的操作。
- 通常用SPMD的形式：map (function, collection)，其中对所有的数据都做function的操作，function可以是很长的一段逻辑，比如一个loop body。collection是一组数据，可以包含多个数据。
- gather/scatter: gather是把本来不连续的数据按照index放到一起，scatter是把本来连续的数据分散开。
The systolic arrays model
- 读内存太慢了，memory bandwidth会成为瓶颈
- 所以要避免不必要的内存读取，尽量只读一次内存，做完所有需要用到这块内存的操作，再写回去。
- 示例：矩阵乘法 https://www.youtube.com/watch?v=2VrnkXd9QR8

一些基础知识

ISPC

ISPC代码调用时会生成多个program instances, 可以利用 programCount 和 programIndex 来获取instance总数和当前instance编号。

uniform 表示在一个SIMD程序块中，变量对所有SIMD通道都是相同的值。仅仅是一种优化，不影响正确性(因为uniform变量只需要加载一次或执行一次，编译器可以做出优化，不加uniform可能造成不必要的重复计算)。

非uniform (varying) 表示变量在不同SIMD通道可能有不同的值。

所以说 programCount 是 uniform, programIndex 是 varying.

ISPC可以通过tasks来实现多核加速，利用多线程。

Contrary to threads, tasks do not have execution context and they are only pieces of work. ISPC编译器接受tasks并自行决定启动多少个threads。

通常我们应该启动比cpu逻辑线程数更多的tasks数量，但也不要太多，否则会有scheduling的overhead。

task自带 taskIndex。

CUDA

host是CPU, device是GPU

__device__: 在device上执行，只能在device中调用

__global__: 在device上执行，只能在host中调用

__host__: 在host上执行且只能在host上调用

cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost)

threads grouped into blocks

需要指明blocks的数量，和每个block中threads的数量。

假设n是总的threads数量, t是每个block中threads的数量。

KernelFunction<<<ceil(n/t), t>>>(args)

每一个thread都会运行同样的kernel，每一个thread由blockID和这个block中的threadID来标识。

Example:

__global__ void vecAddKernel(float* A, float* B, float* C, int n) {
    int i = threadId.x + blockDim.x * blockId.x;
    if (i<n) C[i] = A[i] + B[i];
}
void vecAdd(float* A, float* B, float* C, int n) {
    int size = n * sizeof(float);
    float *d_A, *d_B, *d_C;
  
    cudaMalloc((void **) &d_A, size);
    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
    cudaMalloc((void **) &d_B, size);
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);
    cudaMalloc((void **) &d_C, size);
  
    vecAddKernel<<<ceil(n/256), 256>>>(d_A, d_B, d_C, n);
    
  	cudaMemcpy(C, d_C, size, cudaMemcpyDeviceToHost);
    cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
}

注: 为什么cudaMalloc第一个参数是二级指针，而不直接使用返回值来赋值给指针？

因为 cudaMalloc 的返回值已经用来返回 cudaError_t。

grid和blocks可以是1D, 2D, 3D的。上面这个例子是1D，所以是".x"

2D的例子：假设要把一个WIDTH x WIDTH的矩阵P分成几块。

WIDTH=8, TILE_WIDTH为2的话，就是把8x8的矩阵分成16个小块(grid)，每一个小块大小是2x2(4个thread)。

1
2
3

dim3 dimGrid(WIDTH / TILE_WIDTH, WIDTH / TILE_WIDTH, 1);
dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
MatrixMulKernel<<<dimGrid, dimBlock>>>(args);

每一个thread可以用以下方式来标识

1 2	`row = blockId.y * blockDim.y + threadId.y; column = blockId.x * blockDim.x + threadId.x;`

为什么用两层threads？因为组成多个grid的thread blocks比一个很大的单个thread block更好管理。

GPU有很多很多核心，核心group成SM(streaming multiprocessors)，每一组SM有自己的内存和调度。

GPU不同时启动所有100万个threads，而是把大约1000个thread装进一个block里，并分发给SM。

assign给SM的thread block会使用SM的资源（寄存器和共享内存）。这些资源已经pre-allocated，且由于寄存器数量很多，在切换threads时不需要register flush。

不同的block可以用任何顺序运行，因此不能assume block2在block1之后运行。如果真的要这么做，需要放在不同的kernel里（启动kernel比较耗资源）

同一个block中的thread可以使用 __syncthreads() 来做barrier synchronization。

但是通常不建议使用 __syncthreads()

如何选择合适的block size？

Consideration 1: hardware constraints
- 例如：每一个SM分配小于1536个thread，小于8个block；每一个block小于512个thread
Consideration 2: complexity of each thread
Consideration 3: thread work imbalance.

GPU memory

Global memory很慢，所以同时运行大量线程，线程因为内存IO卡住的时候切换其它线程，这是massive multi-threading (MMT).

这样总的throughput很高，即使每个thread的延迟也很高。

每个SM有自己的scheduler，每个SM存储了所有thread的context(PC, reg等)，所以SM内能做到零开销线程切换。同时，SM scheduler有一个scoreboard追踪哪些thread是blocked/unblocked，所以SM有大约30个核但可以运行大约1000个线程。

Tiled MM是一种进行矩阵乘法内存友好的方法。

CUDA类型关键词

__device__ __shared__ memory: shared; scope: block; lifetime: block
__device__ memory: global; scope: grid; lifetime: application
__device__ __constant__ memory: constant; scope: grid; lifetime: application

Race conditions:

CUDA中难以实现mutex，而且包含critical sections的代码在GPU上本来就运行得不好。

CUDA中有一些原子操作，可以在global或shared memory变量上操作

int atomicInc(int *addr): 加一，返回旧值
int atomicAdd(int *addr, int val): 加val, 返回旧值
int atomicMax(int *addr, int val): 让*addr=max(*addr, val) 并返回旧值
int atomicExch(int *addr1, int val): set
int atomicCAS(int *addr, old, new): Compare and swap.
- if (*addr == old) *addr = new;

Lecture Notes

#C #Lecture Notes #CUDA

[Lecture Notes] CMU 15-618 Parallel Computer Architecture & Programming

https://www.billhu.us/2025/054_cmu_15618/

Author

Bill Hu

Posted on

January 7, 2025

Licensed under

[Lecture Notes] CMU 14-740 Fundamentals of Telecommunication Networks Previous

[Lecture Notes] CMU 15-513 Intro to Computer Systems Next