内容参考自：General-Purpose Graphics Processor Architecture

GPU 硬件通过很宽的 SIMD 硬件挖掘 GPU 程序中的数据级并行。但其 SIMD 硬件并不直接暴露给程序员，程序员使用 SIMT 执行模型：准备好大量的标量线程组给 GPU ，每个线程有自己独立的执行路径，这些标量线程组被 NVIDIA 称为 warps (AMD 称为 wavefronts), 在 GPU 上通过 SIMD 硬件实现，每个线程内部是带锁的

执行模型

对于 discrete GPU ，程序首先执行 CPU 部分代码：

分配 GPU 计算需要的显存空间 (cudaMalloc)
将输入数据初始化到 GPU 显存中 (cudaMemcpy)
进入 GPU 计算 kernel
通过（<<<nblock, thread_per_block>>>）配置线程数(每个 block 将在一个 SM 上运行，因此 block 数不能超过 GPU 中的 SM 数)

NVIDIA 引入 Unified Memory 之后，可以透明地进行 CPU 和 GPU 之间的数据传输，其他架构需要程序员手动管理 GPU 显存的分配和传输

GPU 计算 kernel

GPU 计算 kernel 一般由数千个线程组成，这些线程在 kernel 中的组织结构:

1 个 grid 由 thread blocks (in NVIDIA) / cooperative thread array, CTA (in AMD) 组成
thread blocks / CTA 由 wraps / wavefronts 组成

在程序中 kernel 由一个函数指定( CUDA 中通过 __global__ 关键字指定)

NVIDIA 每个 wrap 由 32 个线程组成 (AMD 每个 wavefront 由 64 个线程组成)

thread blocks / CTA

每个 thread blocks / CTA 内的线程之间可以通过计算核的 SPM 高效通信
- 每个计算核有一个SPM, 供计算核上的所有的 thread block/CTA 使用
- SPM 一般很小，每个计算核的 SPM 一般在 16-64 KB，并作为不同的内存空间提供给程序员
- 在 CUDA 中通过 __shared__ 指定将内存分配到 SPM
- SPM 功能类似于软件管理的 cache, 但计算核中还有硬件管理的 cache
- NVIDIA 称为 shared memory
- AMD GCN 架构将该 SPM 称为 local data store, LDS
  除此之外，所有 compute core 还会共享一个 global data store, GDS
每个 thread blocks / CTA 内的线程可以使用硬件支持的 barrier 指令进行高效同步
- 不同 CTA 之间的线程可以进行通信，但需要通过全局地址空间同步，代价更高

线程 thread

每个线程执行相同的代码
- 这些指令的操作数均为标量操作数
- 但是遵循不同的控制流。
每个线程有自己的 id
- CUDA 中 grid, block, thread 这三个层级都有相应的 id
- 每个线程在其所属的 grid, block 中有唯一的非负坐标 (x,y,z)
  - thread in block: threadIdx.x, threadIdx.y, threadIdx.z
  - max thread id in block: blockDim.x, blockDim.y, blockDim.z
  - block in grid: blockIdx.x, blockIdx.y, blockIdx.z

GPU ISA

NVIDIA

GPU 程序编译过程：

编译： CUDA C 代码 → PTX 汇编代码
(nvcc 编译器)
汇编: PTX 汇编代码 → SASS 二进制代码
(ptxas 汇编器)

Parallel Thread Execution ISA, PTX

类似于 RISC ISA, NVIDIA 有完整的文档公开
- GPGPU-sim 基于 PTX 开发
与 RISC ISA 不同的地方在于：
1. PTX 可以使用无限的寄存器号，但其编译出的 SASS 只能使用有限的寄存器号
2. kernel 参数通过存储常量内存传递，SASS中的非访存指令可以访问这些常量内存，而PTX中的参数则被分配到各自独立的“参数”地址空间中

Streaming ASSembler, SASS

NVIDIA 未提供完整的文档，只提供了指令 opcode，未提到操作数域和指令语义

无法进行自定义修改
不具有后向兼容性
例: NVIDIA 从 Kepler 架构开始不再需要使用计分板进行显式的依赖检查

有一些工作是针对 SASS 的反汇编: decuda
NVIDIA 也提供反汇编工具: cuobjdump

AMD GCN

有完整的硬件级 ISA spec