编程模型基础：从抽象到实战

各位同学好，今天我们聊聊编程模型。说实话，这个概念刚入行时我也觉得挺虚的——不就是写代码吗？但后来踩过几次坑才明白，不理解编程模型，你连硬件都使唤不动。

什么是编程模型？

编程模型，说白了就是你告诉硬件「该干什么」的那套规则。它定义了你怎么组织代码、怎么管理数据、怎么让计算单元协同工作。

我习惯把它比作「厨师和厨房的沟通协议」。你想想看：

• 你告诉厨师「炒个菜」——这是高级指令
• 你告诉厨师「开火、倒油、放菜、翻炒」——这是低级指令
• 编程模型就是决定你用哪种方式沟通的那套约定

在加速器领域，编程模型通常包含三部分：

• 执行模型：线程怎么创建、怎么调度、怎么同步
• 内存模型：数据怎么存放、怎么访问、怎么共享
• 编程接口：你实际写的API、关键字、编译器指令

编程模型与硬件的关系

这里我想强调一点：编程模型不是凭空设计的，它是对硬件架构的抽象。

举个例子。GPU为什么有「线程块」这个概念？因为GPU的SM（流式多处理器）一次只能执行一组线程。CPU为什么用「线程池」？因为CPU核心少，但每个核心能力很强。

我在项目中遇到过一件事：有个同事把CPU上的多线程代码直接搬到GPU上，结果性能反而下降了。为什么？因为CPU的编程模型假设线程是「重量级」的，创建和切换开销大；而GPU的编程模型假设线程是「轻量级」的，成千上万个线程同时跑才有效率。

这张图展示了编程模型与硬件之间的映射关系：

主流加速器编程模型对比

现在市面上主流的加速器编程模型有四个：CUDA、OpenCL、OneAPI、ROCm。我一个个说。

1. CUDA — 老大哥，生态最完善

CUDA是NVIDIA的专属编程模型。说实话，如果你做AI或者高性能计算，CUDA几乎是绕不开的。

它的核心概念：

• 线程层次：thread → warp → block → grid
• 内存层次：寄存器 → 共享内存 → L1/L2 → 全局内存
• 编程方式：用C++扩展，加__global__、__device__关键字

// CUDA示例：向量加法
__global__ void vecAdd(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    vecAdd<<<256, 256>>>(d_a, d_b, d_c, N);
    return 0;
}

2. OpenCL — 跨平台，但用起来有点累

OpenCL是Khronos制定的开放标准。理论上支持CPU、GPU、FPGA、DSP等各种设备。但说实话，用起来比CUDA繁琐不少。

它的核心概念：

• 平台模型：Host + Device(s)
• 执行模型：NDRange → Work-group → Work-item
• 内存模型：全局内存 → 常量内存 → 局部内存 → 私有内存

// OpenCL示例：向量加法（注意上下文代码量）
__kernel void vecAdd(__global float* a, __global float* b, __global float* c) {
    int i = get_global_id(0);
    c[i] = a[i] + b[i];
}

// 主机端需要：创建上下文、编译内核、设置参数、执行... 代码量是CUDA的2-3倍

3. OneAPI — Intel的野心，统一编程模型

OneAPI是Intel推出的统一编程模型。它的口号是「一次编写，到处运行」——支持CPU、GPU、FPGA，甚至其他加速器。

它的核心概念：

• DPC++：基于SYCL的C++扩展，用lambda表达式定义内核
• 统一内存：USM（统一共享内存），简化数据管理
• 设备选择：通过device_selector指定目标硬件

// OneAPI (DPC++) 示例：向量加法
#include <sycl/sycl.hpp>
using namespace sycl;

int main() {
    queue q;
    float* a = malloc_shared<float>(N, q);
    float* b = malloc_shared<float>(N, q);
    float* c = malloc_shared<float>(N, q);
    
    q.parallel_for(range<1>(N), [=](id<1> i) {
        c[i] = a[i] + b[i];
    }).wait();
    
    return 0;
}

4. ROCm — AMD的开源反击

ROCm是AMD的开源GPU计算平台。它兼容HIP（Heterogeneous-compute Interface for Portability），HIP代码可以很方便地移植到CUDA。

它的核心概念：

• HIP：类似CUDA的C++ API，支持AMD和NVIDIA
• ROCclr：运行时层，管理设备、内存、执行
• MIOpen：深度学习库，对标cuDNN

// HIP示例：向量加法（和CUDA几乎一样）
__global__ void vecAdd(float* a, float* b, float* c, int n) {
    int i = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
    if (i < n) c[i] = a[i] + b[i];
}

int main() {
    vecAdd<<<256, 256>>>(d_a, d_b, d_c, N);
    return 0;
}

横向对比总结

特性	CUDA	OpenCL	OneAPI	ROCm
硬件支持	仅NVIDIA	多厂商	Intel为主	AMD为主
编程语言	C++扩展	C99内核	DPC++ (SYCL)	HIP (C++类CUDA)
学习曲线	中等	较陡	中等	低（CUDA经验者）
生态成熟度	★★★★★	★★★☆☆	★★☆☆☆	★★★☆☆
跨平台能力	差	好	较好	中等
性能优化空间	大	中等	中等	大

好了，这一章就到这里。编程模型是加速器编程的「世界观」，理解了它，后面学具体技术就顺了。