GPU 并行计算
了解图形处理器并行计算的基本概念，包括 GPU 架构特点、并行计算模型以及元计算统一系统架构 MUSA 编程语言的核心抽象。
GPU 概述
图形处理器 GPU（Graphics Processing Unit）是一种专为并行计算设计的处理器。GPU 最初用于加速计算机图形渲染，由于其高吞吐量的计算特性，现已广泛应 用于通用并行计算 GPGPU（General-Purpose GPU Computing）领域。
GPU 与 CPU 的架构差异
图形处理器 GPU 和中央处理器 CPU 采用不同的设计目标：
| 特性 | CPU | GPU |
|---|---|---|
| 设计目标 | 低延迟，优化单线程性能 | 高吞吐量，优化并行处理能力 |
| 核数 | 较少（通常 4-64 核） | 大量（数千个计算单元） |
| 缓存设计 | 大容量缓存，复杂控制逻辑 | 较小缓存，更多计算单元 |
| 适用场景 | 串行任务、复杂分支逻辑 | 大规模数据并行计算 |
GPU 和 CPU 硬件逻辑结构对比，如下图所示：
- CPU：将更多晶体管用于缓存和复杂控制逻辑，通过大缓存和低延迟分支预测优化 单线程执行效率。
- GPU：将更多晶体管用于数据处理单元，通过 大规模并行计算 和 线程切换隐藏内存延迟 实现高吞吐量。
以 Intel Xeon 8280 和 MTT S5000 为例，具体型号对比如下所示：
| 指标 | CPU (Intel Xeon 8280) | GPU (MTT S5000) |
|---|---|---|
| 核数 | 28 核 | 4,096 核 |
| 时钟频率 | 2.7 GHz | 1.8 GHz |
| 内存带宽 | 140 GB/s | 448 GB/s |
| 并发线程数 | ~1,792 | ~196,608 |
| 单精度浮点性能 | 4.8 TFLOPS | 14.7 TFLOPS |
应用场景
GPU 并行计算适用于具有高并行度、高计算密度的应用场景。在智能驾驶、智能机器人、新材料发现、脑神经科学以及人工智能时代的科学研究中，GPU 已成为关键的计算基础设施。
| 应用领域 | 典型应用 |
|---|---|
| 科学计算 | 分子动力学模拟、气候建模、流体力学计算 |
| 人工智能 | 深度学习训练与推理、大语言模型 |
| 金融计算 | 蒙特卡洛模拟、期权定价、风险评估 |
| 医学影像 | CT/MRI 图像重建、病理图像分析 |
| 自动驾驶 | 传感器数据融合、实时路径规划 |
GPU 并行计算模型
异构计算架构
MUSA 采用异构计算模型：CPU 作为主机（Host），GPU 作为设备（Device）。Host 和 Device 拥有独立的内存空间，通过 PCIe (高速外围组件互联)等总线进行数据传输。
┌─────────────────┐ ┌─────────────────┐
│ Host │ │ Device │
│ CPU │ │ GPU │
│ ┌───────────┐ │ │ ┌───────────┐ │
│ │ Host │ │ PCIe │ │ Device │ │
│ │ Memory │ │<------->│ │ Memory │ │
│ └───────────┘ │ Bus │ └───────────┘ │
└─────────────────┘ └─────────────────┘
典型执行流程：
- 主机（CPU）：负责顺序逻辑、内存管理、内核（Kernel）启动
- 设备（GPU）：负责大规模并行计算
- 数据传输：通过
musaMemcpy()等 API 在主机和设备间拷贝数据
线程层次结构
MUSA 采用三层线程层次结构组织并行计算：
- Grid（网格）：一次 Kernel 启动的所有线程
- Block（线程块）：线程组，可独立调度执行
- Thread（线程）：单个执行单元
线程层次结构
Grid
├── Block (0, 0)
│ ├── Thread (0, 0)
│ ├── Thread (1, 0)
│ └── ...
├── Block (1, 0)
│ ├── Thread (0, 0)
│ └── ...
└── ...
每个线程通过内置变量获取其在层次结构中的位置：
int thread_id = blockIdx.x * blockDim.x + threadIdx.x;
| 内置变量 | 说明 |
|---|---|
threadIdx | 线程在 Block 内的索引 |
blockIdx | Block 在 Grid 内的索引 |
blockDim | Block 的维度大小 |
gridDim | Grid 的维度大小 |
可扩展编程模型
多核 CPU 和众核 GPU 的出现意味着主流处理器芯片是并行系统。编程模型面临的挑战是如何开发应用软件，能够透明地扩展其并行性，充分利用不断增加的处理器内核。
MUSA 并行编程模型通过三个核心抽象克服这一挑战：
| 抽象 | 说明 |
|---|---|
| 线程组的层次结构 | Grid、Block、Thread 三层组织 |
| 共享内存 | Block 内线程间数据交换 |
| 屏障同步 | Block 内线程协作同步 |
这种设计实现自动可扩展性：
- 线程块可独立调度，可在任意多处理器上执行
- 编译后的程序可在任意数量多处理器上运行
- 运行时系统自动适配物理多处理器数量
无需修改代码，程序即可在不同核心数的 GPU 上透明扩展。
Block 可被调度到 GPU 上的任意 MP（MUSA Processor）执行，执行顺序不确定。如果允许跨 Block 同步，程序将依赖特定的执行顺序，导致无法在不同核心数的 GPU 之间移植。MUSA 编程模型要求 Block 可以按任意顺序、并行或串行执行，从而保证程序的可扩展性。
SIMT 执行模型
GPU 采用单指令多线程 SIMT（Single Instruction Multiple Threads）执行模型：
- 一组线程（Warp）同时执行同一条指令
- 每个线程具有独立的寄存器状态
- 当线程执行路径出现分支时，Warp 序列化执行各分支
Warp 大小：不同 GPU 架构的 Warp 大小不同
- MTT M1000/S4000：Warp = 128 线程
- MTT S5000：Warp = 32 线程
选择 Block 大小时，建议使用 Warp 大小的倍数（如 128、256、512、1024）。
示例：SAXPY 运算（y = a*x + y）
// CPU 串行版本
void saxpy_serial(int n, float a, float *x, float *y) {
for (int i = 0; i < n; ++i)
y[i] = a * x[i] + y[i];
}
// GPU 并行版本——每个线程处理一个元素
__global__ void saxpy_parallel(int n, float a, float *x, float *y) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
y[i] = a * x[i] + y[i];
}
// 启动：数千个线程同时执行
int nblocks = (n + 255) / 256;
saxpy_parallel<<<nblocks, 256>>>(n, 2.0f, x, y);
单指令多线程（SIMT）模型使程序员能够以标量线程方式编写并行代码，同时保持高执行效率。
GPU 内存层次结构
内存类型
GPU 内存系统包含多个层次，各层次在访问速度、容量和可见范围上存在差异：
| 内存类型 | 访问速度 | 容量 | 可见范围 |
|---|---|---|---|
| 寄存器（Register） | 最快（1 周期） | 每线程有限 | 线程私有 |
| 共享内存（Shared Memory） | 快（1-2 周期） | 每 Block 有限 | Block 内共享 |
| 常量内存（Constant Memory） | 中（缓存命中时快） | 较小 | 全局可见 |
| 全局内存（Global Memory） | 慢（100+ 周期） | 大 | 全局可见 |
内存访问优化
高效的 GPU 程序需要合理分配数据到不同内存层次：
- 寄存器：存储线程局部变量，由编译器自动分配
- 共享内存：存储 Block 内线程共享的数据，需程序员显式声明
- 全局内存：存储主机与设备间传输的数据，需显式分配和拷贝
MUSA 编程模型核心抽象
Kernel 函数
Kernel 是在 GPU 上执行的函数，使用 __global__ 修饰符声明：
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
Kernel 函数特点：
- 返回类型必须为
void - 由 Host 调用，在 Device 上执行
- 通过执行配置语法
<<<grid, block>>>指定线程数量
Kernel 启动
Kernel 启动需要指定执行配置参数：
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
| 参数 | 说 明 |
|---|---|
gridSize | Grid 中 Block 的数量 |
blockSize | 每个 Block 中的线程数量 |
sharedMemSize | （可选）动态分配的共享内存大小 |
stream | （可选）执行流标识符 |
内存管理
MUSA 提供与 C 标准库类似的内存管理 API：
| API | 说明 |
|---|---|
musaMalloc(&ptr, size) | 分配设备内存 |
musaMemcpy(dst, src, size, kind) | 在 Host 和 Device 间拷贝数据 |
musaFree(ptr) | 释放设备内存 |
线程同步
Block 内线程可通过屏障同步进行协作：
__syncthreads(); // 等待 Block 内所有线程到达此同步点
同步约束：
- 仅支持 Block 内同步
- 不支持跨 Block 同步（保证程序可移植性）
GPU 程序执行流程
典型的 GPU 程序执行流程如下：
- Host 分配内存并初始化数据
- Host 分配 Device 内存
- Host 将数据拷贝到 Device
- Host 启动 Kernel，在 Device 上并行执行
- Device 完成计算
- Host 将结果拷贝回 Host 内存
- Host 验证结果并释放内存
完整代码示例,参见 快速入门。
