ayayadesu.net

[ blog / entry ]

CS149 Lecture 7：GPU Architecture & CUDA Programming

~13 min read

学习笔记

#CS149.Parallel-Computing
[ contents ] 6 sections

通过 CUDA 的编程语言抽象&实现来讨论 GPU 架构

我们前面知道, 基础的 GPU 架构就是大量的处理单元集群, 通常有特点

  • 多核
  • 每个核相同的 SIMD 执行器 (相同指令进行大量的操作)
  • 每个核上进行多线程操作

GPU history

GPU 最初主要用于加速 3D 图形渲染: 对顶点、三角形和像素进行大量相似的计算, 包括几何变换、光栅化和着色; 这不等同于对大量光线进行追踪

对于材质的渲染效果, 我们知道这可以由着色器 shader program 实现, 例如根据材质、光照和观察方向计算表面颜色; fragment shader 通常输出颜色等属性, 不一定输出光线方向

为什么 GPU 有如此多的计算核心? 实际上这是设计瓶颈导致的, 我们已经知道单个计算单元的 clock speed, power, ILP 已经到达瓶颈, 我们的发展是给计算核心增加更多的晶体管, 这意味着更高的并行度

在世纪初的时, 人们意识到 GPU 这种高效计算处理器可以用于科学计算

GPU compute mode

review 一下我们前面讲到 CPU 是如何运行一个代码的:

  • OS 把程序的 text 加载到内存中
  • OS 选择 CPU 的执行 context
  • OS 介入处理器, 准备执行 context(如设置 registers 内容等)
  • 执行
  • 处理器从环境维护的执行 context 中执行指令

对于 GPU 上的代码, 之前是很复杂的: 早期的 GPU 只能处理 graphics pipeline 上的计算, 于是接口等基建都需要针对问题人为设计 或者把一个计算问题包装为一个图像渲染问题交给 GPU

NVIDIA 在 Tesla 架构时代推出了 CUDA (2007), 为 GPU 硬件提供了计算模式 (non- graphics- specific) 接口, 让用户能够直接指定 GPU 的计算任务, 对应的并行计算平台和编程模型称为 CUDA

CUDA programming language

这里使用 NVIDIA 在 2007 引入的 “C-like” 的 CUDA C/C++ 扩展, 用于表达通过计算模式接口在 GPU 上运行的程序 相对比较底层: CUDA 抽象与现代 GPU 的 capabilities/performance 特征比较接近

CUDA 程序保留了并发线程的层级: 线程 ID 可以是一个至多 3维的列表 由 Grid -> Block -> Thread 的层级来划分

一个 CUDA 程序的代码如

const int Nx = 12;
const int Ny = 6;

dim3 threadsPerBlock(4, 3);
dim3 numBlocks(Nx/threadsPerBlock.x,Ny/threadsPerBlock.y);


// assume A, B, C point to device arrays of Ny rows and Nx columns

// this call will launch 72 CUDA threads:
// 6 thread blocks of 12 threads each
matrixAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);

如代码表述, 每个 Block 的形状为 (4,3), Grid 的形状为 (3,2), 共 6 个 Block、72 个线程. 此处 Nx、Ny 恰好整除 block 尺寸; 一般情况需要向上取整, 并在 kernel 中检查边界

注意这不是 GPU 内处理器的结构, 这是逻辑线程

一般来说 CUDA 代码分为两部分: 一部分为 “Host” code :用于调度总的执行序列, 以及线程的定义等 另一部分为 CUDA kernel : 即进行 SPMD 真正执行的函数, 如

// kernel definition (runs on GPU)
__global__ void matrixAdd(float A[Ny][Nx],
float B[Ny][Nx],
float C[Ny][Nx])
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int j = blockIdx.y * blockDim.y + threadIdx.y;

	C[j][i] = A[j][i] + B[j][i];
}

__global__ 表明这是一个 CUDA kernel function

这个 function 的操作很简单, 通过每个线程获得的 threadIdx 和 blockIdx 来计算矩阵对应位置的加法

一个 CUDA 代码的结构即为上面

  • 由 CPU 执行的 Host code
  • 由 GPU 执行的 Device code (如 SPMD 操作)

对于 CUDA 底层来说, CUDA threads 是显式暴露在程序中的, 如上面, 我们也可以通过 i < Nx && j < Ny 来保护是否越界 CUDA 线程的启动数量是 programmer 指定的, 而不是根据数据数量启动的

这里采用独立显存 GPU 的基本模型: Host code 在 CPU 上运行, 启动 GPU 上的 Device code; host memory 与 device memory 分开管理, 通过显式数据复制交换内容. 这不是所有 CUDA 系统的绝对限制, Unified Memory 或映射的 host memory 也可以提供共享访问

                    CUDA

        HOST                          DEVICE
     +---------+                   +---------+
     |   CPU   |                   |   GPU   |
     |         |   launch kernel   |         |
     | serial  | ----------------> |  SPMD   |
     | code    |                   | threads |
     +----+----+                   +----+----+
          |                             |
          |                             |
     +----v----+   copy data       +----v----+
     | Host    | <---------------> | Device  |
     | Memory  |                   | Global  |
     | (RAM)   |                   | Memory  |
     +---------+                   | (VRAM)  |
                                   +---------+

这涉及到 cudaMemcpy API, 下面是一个示例

float* A = new float[N]; // allocate buffer in host mem
// populate host address space pointer A
for (int i=0; i<N; i++)
	A[i] = (float)i;

size_t bytes = sizeof(float) * N;
float* deviceA;
cudaMalloc(&deviceA, bytes); // allocate buffer in device memory

// populate deviceA
cudaMemcpy(deviceA, A, bytes, cudaMemcpyHostToDevice);

// note: directly accessing deviceA[i] from host is invalid here
// (cannot manipulate contents of deviceA
// directly from host, since deviceA is not a pointer
// into the host’s address space)

实际上的 CUDA 内存抽象是这样的:

每个 thread 有自己的私有数据, 每个 block 有由该 block 的线程共享的 shared memory; device global memory 则可由各个 block 的线程访问. thread 私有变量常放在 registers 中, 溢出等情况会使用 local memory; local 指线程私有的作用域, 不保证物理上在片上

不同的地址空间反映了程序中不同区域的 locality (同一个线程, 同一个 block 的线程等)

这种 locality 的分离提供了计算性能优化的空间

CUDA example: 1D convolution

我们以一维的卷积操作为例子, 如实现

output[i] = (input[i] + input[i+1] + input[i+2]) / 3.f

我们的 Host code 很简单

#define THREADS_PER_BLK 128
int N = 1024*1024; // 本例 N 可以被 THREADS_PER_BLK 整除
float *devInput, *devOutput;
cudaMalloc(&devInput, sizeof(float) * (N+2));
cudaMalloc(&devOutput, sizeof(float) * N);
// 在启动 kernel 前, 将 N+2 个有效输入值复制到 devInput

convolve<<<N/THREADS_PER_BLK, THREADS_PER_BLK>>>(N, devInput, devOutput);

下面保留课堂示例的整除条件: N > 0 且 N % THREADS_PER_BLK == 0, 输入长度为 N+2, 每个 block 恰有 THREADS_PER_BLK 个线程. 代码片段省略错误检查、结果回传和内存释放.

区别即为 kernel 对 convolve 的实现:

一种简单的想法是: 让每个线程计算输出的每个位置即可

#define THREADS_PER_BLK 128

__global__ void convolve(int N, float* input, float* output){
	int index = blockIdx.x * blockDim.x + threadIdx.x;

	float result = 0.0f;
	for (int i=0; i<3; i++){
		result += input[index + i];
	}

	output[index] = result / 3.f;
}

但是这种问题是: 内部输入元素最多被 3 个线程重复读取, 每个 block 共发出 3 × 128 次 global load. 这些访问可能命中 cache, 不能把它们直接等同于 384 次显存事务 一种优化是: 先把需要处理的数据加载到 block 内共享的 memory 提供更快的速度

以下片段替换上面 kernel 中计算 index 之后的部分, 仍要求相同的 block 大小与整除条件:

__shared__ float support[THREADS_PER_BLK+2]; // alloc block 的内存

support[threadIdx.x] = input[index]; // 从 global memory 搬运到 shared memory
if (threadIdx.x < 2){
	support[THREADS_PER_BLK + threadIdx.x] = input[index + THREADS_PER_BLK];
} // 对 boundary 特殊处理需要加载更多的数据

__syncthreads(); // 等待所有线程加载完数据

float result = 0.0f;
for (int i =0; i<3; i++){
	result += support[threadIdx.x + i];
}

output[index] = result / 3.f;

这样每个 block 只需加载 128+2=130 个输入元素. 是否更快还取决于 cache、同步和 shared memory 的开销, 需要实测. 若改成向上取整启动来支持任意 N, 还必须处理末尾 block 的输入、halo 和输出边界, 并保证参与的线程一致到达 __syncthreads().

这里实际上也能看到: CUDA 提供了同步等接口:

  • __syncthreads(): block 内的 barrier, 等待同一 block 的线程到达, 并保证同步前的 shared/global memory 访问对 block 内其他线程可见
  • Atomic operations: 在global 和 per-block 共享内存上提供的 atomic 操作
  • Host/device synchronization: kernel<<<...>>>() 对 host 通常是异步的, launch 返回不代表 GPU 已执行完. Host 需要结果时, 可用 cudaDeviceSynchronize()、stream/event 同步, 或适当的阻塞式结果回传等待完成; kernel 完成表示其所有线程已结束

Summary: CUDA abstractions 执行: 线程层级

  • Bulk launch of many threads
  • Threads 组成 Block, Blocks 组成 Grid

分布式内存地址空间:

  • cudaMemcpy API 负责在 host 和 device 之间的数据的复制
  • 三个不同类型的 device 的内存空间
  • per-thread, per-block(“shared”), per-program(“global”)

在 block 内为线程提供的同步语句与更多同步操作的 atomic 语句

CUDA semantics

CUDA launch 指定许多逻辑线程, 但不会为每个 CUDA thread 创建一个 std::thread() 那样的 OS 线程, 也不意味着所有线程同时驻留在 GPU 上 前面的 CUDA 程序不需要知道核心数 (CUDA API 可以查询 SM 数量等设备属性), CUDA 线程类似于我们前面讨论过的在 data parallel model 中的例子, 只是说有这么多工作可以并行

一个编译好的 CUDA device 二进制文件包含 Program text(instructions) Information about required resources:

  • 每个线程所需的 registers 和 local memory
  • 每个 block 所需的静态 shared memory

每个 block 的线程数由 launch 的 blockDim 指定, 不是一般情况下编译时固定的值; 动态 shared memory 大小也可以在 launch 时指定

完成逻辑线程的任务由 CUDA 调度进行, 在 host code 启动的 kernel 会先交给 GPU 硬件上的 scheduler, 然后 scheduler 再将任务分配给计算核心

CUDA 的核心假设: 每个 thread block 的任务可以以任意顺序进行, 这种任务调度通过动态调度进行

以 NVIDIA V100 SM 为例, 来看它的 “sub-core” 设计

这里的 core 指 SM (Streaming Multiprocessor), 不是单个 CUDA core. SM 中有适合各种运算的 ALU, 以及较大的 register file

在一个 thread block 中, 32 个 threads 叫做一个 warp warp 是 SM 调度和发射指令的基本单位, NVIDIA 将这种执行方式称为 SIMT (Single Instruction Multiple Threads). 同一 warp 内线程沿相同路径执行时, 可以共同执行同一条指令; 若发生分支分歧, 不同路径分开执行并屏蔽不参与的 lanes, 可能降低利用率. 这不意味着整个 GPU 一次只能运行一个 warp

warp 不属于前面 Grid -> Block -> Thread 的基本任务划分层级, 但 CUDA 也暴露了 warpSize、warp shuffle 和 __syncwarp() 等接口; 不能说 warp 完全不属于 CUDA

根据之前的讨论, 对于指令流, SM 可以在多个就绪 warp 之间调度, 用其他 warp 的工作隐藏访存等延迟, 提高利用率

Running a CUDA program on a GPU

大致的顺序如下:

  1. host 把命令发送给 CUDA device, 如执行某个 kernel
  2. scheduler 将 block 映射到 SM
  3. scheduler 继续给 blocks 映射执行 context, 通过动态调度
  4. 如果有 block 在对应 core 上完成任务后, scheduler 继续把新的任务映射过去

Why must CUDA allocate execution contexts for all threads in a block? 例如考虑一个包含 256 个 CUDA threads 的 block, 假设 SM 的执行宽度只能同时处理其中 128 个线程的工作. 为什么不能让前 128 个线程全部完成后, 再开始后 128 个?

因为这样在 kernel 中的 __syncthreads 的线程级管理代码会直接 deadlock, 我们根本没有启动剩下的线程 所以 CUDA 必须为驻留 block 中所有线程保留执行状态, 再在它们的 warps 之间交错推进. 执行宽度不等于可驻留线程数; 若 block 的线程数或资源需求超过设备限制, launch 会失败, 不能靠把同一 block 分成互不重叠的两批来执行

这实际上就是 semantics 和 implementation 的区别:

block 是资源分配的单位, warp 是实际调度执行的单位

CUDA implementation 的 abstraction:

  • 普通 launch 的 thread blocks 可以以任意顺序调度 (不能依赖其他 block 同时驻留或先执行)
  • CUDA threads in the same block run concurrently (live at the same time): 一个 block 中的 threads 的生命周期重叠
  • CUDA implementation: GPU 执行的单位是 warp, 但调度器分配资源的单位是 block

我们考虑另一个例子: 让一个程序绘制直方图 所有 CUDA threads atomically 更新全局内存中的共享变量

事实上, 我们并未假定 CUDA blocks 之间是独立的, 只是说他们可以以任意顺序调度, 对同一计数器的 atomicAdd 保证更新不会丢失, 而且不要求某个 block 先运行; 这不代表 atomics 自动消除所有数据竞争或充当全局 barrier

考虑一个 single SM GPU, 资源只允许同时驻留一个 block, 但 grid 中有两个 blocks 考虑给一个 block 先计算, 然后另一个 block 等待前一个完成, 再计算 这不是一个合理的 CUDA code, CUDA program 本身并不保证哪个 block 先进行计算, 这极易导致 dead lock

课件最后的 persistent threads 思路, 是启动少量 blocks 并让它们反复从工作队列取任务. 但仅手写 block 数量, 并不能在普通 launch 中强制所有 blocks 同时驻留; 这依赖具体硬件和 kernel 的资源占用, 不能当作可移植的同步保证. 若确实需要 grid 范围的 barrier, 可在支持的设备上使用 cooperative launch 配合 Cooperative Groups 的 grid.sync(), 并满足驻留资源限制; 或把阶段拆成同一 stream 中顺序启动的多个 kernels

CUDA summary

  • Execution semantics: 按照 data parallel model, 将任务分解到各个 thread blocks 上. Thread blocks 中的 threads 同时存在. 当然有不同的执行模型, 在面对一个并行程序系统的时候时刻记住这一点
  • Memory semantics: 分布式地址空间 (host/device) , 线程 local/线程 block 级/ 全局 变量
  • Key implementation details: thread block 和 warp 的区别

核对参考: CUDA Programming Model、Asynchronous Execution、Cooperative Groups。

← Back to blog