「异构计算与 GPU 卸载编译」

1. 异构计算的编译挑战

一句话总结： 异构编译器要为一个程序生成两套指令流与两层内存管理，还要在两者之间插入数据搬运与同步，这比单目标编译多出一整类决策。

传统编译器的假设很简洁：一个程序、一种指令集、一层内存。编译器把源语言翻译成目标机器的指令，寄存器分配器决定值放在哪个寄存器，运行时管理一层堆内存。优化器在这个模型里工作得很好。

异构计算打破了这个假设。程序要在 CPU 与加速器（GPU、FPGA、DSP）上协同执行：

// 一个典型的卸载程序结构
void saxpy(float a, float *x, float *y, int n) {
    #pragma omp target teams distribute parallel for \
            map(to: x[0:n]) map(tofrom: y[0:n])
    for (int i = 0; i < n; i++) y[i] = a * x[i] + y[i];
}

这段代码在编译器眼里意味着：

1. 两套代码：CPU 侧要生成「分配设备内存、拷贝数据、启动内核、等待完成」的宿主代码；设备侧要生成真正做 y[i] = a*x[i] + y[i] 的内核。
2. 两层内存：x 与 y 在主机内存里有一份，在设备内存里也有一份。编译器要决定哪些数据需要拷贝、什么时候拷贝、拷贝后谁拥有最新版本。
3. 两个并行模型：CPU 侧是线程与 SIMD，GPU 侧是线程块与网格。同一段循环要映射到两套完全不同的层级结构上。
4. 同步点：内核启动是异步的，主机代码必须知道何时插入等待，才能保证后续读取看到正确数据。

一个经常被引用的经验数字是：在典型的 GPU 卸载程序里，数据搬运的时间可能占端到端时间的 50% 以上。这意味着编译器在数据搬运上的决策，其重要性不亚于内核本身的优化。这也解释了为什么现代卸载编译器把大量精力放在「减少拷贝、重叠拷贝与计算、以及让拷贝更高效」上。

2. 卸载指令与编程模型

一句话总结： 卸载编程模型分两大派：指令式（OpenMP target、OpenACC）由编译器推断映射，内核式（CUDA、HIP、SYCL）由程序员显式控制层级与内存。

2.1 指令式卸载

一句话总结： OpenMP target 与 OpenACC 用编译指示描述「这段代码在哪里跑、哪些数据要搬」，编译器负责生成宿主胶水代码与设备内核。

// OpenMP target：把循环卸载到设备
void vec_add(double *a, double *b, double *c, int n) {
    #pragma omp target teams distribute parallel for \
            num_teams(108) thread_limit(256) \
            map(to: a[0:n], b[0:n]) map(from: c[0:n])
    for (int i = 0; i < n; i++) c[i] = a[i] + b[i];
}

这段代码里的三个关键部分：

• target 表示「这段区域在设备上执行」。
• teams distribute parallel for 描述设备的层级映射：teams 对应 GPU 的线程块网格，parallel for 对应块内的线程并行。
• map(to: ...)、map(from: ...)、map(tofrom: ...) 描述数据方向：只拷进去、只拷回来、双向。

// OpenACC 的等价写法：kernels 让编译器自己推断并行与数据
void vec_add(double *a, double *b, double *c, int n) {
    #pragma acc kernels
    for (int i = 0; i < n; i++) c[i] = a[i] + b[i];
}

kernels 与 parallel 的区别值得注意：kernels 把并行化决策交给编译器（它可能选择不并行），parallel 则强制并行化。前者更安全，后者更可预测。

# 编译 OpenMP target 卸载程序并查看生成的设备代码
gcc -fopenmp -foffload=nvptx-none -O3 saxpy.c -o saxpy
clang -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -O3 saxpy.c -o saxpy

2.2 内核式卸载

一句话总结： CUDA 与 SYCL 让程序员显式写出网格维度、共享内存与拷贝方向，控制力最强，代价是代码与硬件绑定。

// CUDA 的等价实现：显式描述层级与内存
__global__ void vec_add_kernel(const double *a, const double *b, double *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // 全局线程索引
    if (i < n) c[i] = a[i] + b[i];
}

void vec_add(double *a, double *b, double *c, int n) {
    double *da, *db, *dc;
    cudaMalloc(&da, n * 8); cudaMalloc(&db, n * 8); cudaMalloc(&dc, n * 8);
    cudaMemcpy(da, a, n * 8, cudaMemcpyHostToDevice);   // 显式拷入
    cudaMemcpy(db, b, n * 8, cudaMemcpyHostToDevice);
    vec_add_kernel<<<(n + 255) / 256, 256>>>(da, db, dc, n);
    cudaMemcpy(c, dc, n * 8, cudaMemcpyDeviceToHost);   // 显式拷回
    cudaFree(da); cudaFree(db); cudaFree(dc);
}

// SYCL 用 C++ 抽象描述同样的东西，可在多后端间移植
void vec_add(queue &q, double *a, double *b, double *c, int n) {
    buffer<double, 1> ba(a, range<1>(n)), bb(b, range<1>(n)), bc(c, range<1>(n));
    q.submit([&](handler &h) {
        accessor A(ba, h, read_only), B(bb, h, read_only), C(bc, h, write_only);
        h.parallel_for(range<1>(n), [=](id<1> i) { C[i] = A[i] + B[i]; });
    });   // buffer 析构时自动拷回主机
}

工程上的选择往往是用指令式写可移植的部分，用内核式写热点。这也是 OpenMP 5.x 引入 dispatch 与 interop 的原因：允许在同一个程序里混用两种模型。

3. 并行模型映射

一句话总结： 编译器要把源语言的循环结构映射到设备的执行层级上，映射质量决定了并行度、占用率与访存效率。

3.1 循环到层级的映射

一句话总结： 嵌套循环通常映射为「外层循环对应线程块、内层循环对应线程」的结构，编译器要检查依赖、选择分块大小、并决定是否需要同步。

GPU 的执行层级是三层：网格（grid）包含多个线程块（block），每个块包含多个线程（thread），线程以 32 个为一组构成 warp（NVIDIA）或 wavefront（AMD）一起执行。

// 源：for (i) for (j) A[i][j] = f(B[i][j]);
// 外层循环映射为 grid，内层循环映射为 block 内的线程
__global__ void kernel(double *A, double *B, int M, int N) {
    int i = blockIdx.x, j = threadIdx.x;   // 外层索引、内层索引
    if (i < M && j < N) A[i * N + j] = f(B[i * N + j]);
}

映射的合法性依赖于依赖分析：如果内层循环存在跨迭代依赖（比如 A[j] = A[j-1] + 1），直接映射为线程会导致数据竞争。编译器必须证明无依赖，或者插入同步。

# 卸载合法性检查的简化逻辑
def can_offload(nest, deps):
    for d in deps:
        if d.kind in ("RAW", "WAR", "WAW") and d.carried_by(nest.outer):
            return False, "外层循环存在跨迭代依赖，不能映射为线程块"
        if d.kind in ("RAW", "WAR", "WAW") and d.carried_by(nest.inner):
            return True, "需块内同步或私有化"     # 内层依赖可局部解决
    return True, "无依赖，可自由映射"

// 通过私有化解决内层依赖：把原地更新改写成读一份、写另一份
// 源（有依赖）：     for (j) { tmp = A[j]; A[j] = A[j-1] + tmp; }
// 私有化后（无依赖）：for (j) { double tmp = B[j]; C[j] = B[j-1] + tmp; }

3.2 分块与共享内存

一句话总结： 把循环分块并把块内复用的数据搬进共享内存，是 GPU 上最有效的访存优化之一，编译器能否自动做这件事取决于依赖分析与代价模型。

GPU 的内存层级里，共享内存（shared memory）是块内线程共享的低延迟暂存区，延迟比全局内存低一到两个数量级。矩阵乘法的分块是最经典的例子：

// 分块矩阵乘法：把 A、B 的一块搬进共享内存复用
#define TILE 16
__global__ void matmul_tiled(const float *A, const float *B, float *C, int N) {
    __shared__ float As[TILE][TILE], Bs[TILE][TILE];
    int row = blockIdx.y * TILE + threadIdx.y, col = blockIdx.x * TILE + threadIdx.x;
    float acc = 0.0f;
    for (int t = 0; t < N / TILE; t++) {
        As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col];
        __syncthreads();                 // 等所有线程搬完
        for (int k = 0; k < TILE; k++) acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();                 // 等用完，避免下一轮覆盖
    }
    C[row * N + col] = acc;
}

分块把全局访存从 O(n^3) 降到 O(n^3 / tile)，收益随 tile 增大，但共享内存容量（典型 48KB）给出了上限：

代价模型要同时权衡两件事：tile 越大访存越少，但 2 * tile * tile * 4 字节的共享内存占用也越大，超出 48KB 就不可行。自动做分块对编译器来说很难：它需要判断哪些数据会被块内复用、复用距离有多远、共享内存容量是否够、以及同步点的插入位置。生产级编译器（如 Polly、TVM、Triton）通常在特定模式上做这件事，而不是通用地做。

4. 内存层级与数据搬运

一句话总结： 数据在主机与设备之间有独立地址空间时，编译器必须插入显式拷贝；统一内存用页错误按需迁移，简化了编程但把开销藏在了运行时。

4.1 显式拷贝与统一内存

一句话总结： 显式拷贝让程序员精确控制搬运时机与方向，统一内存让硬件在访问时按页迁移，前者可预测、后者易用。

// 显式拷贝：方向明确，时机可控
void explicit_copy(float *h_x, int n) {
    float *d_x;
    cudaMalloc(&d_x, n * 4);
    cudaMemcpy(d_x, h_x, n * 4, cudaMemcpyHostToDevice);  // 拷贝进
    kernel<<<n / 256, 256>>>(d_x, n);
    cudaMemcpy(h_x, d_x, n * 4, cudaMemcpyDeviceToHost);  // 拷贝回
    cudaFree(d_x);
}

// 统一内存：一次分配，双方都能访问，硬件按需迁移
void unified_memory(float *h_x, int n) {
    float *u_x; cudaMallocManaged(&u_x, n * 4);   // 统一地址空间
    memcpy(u_x, h_x, n * 4);                      // 主机写
    kernel<<<n / 256, 256>>>(u_x, n);             // 设备读，触发页迁移
    cudaDeviceSynchronize();
    memcpy(h_x, u_x, n * 4);                      // 主机读，页迁移回
}

用 nsys profile --stats=true ./app 可以观察 kernel 时间与 memcpy 时间的占比；若 memcpy 超过 30%，优先考虑减少拷贝或改用固定内存。

统一内存的陷阱在于页错误：设备访问一个还在主机内存的页时，会触发一次页错误把页迁移过来。如果访问模式随机且分散，页会来回迁移，性能可能比显式拷贝差好几倍。缓解手段是预取（prefetch）：

// 用预取把页迁移提前，避免运行中频繁页错误
cudaMemPrefetchAsync(u_x, n * 4, device_id, stream);
kernel<<<grid, block, 0, stream>>>(u_x, n);
cudaMemPrefetchAsync(u_x, n * 4, cudaCpuDeviceId, stream);

4.2 拷贝与计算的重叠

一句话总结： 把数据分成多块、用多个流分别拷贝与计算，可以让 PCIe 传输与内核执行并行，端到端时间接近两者中较大的那个。

// 分块流水线：拷贝第 i+1 块的同时计算第 i 块（主机内存须 pinned，否则退化为同步）
for (int i = 0; i < nchunks; i++) {
    int off = i * chunk, s = i % nstreams;
    cudaMemcpyAsync(d_a + off, h_a + off, chunk * 4, cudaMemcpyHostToDevice, stream[s]);
    kernel<<<(chunk + 255) / 256, 256, 0, stream[s]>>>(d_a + off, chunk);
    cudaMemcpyAsync(h_out + off, d_out + off, chunk * 4, cudaMemcpyDeviceToHost, stream[s]);
}

5. 主机与设备协同

一句话总结： 卸载后的程序由宿主代码与设备内核交替执行，编译器要正确插入异步启动、事件依赖与等待，才能在保持正确性的同时获得重叠收益。

5.1 异步与同步

一句话总结： 内核启动默认异步，宿主代码必须显式等待或通过事件建立依赖，否则会出现「读到旧数据」或「拷贝未完成就计算」的错误。

// 同步粒度：cudaDeviceSynchronize（全部）/ cudaStreamSynchronize（单流）/ cudaEventSynchronize（单事件）
kernel<<<grid, block>>>(d_x);   // 内核启动默认异步，立即返回
// 跨流依赖：stream2 等 stream1 记录的事件 done，之后才能读 d_x
cudaEvent_t done; cudaEventCreate(&done); cudaEventRecord(done, stream1);
cudaStreamWaitEvent(stream2, done, 0);
kernel_b<<<grid, block, 0, stream2>>>(d_x);

宿主代码的生成是模板化的：device_alloc 分配、device_memcpy_to 拷入、<<<grid, block>>> 启动内核、device_memcpy_from 拷回、device_free 释放。编译器为每段卸载区域套一遍这个模板，再按需插入事件与等待。

5.2 卸载粒度的选择

// 循环级卸载：每轮都拷进拷出，数据在主机与设备之间来回搬
for (int t = 0; t < T; t++) {
    #pragma omp target map(tofrom: a[0:n])
    for (int i = 0; i < n; i++) a[i] += 1.0;
}

// 区域级卸载：target data 建立设备数据环境，内层 target 复用映射
#pragma omp target data map(tofrom: a[0:n])
{
    for (int t = 0; t < T; t++) {
        #pragma omp target
        for (int i = 0; i < n; i++) a[i] += 1.0;   // 无拷贝
    }
}

这个例子说明了卸载编译里最重要的一个工程原则：把 target data 提到循环外面。它带来的性能差异常常是数倍——因为省掉的是每轮迭代的整段数组拷贝，而不是几条指令。

6. 编译器实现要点

一句话总结： 卸载编译器的实现分三段：前端识别可卸载区域并收集数据映射，中端做设备无关的并行化与优化，后端为设备生成 SIMT 代码。

; LLVM IR 里的卸载表示（简化示意）：卸载区域被包装成一个函数
define void @saxpy_offload_region(ptr %x, ptr %y, float %a, i32 %n) {
  %i = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()   ; 内核体：并行化后的循环
  ret void
}
; 宿主侧调用被替换为卸载描述 llvm.openmp.target(ptr @saxpy_offload_region, ...)

设备后端的寄存器分配与 CPU 有很大不同。GPU 上寄存器数量是占用率（occupancy）的决定因素：每个线程用的寄存器越多，一个 SM 上能同时驻留的线程就越少，延迟隐藏能力就越差。因此设备编译器常在寄存器压力与占用率之间做取舍，甚至故意溢出到本地内存来降低寄存器数。

// 占用率与寄存器数的关系（每个 SM 有 65536 个寄存器）
// 每线程 32 个寄存器：65536 / 32 = 2048 线程（通常受上限 1536 限制）
// 每线程 128 个寄存器：65536 / 128 = 512 线程
// 后者占用率只有前者的 1/3，延迟隐藏能力显著下降

用 nvcc -Xptxas -v 可以打印每个内核的寄存器数与共享内存用量（形如 Used 28 registers, 512 bytes smem），ncu --set full 则给出实际占用率与访存效率。

7. 性能调优与陷阱

一句话总结： 异构程序的性能瓶颈通常不在内核指令数，而在数据搬运、访存模式、分支发散与占用率，调优要先用工具定位再动手。

// 非合并访存 vs 合并访存
// 差：线程 0 访问 0，线程 1 访问 N，事务分散
__global__ void bad(float *a, int N) { a[threadIdx.x * N] = 1.0f; }
// 好：相邻线程访问相邻地址，一个 warp 的事务合并成少数几次
__global__ void good(float *a) {
    a[blockIdx.x * blockDim.x + threadIdx.x] = 1.0f;
}

// 分支发散：warp 内一半线程走 if，一半走 else，两条路径都要执行
__global__ void divergent(const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    y[i] = x[i] > 0 ? sqrtf(x[i]) : 0.0f;   // 只有部分线程真正需要 sqrt
}
// 谓词化：把条件选择改写成无分支表达式，全 warp 走同一路径
__global__ void predicated(const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    y[i] = sqrtf(fmaxf(x[i], 0.0f)) * (x[i] > 0);
}

# 用 ncu 定位瓶颈：Memory/Compute Throughput 接近 100% 说明对应资源受限
# Achieved Occupancy 与理论值差距大说明延迟隐藏不足
# Warp Execution Efficiency 低于 100% 说明有分支发散
ncu --set full --launch-count 1 ./app

一个常见的误解是「把 CPU 代码直接加上 #pragma omp target 就能获得加速」。实际上，GPU 与 CPU 的性能特征差异极大：CPU 靠大缓存与乱序执行掩盖延迟，GPU 靠海量线程隐藏延迟。一段在 CPU 上跑得很快的代码（比如指针追逐、递归、频繁小分配），在 GPU 上可能慢几十倍。卸载编译器的价值在于降低尝试成本，但它不能改变算法与数据布局对硬件的适配性——这一点必须由程序员判断。

8. 总结

异构计算的编译问题，本质上是把一个程序拆成两个程序、再让它们协同。这个拆分的每一步都引入了新的决策点：哪些代码卸载、数据什么时候搬、同步点插在哪里、寄存器与占用率如何取舍。这些决策没有一个能靠局部规则完美解决，因此现代卸载编译器普遍采用「程序员给提示、编译器做推断、运行时做调度」的三层分工。理解了这三层各自能做什么，也就理解了为什么异构程序的性能优化至今仍然高度依赖人工——工具能消除低效，但无法替你判断算法是否适合加速器。至此，本专题从编译器的前端、中端、后端，一路走到运行时的动态优化与异构目标，形成了一条从源码到机器的完整链路。

延伸阅读

• 交叉编译与多目标
• LLVM 与 LLVM IR 后端
• 向量化与 SIMD
• 目标代码生成
• 运行时内存管理