技术文档
GPU 编译器技术

GPU 编译器技术

GPU 编译器技术

GPU 编译器是现代并行计算的核心技术，负责将高级语言代码转换为GPU可执行的机器码。本章节深入探讨CUDA、OpenCL、ROCm等主流GPU编程模型的编译技术。

📋 章节概览

GPU编译器简介
并行计算编译基础
CUDA 编译器
NVIDIA GPU 编译技术
OpenCL 编译器
跨平台并行计算
ROCm 编译器
AMD GPU 编译技术
GPU 优化
性能优化技术
内存管理
GPU 内存层次结构
性能分析
性能调优工具

🎯 学习目标

通过本章节的学习，你将掌握：

  • GPU 架构理解: 深入了解现代GPU的硬件架构
  • 编译流程: 掌握GPU代码的完整编译过程
  • 性能优化: 学会GPU程序的性能分析和优化
  • 内存管理: 理解GPU内存层次和优化策略
  • 并行模式: 掌握各种GPU并行编程模式

🏗️ GPU 编译器架构

CUDA 编译流程 

CUDA C/C++  →  [NVCC]  →  PTX  →  [Driver]  →  SASS
     ↓              ↓           ↓              ↓
  源代码        前端编译    中间表示      机器码

OpenCL 编译流程 

OpenCL C  →  [Clang]  →  SPIR-V  →  [Runtime]  →  ISA
    ↓            ↓          ↓            ↓         ↓
  内核代码    前端编译   中间表示    运行时编译   机器码

ROCm 编译流程 

HIP/OpenCL  →  [HCC/Clang]  →  AMDGCN  →  [ROCm]  →  GCN ISA
     ↓              ↓            ↓          ↓          ↓
   源代码        前端编译     LLVM IR   后端编译    机器码

🔧 开发环境搭建

CUDA 开发环境 

# 安装 CUDA Toolkit
wget https://developer.download.nvidia.com/compute/cuda/12.0.0/local_installers/cuda_12.0.0_525.60.13_linux.run
sudo sh cuda_12.0.0_525.60.13_linux.run

# 设置环境变量
export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH

# 验证安装
nvcc --version
nvidia-smi

OpenCL 开发环境 

# 安装 OpenCL 头文件和库
sudo apt-get install opencl-headers ocl-icd-opencl-dev

# 安装供应商特定的 ICD
# NVIDIA
sudo apt-get install nvidia-opencl-icd
# AMD
sudo apt-get install mesa-opencl-icd
# Intel
sudo apt-get install intel-opencl-icd

# 验证安装
clinfo

ROCm 开发环境 

# 添加 ROCm 仓库
wget -qO - https://repo.radeon.com/rocm/rocm.gpg.key | sudo apt-key add -
echo 'deb [arch=amd64] https://repo.radeon.com/rocm/apt/debian/ ubuntu main' | sudo tee /etc/apt/sources.list.d/rocm.list

# 安装 ROCm
sudo apt update
sudo apt install rocm-dkms

# 设置环境变量
export PATH=/opt/rocm/bin:$PATH
export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH

# 验证安装
rocminfo
hipcc --version

💻 第一个GPU程序

CUDA 示例 

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const int N = 1024;
    size_t size = N * sizeof(float);
    
    // 分配主机内存
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    
    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = i * 2;
    }
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);
    
    // 拷贝数据到设备
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 启动内核
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N);
    
    // 拷贝结果回主机
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %.2f\n", i, h_c[i]);
    }
    
    // 清理内存
    free(h_a); free(h_b); free(h_c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    
    return 0;
}

编译和运行 

# 编译 CUDA 程序
nvcc -o vector_add vector_add.cu

# 运行程序
./vector_add

# 查看 PTX 中间代码
nvcc -ptx vector_add.cu
cat vector_add.ptx

# 查看 SASS 汇编代码
nvcc -cubin vector_add.cu
cuobjdump -sass vector_add.cubin

🚀 GPU 架构特点

SIMT 执行模型 

Warp (32 threads)
├── Thread 0  ┐
├── Thread 1  │ 执行相同指令
├── Thread 2  │ (SIMT)
├── ...       │
└── Thread 31 ┘

内存层次结构 

全局内存 (Global Memory)     - 大容量，高延迟
    ↑
共享内存 (Shared Memory)     - 小容量，低延迟
    ↑  
寄存器 (Registers)          - 最快访问
    ↑
常量内存 (Constant Memory)   - 只读，有缓存
纹理内存 (Texture Memory)   - 只读，空间局部性优化

📊 性能分析工具

NVIDIA 工具链 

# Nsight Compute - 内核性能分析
ncu --set full ./vector_add

# Nsight Systems - 系统级性能分析  
nsys profile ./vector_add

# nvprof - 传统性能分析工具
nvprof ./vector_add

AMD 工具链 

# ROCProfiler - AMD GPU 性能分析
rocprof ./hip_program

# ROCTracer - API 调用跟踪
roctracer ./hip_program

🔗 相关资源