Select a result to preview

CUDA1_CUDA编程基础

CUDA（一）：CUDA 编程基础

文章来源：知乎专栏 - CUDA（一）：CUDA 编程基础
文章ID：645330027
发布时间：编辑于 2025-07-08 14:25・上海

一、GPU 与 CUDA 架构深度解析

1.1 GPU 与 CPU 的设计哲学差异

处理器的两个核心指标是延迟（Latency）和吞吐量（Throughput）：

CPU：延迟导向设计

CPU 的架构特点：

  1. 多级高速缓存结构：提升指令访问存储速度
  2. 复杂控制单元：
  1. 强大的运算单元（Core）：整型和浮点型复杂运算速度快

设计目标：减少单条指令的执行延迟

GPU：吞吐导向设计

/

性能对比与适用场景

性能差异：

GPU 适用场景：

  1. 计算密集型任务：数值计算比例远大于内存操作，内存访问延时可被计算掩盖
  2. 数据并行任务：大任务可拆解为执行相同指令的小任务，对复杂流程控制需求较低

1.2 CUDA 硬件架构层次

CUDA（Compute Unified Device Architecture）是支持 GPU 通用计算的平台和编程模型，提供 C/C++ 语言扩展和用于编程管理 GPU 的 API。

硬件层次结构（从底层到顶层）

1. SP（Stream Processor，线程处理器）

2. SM（Streaming Multiprocessor，多核处理器）

3. GPU（Device）

内存层次总结

从线程视角看内存访问权限：

1.3 CUDA 软件抽象模型

硬件与软件的对应关系：

Thread Block（线程块）特性

线程块是软件侧的基本执行单位，具有以下特点：

二、CUDA 编程核心要素

2.1 Kernel 函数与线程组织

Kernel 定义：

线程索引与线程 ID 的关系

2.2 矩阵加法示例：单 Block 版本

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

关键点：

2.3 矩阵加法示例：多 Block 版本

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

关键改进：

内置变量：

2.4 线程同步机制

**__syncthreads() 函数**：

重要限制：

2.5 Warp（线程束）：GPU 执行的最小单元

SIMT 架构（Single-Instruction, Multiple-Thread）

核心概念：

Warp 定义：

实践要点

关键优化建议：

Kernel 执行流程

当一个 kernel 被执行时：

  1. Grid 中的线程块被分配到 SM 上
  2. 一个线程块的 thread 只能在一个 SM 上调度
  3. 一个 SM 通常可以调度多个线程块
  4. 大量 thread 可能被分配到不同的 SM 上
  5. 每个 thread 拥有独立的程序计数器和状态寄存器
  6. 每个 thread 用自己的数据执行指令（SIMT）

硬件限制：

文章概要

本文是关于CUDA（一）：CUDA 编程基础的技术文章，属于高性能计算与AI基础设施系列。

核心内容

一、GPU 与 CUDA 结构

1.1 GPU 再认识

对于处理器而言，有2个指标是最主要的：延迟 和吞吐量
。延迟，是指从发出指令到最终返回结果中间经历的时间间隔。而吞吐量，就是单位之间内处理的指令的条数。下面将主要从这两个方面来比较GPU和CPU：

下图左是 CPU 的示意图，有以下几个特点：

  1. CPU 中包含了多级高速的缓存结构。 这样提升了指令访问存储的速度。
  2. CPU 中包含了很多控制单元。 具体有2种，一个是分支预测机制，另一个是流水线前传机制。
  3. CPU 的运算单元 (Core) 强大，整型浮点型复杂运算速度快。

基于以上三点，CPU 在设计时的导向就是减少指令的时延，被称之为延迟导向设计 。

下图右是 GPU 的示意图，有以下几个特点：

  1. GPU 中虽有缓存结构但是数量少。 因为要减少指令访问缓存的次数。
  2. GPU 中控制单元非常简单。 控制单元中没有分支预测机制和数据转发机制，对于复杂的指令运算就会比较慢。
  3. GPU 的运算单元 (Core) 非常多，采用长延时流水线以实现高吞吐量。 每一行的运算单元的控制器只有一个，意味着每一行的运算单元使用的指令是相同的，不同的是它们的数据内容。那么这种整齐划一的运算方式使得 GPU 对于那些控制简单但运算高效的指令的效率显著增加。

基于此，可以看到 GPU 在设计过程中以一个原则为核心：增加简单指令的吞吐，这称 GPU 为吞吐导向设计。

GPU vs CPU

由于设计原则不同，二者擅长的场景有所不同：

进一步可以具体化适合 GPU 的场景：

  1. 计算密集：数值计算的比例要远大于内存操作，因此内存访问的延时可以被计算掩盖。
  2. 数据并行：大任务可以拆解为执行相同指令的小任务，因此对复杂流程控制的需求较低。

2.2 CUDA 结构

CUDA (Compute Unified Device Architecture)是支持 GPU 通用计算的平台和编程模型，提供 C/C++ 语言扩展和
用于编程和管理 GPU的API。

从硬件的角度来讲，CUDA 内存模型的最基本的单位就是
[SP](https://zhida.zhihu.com/search?content_id=231537480&content_type=Article&match_order=1&q=SP&zd_token=eyJhbGciOiJIUzI1NiIsInR5cCI6IkpXVCJ9.eyJpc3MiOiJ6aGlkYV9zZXJ2ZXIiLCJleHAiOjE3NzgxMzY4NzAsInEiOiJTUCIsInpoaWRhX3NvdXJjZSI6ImVudGl0eSIsImNvbnRlbnRfaWQiOjIzMTUzNzQ4MCwiY29udGVudF90eXBlIjoiQXJ0aWNsZSIsIm1hdGNoX29yZGVyIjoxLCJ6ZF90b2tlbiI6bnVsbH0.jqrEBLK7YokxKnurB6qTFjSv-
wXi8ZfsI-99O_rvIhs&zhida_source=entity) (线程处理器)。每个线程处理器 (SP) 都用自己的
registers (寄存器) 和 local memory (局部内存)
。寄存器和局部内存只能被自己访问，不同的线程处理器之间是彼此独立的。

由多个线程处理器 (SP) 和一块共享内存所构成的就是 SM (多核处理器)
(灰色部分)。多核处理器里边的多个线程处理器是互相并行的，是不互相影响的。每个多核处理器 (SM) 内都有自己的 shared memory
(共享内存)，shared memory 可以被线程块内所有线程访问。

再往上，由这个 SM (多核处理器) 和一块全局内存，就构成了 GPU。一个 GPU 的所有 SM 共有一块 global memory
(全局内存)，不同线程块的线程都可使用。

上面这段话可以表述为：每个 thread 都有自己的一份 register 和 local memory 的空间。同一个 block 中的每个 thread
则有共享的一份 sha

三、实践：PyTorch 自定义 CUDA 算子

本节实现一个简单的 CUDA 算子（两个 n×n tensor 相加），并通过 PyTorch 调用。

3.1 算子设计与实现

设计规格

索引映射关键

每个线程需要从线程索引映射到全局线性内存索引：

CUDA 实现代码

__global__ void MatAdd(float* c,
                       const float* a,
                       const float* b,
                       int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    int idx = j*n + i;
    if (i < n && j < n)
        c[idx] = a[idx] + b[idx];
}

void launch_add2(float* c,
                 const float* a,
                 const float* b,
                 int n) {
    dim3 block(16, 16);
    dim3 grid(n/block.x, n/block.y);
    
    MatAdd<<<grid, block>>>(c, a, b, n);
}

关键点：

3.2 Torch C++ 封装

CUDA kernel 函数 PyTorch 不能直接调用，需要提供接口（add2_ops.cpp）：

#include <torch/extension.h>
#include "add2.h"

void torch_launch_add2(torch::Tensor &c,
                       const torch::Tensor &a,
                       const torch::Tensor &b,
                       int64_t n) {
    launch_add2((float *)c.data_ptr(),
                (const float *)a.data_ptr(),
                (const float *)b.data_ptr(),
                n);
}

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
    m.def("torch_launch_add2",
          &torch_launch_add2,
          "add2 kernel warpper");
}

TORCH_LIBRARY(add2, m) {
    m.def("torch_launch_add2", torch_launch_add2);
}

工作流程：

  1. torch_launch_add2 接收 C++ 版本的 torch tensor
  2. 转换为 C++ 指针数组
  3. 调用 CUDA 函数 launch_add2 执行核函数
  4. 使用 pybind11 封装，生成 Python 可调用的 .so 库

Torch 使用 CUDA 算子的三步骤

  1. 编写 CUDA 算子和对应的调用函数
  2. 编写 torch cpp 函数建立 PyTorch 和 CUDA 之间的联系，用 pybind11 封装
  3. 用 PyTorch 的 cpp 扩展库进行编译和调用

3.3 三种编译调用方法

方法一：JIT（Just-In-Time）编译

特点：Python 代码运行时即时编译 cpp 和 cuda 文件

from torch.utils.cpp_extension import load
cuda_module = load(name="add2",
                   extra_include_paths=["include"],
                   sources=["kernel/add2_ops.cpp", "kernel/add2_kernel.cu"],
                   verbose=True)

cuda_module.torch_launch_add2(cuda_c, a, b, n)

执行命令：

python run_time.py --compiler jit

性能结果：

方法二：SETUP 编译

特点：通过 Setuptools 预编译

from setuptools import setup
from torch.utils.cpp_extension import BuildExtension, CUDAExtension

setup(
    name="add2",
    include_dirs=["include"],
    ext_modules=[
        CUDAExtension(
            "add2",
            ["kernel/add2_ops.cpp", "kernel/add2_kernel.cu"],
        )
    ],
    cmdclass={
        "build_ext": BuildExtension
    }
)

编译过程：

python setup.py install

核心操作：

[1/2] nvcc -c add2_kernel.cu -o add2_kernel.o
[2/2] c++ -c add2.cpp -o add2.o
x86_64-linux-gnu-g++ -shared add2.o add2_kernel.o -o add2.cpython-37m-x86_64-linux-gnu.so

调用方式：

import torch
import add2
add2.torch_launch_add2(c, a, b, n)

执行命令：

python run_time.py --compiler setup

性能结果：

方法三：CMAKE 编译

特点：使用 CMake 构建系统

cmake_minimum_required(VERSION 3.1 FATAL_ERROR)
set(CMAKE_CUDA_COMPILER "/usr/local/cuda/bin/nvcc")
project(add2 LANGUAGES CXX CUDA)

find_package(Python REQUIRED)
find_package(CUDA REQUIRED)

execute_process(
    COMMAND
        ${Python_EXECUTABLE} -c
            "import torch.utils; print(torch.utils.cmake_prefix_path)"
    OUTPUT_STRIP_TRAILING_WHITESPACE
    OUTPUT_VARIABLE DCMAKE_PREFIX_PATH)

set(CMAKE_PREFIX_PATH "${DCMAKE_PREFIX_PATH}")

find_package(Torch REQUIRED)
find_library(TORCH_PYTHON_LIBRARY torch_python PATHS "${TORCH_INSTALL_PREFIX}/lib")

include_directories(/usr/include/python3.7)
include_directories(../include)

set(SRCS ../kernel/add2_ops.cpp ../kernel/add2_kernel.cu)
add_library(add2 SHARED ${SRCS})

target_link_libraries(add2 "${TORCH_LIBRARIES}" "${TORCH_PYTHON_LIBRARY}")

编译命令：

mkdir build
cd build
cmake ..
make

注意事项：

调用方式：

import torch
torch.ops.load_library("build/libadd2.so")
torch.ops.add2.torch_launch_add2(c, a, b, n)

执行命令：

python run_time.py --compiler cmake

性能结果：

四、核心知识点总结

4.1 架构理解

  1. GPU vs CPU：
  1. 硬件层次：SP → SM → GPU
  2. 软件层次：Thread → Block → Grid
  3. 内存层次：Register/Local Memory → Shared Memory → Global Memory

4.2 编程要点

  1. Kernel 函数：使用 __global__ 修饰，由多个线程并行执行
  2. 索引计算：globalIdx = blockIdx * blockDim + threadIdx
  3. 线程同步：__syncthreads() 仅同步块内线程
  4. Warp 优化：Block 大小设为 32 的倍数
  5. 边界检查：防止线程索引越界

4.3 PyTorch 集成

  1. 三层结构：
  1. 三种编译方式：

4.4 性能观察

从测试结果看，H100 相比 V100 性能提升显著（约 150 倍），体现了新一代 GPU 架构的优势。

参考资料

  1. NVIDIA CUDA C Programming Guide
  2. CUDA 编程上手指南(一)：CUDA C 编程及 GPU 基本知识
  3. godweiyang/NN-CUDA-Example
  4. PyTorch 自定义 CUDA 算子教程（一）
  5. PyTorch 自定义 CUDA 算子教程（二）

渺万里层云，千山暮雪，只影向谁去？ —— 元好问《摸鱼儿·雁丘词》