CUDA1_CUDA编程基础
CUDA(一):CUDA 编程基础
文章来源:知乎专栏 - CUDA(一):CUDA 编程基础
文章ID:645330027
发布时间:编辑于 2025-07-08 14:25・上海
一、GPU 与 CUDA 架构深度解析
1.1 GPU 与 CPU 的设计哲学差异
处理器的两个核心指标是延迟(Latency)和吞吐量(Throughput):
- 延迟:从发出指令到返回结果的时间间隔
- 吞吐量:单位时间内处理的指令条数
CPU:延迟导向设计
CPU 的架构特点:
- 多级高速缓存结构:提升指令访问存储速度
- 复杂控制单元:
- 分支预测机制(Branch Prediction)
- 流水线前传机制(Pipeline Forwarding)
- 强大的运算单元(Core):整型和浮点型复杂运算速度快
设计目标:减少单条指令的执行延迟
GPU:吞吐导向设计
/
性能对比与适用场景
性能差异:
- CPU:连续计算部分,单条复杂指令延迟比 GPU 快 10 倍以上
- GPU:并行计算部分,单位时间内执行指令数量比 CPU 多 10 倍以上
GPU 适用场景:
- 计算密集型任务:数值计算比例远大于内存操作,内存访问延时可被计算掩盖
- 数据并行任务:大任务可拆解为执行相同指令的小任务,对复杂流程控制需求较低
1.2 CUDA 硬件架构层次
CUDA(Compute Unified Device Architecture)是支持 GPU 通用计算的平台和编程模型,提供 C/C++ 语言扩展和用于编程管理 GPU 的 API。
硬件层次结构(从底层到顶层)
1. SP(Stream Processor,线程处理器)
- CUDA 内存模型的最基本单位
- 每个 SP 拥有独立的:
- Registers(寄存器)
- Local Memory(局部内存)
- 特点:寄存器和局部内存只能被自己访问,不同 SP 之间彼此独立
2. SM(Streaming Multiprocessor,多核处理器)
- 由多个 SP + 共享内存构成
- 每个 SM 拥有:
- Shared Memory(共享内存):可被线程块内所有线程访问
- 特点:SM 内的多个 SP 互相并行,互不影响
3. GPU(Device)
- 由多个 SM + 全局内存构成
- 所有 SM 共享:
- Global Memory(全局内存):所有线程块的线程都可访问
内存层次总结
从线程视角看内存访问权限:
- 每个 thread 拥有独立的 register 和 local memory
- 同一 block 中的所有 thread 共享 shared memory
- 所有 thread(包括不同 block)共享 global memory
- 不同 grid 拥有各自的 global memory
1.3 CUDA 软件抽象模型
硬件与软件的对应关系:
- SP(线程处理器) ↔ Thread(线程)
- SM(多核处理器) ↔ Thread Block(线程块)
- Device(设备端) ↔ Grid(线程块组合体)
Thread Block(线程块)特性
线程块是软件侧的基本执行单位,具有以下特点:
- 块内协作:线程通过共享内存、原子操作和屏障同步进行协作
- 块间独立:不同块中的线程不能协作
- 执行独立性:线程块必须能以任意顺序(并行或串行)执行
二、CUDA 编程核心要素
2.1 Kernel 函数与线程组织
Kernel 定义:
- CUDA C++ 定义的基本函数执行单元
- 使用
__global__修饰符声明 - 调用时由 N 个不同的 CUDA 线程并行执行 N 次(而非像常规 C++ 函数只执行一次)
线程索引与线程 ID 的关系
- 一维块:线程索引 = 线程 ID
- 二维块(大小 D_x × D_y):索引 (x, y) 的线程 ID =
x + y * D_x - 三维块(大小 D_x × D_y × D_z):索引 (x, y, z) 的线程 ID =
x + y * D_x + z * D_x * D_y
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);
...
}
关键点:
- 使用
threadIdx.x和threadIdx.y获取线程在块内的索引 - 单个 block 包含 N×N 个线程
- 每个线程处理一个矩阵元素
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);
...
}
关键改进:
- 每个 block 使用 16×16 个线程
- 全局索引计算:
blockIdx * blockDim + threadIdx - 边界检查:
if (i < N && j < N)防止越界访问
内置变量:
blockIdx:block 在 grid 中的索引(可以是一维、二维或三维)blockDim:block 的尺寸(每个维度的线程数)threadIdx:thread 在 block 中的索引
2.4 线程同步机制
**__syncthreads() 函数**:
- 作用:块内线程同步屏障
- 行为:块中所有线程必须等待到达此点,然后才能继续执行
- 用途:协调内存访问,确保共享内存的数据一致性
重要限制:
- 只能同步同一 block 内的线程
- 不同 block 的线程无法同步
2.5 Warp(线程束):GPU 执行的最小单元
SIMT 架构(Single-Instruction, Multiple-Thread)
核心概念:
- GPU 的每一行由 1 个控制单元 + 若干计算单元组成
- 所有计算单元执行相同的控制指令,但处理不同数据
- 这是典型的"单指令多数据流"(SIMD)机制
Warp 定义:
- 一个 warp 包含 32 个并行 thread
- 这些 thread 以不同数据资源执行相同指令
- Warp 是 SM 采用 SIMT 架构的最基本执行单元
- Warp 本质上是线程在 GPU 上运行的最小单元
实践要点
关键优化建议:
- Block 的 thread 数量应设置为 32 的倍数(因为 warp 大小为 32)
- 这样可以避免 warp 资源浪费,提高执行效率
Kernel 执行流程
当一个 kernel 被执行时:
- Grid 中的线程块被分配到 SM 上
- 一个线程块的 thread 只能在一个 SM 上调度
- 一个 SM 通常可以调度多个线程块
- 大量 thread 可能被分配到不同的 SM 上
- 每个 thread 拥有独立的程序计数器和状态寄存器
- 每个 thread 用自己的数据执行指令(SIMT)
硬件限制:
- 每个块的线程数量有限(A100 上最多 1024 个线程)
- 原因:块的所有线程必须驻留在同一个 SM 核心上,共享该核心的有限内存资源
文章概要
本文是关于CUDA(一):CUDA 编程基础的技术文章,属于高性能计算与AI基础设施系列。
核心内容
一、GPU 与 CUDA 结构
1.1 GPU 再认识
对于处理器而言,有2个指标是最主要的:延迟 和吞吐量
。延迟,是指从发出指令到最终返回结果中间经历的时间间隔。而吞吐量,就是单位之间内处理的指令的条数。下面将主要从这两个方面来比较GPU和CPU:
下图左是 CPU 的示意图,有以下几个特点:
- CPU 中包含了多级高速的缓存结构。 这样提升了指令访问存储的速度。
- CPU 中包含了很多控制单元。 具体有2种,一个是分支预测机制,另一个是流水线前传机制。
- CPU 的运算单元 (Core) 强大,整型浮点型复杂运算速度快。
基于以上三点,CPU 在设计时的导向就是减少指令的时延,被称之为延迟导向设计 。
下图右是 GPU 的示意图,有以下几个特点:
- GPU 中虽有缓存结构但是数量少。 因为要减少指令访问缓存的次数。
- GPU 中控制单元非常简单。 控制单元中没有分支预测机制和数据转发机制,对于复杂的指令运算就会比较慢。
- GPU 的运算单元 (Core) 非常多,采用长延时流水线以实现高吞吐量。 每一行的运算单元的控制器只有一个,意味着每一行的运算单元使用的指令是相同的,不同的是它们的数据内容。那么这种整齐划一的运算方式使得 GPU 对于那些控制简单但运算高效的指令的效率显著增加。
基于此,可以看到 GPU 在设计过程中以一个原则为核心:增加简单指令的吞吐,这称 GPU 为吞吐导向设计。
GPU vs CPU
由于设计原则不同,二者擅长的场景有所不同:
- CPU 在连续计算部分,延迟优先,CPU 比 GPU 单条复杂指令延迟快10倍以上。
- GPU 在并行计算部分,吞吐优先,GPU 比 CPU 单位时间内执行指令数量10倍以上。
进一步可以具体化适合 GPU 的场景:
- 计算密集:数值计算的比例要远大于内存操作,因此内存访问的延时可以被计算掩盖。
- 数据并行:大任务可以拆解为执行相同指令的小任务,因此对复杂流程控制的需求较低。
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 算子设计与实现
设计规格
- 功能:两个形状为 n×n 的 tensor 相加
- Block 结构:二维,每个 block 有 16×16 个线程
- Grid 结构:二维,共有 (n/16) × (n/16) 个 block
索引映射关键
每个线程需要从线程索引映射到全局线性内存索引:
- 全局行索引:
i = blockIdx.x * blockDim.x + threadIdx.x - 全局列索引:
j = blockIdx.y * blockDim.y + threadIdx.y - 线性内存索引:
idx = j * n + i
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);
}
关键点:
MatAdd是 kernel 函数,运行在 GPU 端launch_add2是 CPU 端的执行函数,调用 kernel- Kernel 调用是异步的,调用完后控制权立即返回给 CPU
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);
}
工作流程:
torch_launch_add2接收 C++ 版本的 torch tensor- 转换为 C++ 指针数组
- 调用 CUDA 函数
launch_add2执行核函数 - 使用 pybind11 封装,生成 Python 可调用的 .so 库
Torch 使用 CUDA 算子的三步骤
- 编写 CUDA 算子和对应的调用函数
- 编写 torch cpp 函数建立 PyTorch 和 CUDA 之间的联系,用 pybind11 封装
- 用 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
性能结果:
- V100:Cuda time: 2443.504us, Torch time: 2450.132us
- H100:Cuda time: 16.427us, Torch time: 15.330us
方法二: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
性能结果:
- V100:Cuda time: 2445.340us, Torch time: 2449.226us
- H100:Cuda time: 13.733us, Torch time: 14.949us
方法三: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
注意事项:
- cpp 端使用
TORCH_LIBRARY进行封装 - 生成
libadd2.so动态链接库
调用方式:
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
性能结果:
- V100:Cuda time: 2454.185us, Torch time: 2445.102us
- H100:Cuda time: 18.907us, Torch time: 16.665us
四、核心知识点总结
4.1 架构理解
- GPU vs CPU:
- CPU:延迟导向,适合复杂控制流
- GPU:吞吐导向,适合数据并行和计算密集型任务
- 硬件层次:SP → SM → GPU
- 软件层次:Thread → Block → Grid
- 内存层次:Register/Local Memory → Shared Memory → Global Memory
4.2 编程要点
- Kernel 函数:使用
__global__修饰,由多个线程并行执行 - 索引计算:
globalIdx = blockIdx * blockDim + threadIdx - 线程同步:
__syncthreads()仅同步块内线程 - Warp 优化:Block 大小设为 32 的倍数
- 边界检查:防止线程索引越界
4.3 PyTorch 集成
- 三层结构:
- CUDA Kernel(.cu 文件)
- C++ 封装(.cpp 文件,使用 pybind11)
- Python 调用
- 三种编译方式:
- JIT:开发调试方便,首次运行慢
- Setup:标准 Python 包安装方式
- CMake:灵活性高,适合复杂项目
4.4 性能观察
从测试结果看,H100 相比 V100 性能提升显著(约 150 倍),体现了新一代 GPU 架构的优势。
参考资料
- NVIDIA CUDA C Programming Guide
- CUDA 编程上手指南(一):CUDA C 编程及 GPU 基本知识
- godweiyang/NN-CUDA-Example
- PyTorch 自定义 CUDA 算子教程(一)
- PyTorch 自定义 CUDA 算子教程(二)
渺万里层云,千山暮雪,只影向谁去? —— 元好问《摸鱼儿·雁丘词》