AI infra：并行编程与 CUDA 编程入门
ai infra, cuda

AI infra：并行编程与 CUDA 编程入门

主要介绍并行计算与 CUDA 编程的基础概念。文章首先区分串行、并发与并行三种程序执行方式，并说明 CPU 多核并行和 GPU 大规模线程并行的差异。随后以向量加法为例，讲解 CUDA 程序从数据准备、显存申请、数据传输、核函数执行到结果拷回的完整流程，并进一步介绍 Grid、Block、Thread 线程层级、SIMT 执行模型、GPU 硬件单元以及 nvcc 编译流程。最后结合 Nsight Systems 分析程序性能，引出主机与设备间数据传输优化的重要性。

189次点击8分钟阅读

什么是并行

计算机中程序的运行方式主要有三种：

串行、并发、并行
  • 串行：同一时刻只有一个任务在执行，必须等前一个任务彻底结束，后一个任务才能开始。
  • 并发：多个任务都已经开始，但同一时刻只有一个任务在执行，往往通过轮流转换的方式来实现。
  • 并行：多个任务在同一时刻同时执行。

计算机如何并行

计算机中的运算部件可以并行的主要有两种：CPU 和 GPU。

CPU与GPU对比

CPU 的并行主要通过多核来执行。一台现代的 CPU 通常有多个核心，每个核心都能独立执行指令，这样可以把不同的任务分配到不同的核心上同时运行。如操作系统、Web 服务等。

GPU 的并行是通过将一个大任务拆分成海量相似的小任务，让许多计算单元同时做同样模式的工作。GPU 的内部拥有大量更轻量的执行单元，通过牺牲单个线程的灵活性换来极强的批量吞吐能力。

CUDA 编程

  • C/C++语法：宿主语言为 C/C++，使用范围广，应用领域对口，新语法的学习门槛较低。
  • 与 CPU 端协作：CPU 负责整理程序执行和处理逻辑等。
  • SIMT 模式：Single Instruction, Multiple Threads，只需要编写一个指令，一个指令可同时被多个线程执行。
  • 自动调度：根据设定的执行参数自动调度资源（优化）。

GPU 硬件单元

GPU 的硬件单元指的是流式多处理器（Streaming Multiprocessor， SM），是 GPU 内部的核心工作单元，GPU 会将所有的线程分配到一个个 SM 上执行。其中单个 SM 中又主要有 CUDA Core 和 Tensor Core：

SM

CUDA Core 通常处理的是一个线程的一条普通算术指令，如整数运算、浮点运算、加减乘除、逻辑判断、地址计算以及一些通用指令等，而 Tensor Core 主要用于专门加速矩阵乘加操作的，其最典型的操作即：

D=A×B+CD = A \times B + C

也就是矩阵乘法再加累加。这类操作在深度学习里极其常见，因为全连接层本质上是矩阵乘法，卷积在很多实现里也会转成矩阵乘法，Transformer 里的 attention、MLP 也高度依赖 GEMM。

CUDA 编程运行流程

CUDA代码执行流程

CPU 首先进行程序初始化、数据准备等工作，而后将数据从主机内存拷贝至设备全局内存。而后 SM 读取全局内存中的数据，并行执行计算任务，再将结果写回全局内存，最后将结果从设备全局内存拷贝回主机内存。

CPU 加法

CPU 加法的过程主要如下：准备和初始化数据，定义加法函数，依靠循环来进行所有的元素加法，调用函数后验证结果。

CPU加法
#include <stdio.h>
#include <stdlib.h>
#include <vector>

using namespace std;

void add_cpu(vector<float>& c, const vector<float>& a, const vector<float>& b){
    for(size_t i = 0; i < c.size(); i++){
        c[i] = a[i] + b[i];
    }
}

int main(){
    const size_t SIZE = 1 << 20;
    std::vector<float> a(SIZE, 1.0f);
    std::vector<float> b(SIZE, 2.0f);
    std::vector<float> c(SIZE, 0.0f);
    add_cpu(c, a, b);
    return 0;
}

CPU 加法→GPU 加法

CPU先准备和初始化数据：

GPU加法-准备数据
const size_t SIZE = 1 << 20; // 1M elements
size_t bytes = SIZE * sizeof(float);

std::vector<float> h_a(SIZE, 1.0f);
std::vector<float> h_b(SIZE, 2.0f);
std::vector<float> h_c(SIZE, 0.0f);

然后将数据传输到 GPU，类似于 C 语言中的申请堆内存，GPU 中的内存也需要先申请再使用：

申请GPU内存
float *d_a, *d_b, *d_c;
CUDA_CHECK(cudaMalloc((void**)&d_a, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_b, bytes));
CUDA_CHECK(cudaMalloc((void**)&d_c, bytes));

内存申请成功后传输数据：

传输数据至设备
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));

而后 GPU 从全局内存中读取并计算后写回。这里的计算是要写在 GPU 上运行的加法函数并进行调用，在 CPU 的加法运算中使用了循环，而在 GPU 中要使用的是 SIMT，SIMT 指挥每个线程，需要组织结构和编号。在 CUDA 中主要是 Grid→Block→Thread 的层级关系。

  • Thread 是最小的执行单位，即一个线程，通常负责处理一小份数据，每个 Thread 中通常有自己的局部变量、寄存器上下文、程序计数状态和 threadIdx。
  • Block 是一组 Thread 的集合，CUDA 不会让单个线程独自运行，而是将线程先组织成 Block，一个 Block 的线程中有几个比较重要的特性：它们可以通过 shared memory 共享数据，可以用 __syncthreads() 做块内同步，还会被分配到同一个 SM 上执行，SM 上可以同时驻留多个 Block。不同的 Block 之间一般不直接同步。Block 还有自己的索引：blockIdx.x,blockIdx.y,blockIdx.z，每个 Block 内部的线程也有自己的局部编号：threadIdx.x,threadIdx.y,threadIdx.z。Block 的尺寸在进行 kernel lauch 时指定，如 dim3 blockDim(256) 和 dim3 blockDim(16,16) 都是表示每个 block 有 256 个线程。
  • Grid 是一次 kernel 启动时，所有 block 的总集合。即每调用一次 kernel<<gridDim, blockDim>>(...); 里面的 gridDim 表示 Block 的数量，blockDim 表示每个 Block 中有多少个 Thread。

因此就需要定义 Block 的数量和大小来指挥线程并行计算，还要定义 GPU 上的加法函数，结合定义的信息调用 GPU 加法函数。

启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;
    // 启动 kernel
    vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, SIZE);
    // 检查 kernel 启动错误
    CUDA_CHECK(cudaGetLastError());
    // 等待 GPU 执行完成
    CUDA_CHECK(cudaDeviceSynchronize());

其中核函数定义为：

kernel函数定义
__global__ void vec_add(const float* a, const float* b, float* c, size_t n) {
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

其中的 dim3 是 CUDA 表示线程层级结构的类型，<<>> 中的内容传递线程层级信息给核函数，核函数（kernel）即设备侧的入口函数，__global__ 表示这是一个核函数，最后将结果拷回主机：

拷贝结果至主机
// 把结果拷回主机
    CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));
    // 简单验证
    for (size_t i = 0; i < 10; ++i) {
        printf("h_c[%zu] = %f\n", i, h_c[i]);
    }
    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));
    return 0;

完整的代码如下：

CUDA并行加法完整代码
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <vector>

#define CUDA_CHECK(call)                                                       \
    do {                                                                       \
        cudaError_t err__ = (call);                                            \
        if (err__ != cudaSuccess) {                                            \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__,   \
                    cudaGetErrorString(err__));                                \
            exit(EXIT_FAILURE);                                                \
        }                                                                      \
    } while (0)

// GPU kernel：逐元素相加
__global__ void vec_add(const float* a, const float* b, float* c, size_t n) {
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    const size_t SIZE = 1 << 20; // 1M elements
    size_t bytes = SIZE * sizeof(float);

    std::vector<float> h_a(SIZE, 1.0f);
    std::vector<float> h_b(SIZE, 2.0f);
    std::vector<float> h_c(SIZE, 0.0f);

    float *d_a, *d_b, *d_c;
    CUDA_CHECK(cudaMalloc((void**)&d_a, bytes));
    CUDA_CHECK(cudaMalloc((void**)&d_b, bytes));
    CUDA_CHECK(cudaMalloc((void**)&d_c, bytes));

    CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));

    // 配置 kernel 启动参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;

    // 启动 kernel
    vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, SIZE);

    // 检查 kernel 启动错误
    CUDA_CHECK(cudaGetLastError());

    // 等待 GPU 执行完成
    CUDA_CHECK(cudaDeviceSynchronize());

    // 把结果拷回主机
    CUDA_CHECK(cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost));

    // 简单验证
    for (size_t i = 0; i < 10; ++i) {
        printf("h_c[%zu] = %f\n", i, h_c[i]);
    }

    CUDA_CHECK(cudaFree(d_a));
    CUDA_CHECK(cudaFree(d_b));
    CUDA_CHECK(cudaFree(d_c));

    return 0;
}

最后通过以下命令编译运行即可：

nvcc编译运行命令
nvcc add_cuda.cu -o add_cuda
./add_cuda

nvcc 是 CUDA 的编译器，是 CUDA Toolkit 的一部分。其中编译的流程大致如下：

CUDA代码编译流程

nvcc 将同一份 .cu 源文件拆分成两条路径来处理：一条是 host 路径，用于生成 CPU 侧的目标代码，另一条是 device 路径，用于生成 GPU 侧的目标代码。具体而言，nvcc 会先对源文件中的 device 部分进行处理，生成 PTX 或 cubin 这类设备端中间结果，并将其封装进 fatbinary；随后再对 host 部分重新预处理，把 CUDA 的扩展语法转换成宿主编译器能够理解的标准 C/C++ 形式，同时把前面生成的 fatbinary 嵌入进去，最后再交给 g++、clang++ 或 MSVC 这样的 host compiler 生成 host object。

Nsight Systems（CLI）

英伟达提供了一个系统级性能分析工具 Nsight Systems，便于对编写的 CUDA 程序进行性能评估。

nsys性能评估
nsys profile -t cuda,nvtx,osrt -o add_cuda -f true ./add_cuda
nsys stats add_cuda.nsys-rep

通过以上命令就能够对运行的程序进行性能评估。观察得到的结果是 HtoD 和 DtoH 的耗时远远大于核函数的时间。设定的指标为：

总耗时=TH2D+Tkernel+TD2H总耗时= T_{H 2 D} + T_{kernel} + T_{D 2H}

要进一步优化，就需要从 HtoD 和 DtoH 这两个时间入手。

设备信息

一般而言，启用越多的线程，就能获得越大的并行度，越大的并行度也就能获取越好的性能。一个设备最多能够启用的线程数量可以通过网络/官网查询获取，也可以通过代码获取，NVIDIA 提供了 cudaDeviceProp 来通过代码查询。

cudaDeviceProp获取设备信息
#include <iostream>
#include <cuda_runtime.h>

int main() {
    int device_id = 0;                 // 想查询的 GPU 编号
    cudaDeviceProp prop;

    cudaError_t err = cudaGetDeviceProperties(&prop, device_id);
    if (err != cudaSuccess) {
        std::cerr << "cudaGetDeviceProperties failed: "
                  << cudaGetErrorString(err) << std::endl;
        return 1;
    }

    std::cout << "Device name: " << prop.name << std::endl;
    std::cout << "Total global memory: " << prop.totalGlobalMem << std::endl;
    std::cout << "SM count: " << prop.multiProcessorCount << std::endl;
    std::cout << "Max threads per block: " << prop.maxThreadsPerBlock << std::endl;
    std::cout << "Max threads per SM: " << prop.maxThreadsPerMultiProcessor << std::endl;
    std::cout << "Max threads dim: (" << prop.maxThreadsDim[0] << ", "
              << prop.maxThreadsDim[1] << ", " << prop.maxThreadsDim[2] << ")"
              << std::endl;
    std::cout << "Compute capability: "
              << prop.major << "." << prop.minor << std::endl;

    return 0;
}

相关文章

AI infra：内存模型与规约优化

AI infra：内存模型与规约优化

2026/05/20ai infra13618分钟阅读
AI infra：性能模型与逐元素优化

AI infra：性能模型与逐元素优化

2026/05/19ai infra13412分钟阅读