CUDA 编程模型——把矩阵乘法映射到 GPU / The CUDA Programming Model for Mapping Matrix Multiplication to GPUs
📅 创建时间：2026-06-02 🏷️ 标签：#CUDA #kernel #grid #block #thread #nvcc 📚 前置知识：[[03-gpu-architecture]]（GPU 架构） 📚 相关知识：[[05-cuda-kernel-and-memory]]（CUDA 内存） [[06-cuda-optimization]]（性能优化）
先抓住直觉
CUDA 的核心不是新语法，而是把一份工作拆给许多编号不同的线程。Host 负责准备和发号施令，Device 负责批量计算；每个线程根据自己的编号找到要处理的数据。
- 必须理解：Host/Device 分工；Grid、Block、Thread 的包含关系；线程索引如何映射到数据。
- 用到再查：函数修饰符全集、CMake 配置和错误码。
- 最小目标：先看懂向量加法，再进入矩阵乘法；不要一开始就手写高性能 GEMM。
场景：第一个 CUDA 程序——向量加法
┌─────────────────────────────────────────────────────────────┐
│ │
│ 你已经理解了 GPU 的架构： │
│ - 6912 个 CUDA 核心 │
│ - SM → Warp → Thread 的调度层次 │
│ - SIMT 执行模型 │
│ │
│ 现在，你要把这个简单的 C++ 代码改成 CUDA 版本： │
│ │
│ // C++ 版本 │
│ for (int i = 0; i < N; i++) { │
│ c[i] = a[i] + b[i]; │
│ } │
│ │
│ 目标：在 GPU 上并行执行这段计算 │
│ │
└─────────────────────────────────────────────────────────────┘第1节：Host 和 Device——CUDA 的核心概念
两种代码，两种运行环境
┌─────────────────────────────────────────────────────────────┐
│ CUDA 程序的两部分 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Host（C++ 代码，运行在 CPU 上）： │
│ - 数据准备：把数据从 CPU 内存拷贝到 GPU 显存 │
│ - 调用 kernel：告诉 GPU 要执行什么计算 │
│ - 结果回收：把结果从 GPU 显存拷回 CPU 内存 │
│ - 控制流：for/if 等逻辑 │
│ │
│ Device（CUDA 代码，运行在 GPU 上）： │
│ - 具体计算：每个线程执行一小部分计算 │
│ - 线程协作：同 SM 内线程共享数据 │
│ │
│ 形象理解： │
│ Host = 包工头（CPU）：分配任务、准备材料、收集结果 │
│ Device = 建筑队（GPU）：大规模并行施工 │
│ │
└─────────────────────────────────────────────────────────────┘代码分布示例
cpp
// ================================================================
// main.cpp（C++，Host 端）
// ================================================================
#include <cuda_runtime.h> // CUDA 运行时 API
#include <stdio.h>
int main() {
const int N = 1000000;
float *h_a, *h_b, *h_c; // h_ = host 内存（CPU）
float *d_a, *d_b, *d_c; // d_ = device 内存（GPU）
// Step 1: 分配 Host 内存
h_a = (float*)malloc(N * sizeof(float));
h_b = (float*)malloc(N * sizeof(float));
h_c = (float*)malloc(N * sizeof(float));
// Step 2: 初始化数据
for (int i = 0; i < N; i++) {
h_a[i] = i * 1.0f;
h_b[i] = i * 2.0f;
}
// Step 3: 分配 Device 内存（GPU 显存）
cudaMalloc(&d_a, N * sizeof(float));
cudaMalloc(&d_b, N * sizeof(float));
cudaMalloc(&d_c, N * sizeof(float));
// Step 4: CPU 数据拷贝到 GPU
cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);
// Step 5: 调用 kernel（启动 GPU 计算）
// 解释见下一节
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
// Step 6: GPU 结果拷贝回 CPU
cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);
// Step 7: 验证结果
for (int i = 0; i < 10; i++) {
printf("c[%d] = %f\n", i, h_c[i]);
}
// Step 8: 释放内存
free(h_a); free(h_b); free(h_c);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}cpp
// ================================================================
// vector_add.cuh（CUDA 头文件）
// ================================================================
#ifndef VECTOR_ADD_CUH
#define VECTOR_ADD_CUH
// __global__ 声明的函数从 CPU 调用，在 GPU 上执行
void add(float* a, float* b, float* c, int N);
#endifcpp
// ================================================================
// vector_add.cu（CUDA 源代码，Device 端）
// ================================================================
// CUDA 头文件
#include "vector_add.cuh"
#include <cuda_runtime.h>
// __global__: 从 CPU 调用，在 GPU 上运行的函数（称为 "kernel"）
// 这是 CUDA 编程的核心：一个 kernel 被 N 个线程并行执行
__global__
void add_kernel(float* a, float* b, float* c, int N) {
// 每个线程计算一个元素的结果
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 边界检查：防止数组越界
if (i < N) {
c[i] = a[i] + b[i];
}
}
// 包装函数（Host 端调用这个，它再调用 kernel）
void add(float* a, float* b, float* c, int N) {
// 计算 grid 和 block 的维度
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 启动 kernel：<<<blocksPerGrid, threadsPerBlock>>>
add_kernel<<<blocksPerGrid, threadsPerBlock>>>(a, b, c, N);
// 注意：kernel 启动是异步的！
// CPU 不会等待 GPU 完成就继续执行下一行
// 下一行立即执行，除非调用 cudaMemcpy 或 cudaDeviceSynchronize
}编译和运行
bash
# 编译
nvcc vector_add.cu main.cpp -o vector_add
# 运行（需要机器上有 NVIDIA GPU）
./vector_add
# 如果没有 GPU，可以模拟运行（需要 NVIDIA 开发驱动）
# CUDA Toolkit 自带 cuda-gdb 和 compute-sanitizer第2节：Grid / Block / Thread——CUDA 的三层线程组织
线程层次结构
┌─────────────────────────────────────────────────────────────┐
│ CUDA 线程层次结构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Grid（整个 GPU） │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ Block 0 Block 1 Block 2 Block 3 ... │ │
│ │ ┌─────┐ ┌─────┐ ┌─────┐ ┌─────┐ │ │
│ │ │ T0 │ │ T0 │ │ T0 │ │ T0 │ │ │
│ │ │ T1 │ │ T1 │ │ T1 │ │ T1 │ │ │
│ │ │ ... │ │ ... │ │ ... │ │ ... │ │ │
│ │ │ T255 │ │ T255 │ │ T255 │ │ T255 │ │ │
│ │ └─────┘ └─────┘ └─────┘ └─────┘ │ │
│ │ 256线程 256线程 256线程 256线程 │ │
│ │ │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ Grid 的维度：gridDim.x, gridDim.y, gridDim.z │
│ Block 的维度：blockDim.x, blockDim.y, blockDim.z │
│ │
└─────────────────────────────────────────────────────────────┘线程索引计算
cpp
// kernel 内部，每个线程通过内置变量知道自己负责哪个数据
// 内置变量速查：
// threadIdx.x/y/z 当前线程在线程块内的索引
// blockIdx.x/y/z 当前线程块在 grid 内的索引
// blockDim.x/y/z 每个线程块的维度
// gridDim.x/y/z grid 的维度
// 计算全局线程索引（唯一标识每个线程）
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 二维 Grid + 二维 Block 的例子：
// gridDim = (Gx, Gy), blockDim = (Bx, By)
// 全局索引 = (by * Bx + bx) * (Gy * Gx) + (ty * Bx + tx)
// 但实际代码中通常只用一维就够用了维度设置规则
┌─────────────────────────────────────────────────────────────┐
│ 维度设置规则 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Block 维度（每个 block 的线程数）： │
│ - 最大值：1024（硬件限制） │
│ - 常用值：128, 256, 512, 1024 │
│ - 推荐值：256（经验最优） │
│ │
│ Block 内线程数最好是 32 的倍数（Warp 大小） │
│ 因为 warp = 32 threads，必须整 warp 调度 │
│ 如果不是 32 的倍数，会浪费部分硬件 │
│ │
│ Grid 维度（block 的数量）： │
│ - 最大值：2^31 - 1（x方向）、65535（y）、65535（z） │
│ - Grid 大小 = block 数 × 每 block 线程数 = 总线程数 │
│ - 总线程数应该 >= GPU 核心数（通常远大于） │
│ │
│ 示例： │
│ A100 有 6912 个核心，N = 1000000 │
│ threadsPerBlock = 256 │
│ blocksPerGrid = ceil(1000000 / 256) = 3907 │
│ 总线程数 = 3907 × 256 = 1,000,192 > 6912 ✅ │
│ │
└─────────────────────────────────────────────────────────────┘第3节：CUDA 关键字——三种函数的区别
__global__ / __device__ / __host__
┌─────────────────────────────────────────────────────────────┐
│ CUDA 函数修饰符 │
├─────────────────────────────────────────────────────────────┤
│ │
│ __global__ void my_kernel(...) │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：CPU（Host） │ │
│ │ 执行位置：GPU（Device） │ │
│ │ 用途：kernel 入口（必须用这个声明） │ │
│ │ 注意：返回类型必须是 void │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ __device__ float my_device_func(float x) │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：GPU（Device） │ │
│ │ 执行位置：GPU（Device） │ │
│ │ 用途：GPU 内部调用的辅助函数 │ │
│ │ 注意：只能被 __global__ 或 __device__ 函数调用 │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ __host__ void my_host_func(...) ← 默认，省略不写 │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：CPU（Host） │ │
│ │ 执行位置：CPU（Host） │ │
│ │ 用途：普通 C++ 函数 │ │
│ │ 注意：可以同时有 __host__ 和 __device__（两边都编译）│ │
│ └───────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘变量修饰符
┌─────────────────────────────────────────────────────────────┐
│ CUDA 变量修饰符 │
├─────────────────────────────────────────────────────────────┤
│ │
│ __device__ ← 默认，全局内存，GPU 上 │
│ - 全局内存变量，所有线程可见 │
│ - 通过 cudaMalloc/cudaFree 管理 │
│ │
│ __shared__ ← 共享内存，同 SM 内线程 │
│ - 声明在 kernel 内部（或者用 shared 限定符） │
│ - 同一 block 内所有线程共享 │
│ - 速度极快（~1ns，接近寄存器） │
│ │
│ __constant__ ← 常量内存，只读 │
│ - 所有线程只读 │
│ - 编译时常量或运行时通过 cudaMemcpyToSymbol 设置 │
│ - 有硬件缓存（广播优化） │
│ │
│ __managed__ ← 统一内存（CUDA 6.0） │
│ - CPU 和 GPU 共享同一地址空间 │
│ - 自动在需要时迁移数据 │
│ │
└─────────────────────────────────────────────────────────────┘第4节：矩阵乘法的 CUDA 实现——从向量到矩阵
最简单的矩阵乘法
目标：C = A × B
A: M×K 矩阵
B: K×N 矩阵
C: M×N 矩阵cpp
// matrix_mul.cu
__global__
void matrix_mul_kernel(float* A, float* B, float* C,
int M, int N, int K) {
// 每个线程计算 C 的一个元素 C[row][col]
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < M && col < N) {
float sum = 0.0f;
for (int k = 0; k < K; k++) {
sum += A[row * K + k] * B[k * N + col];
}
C[row * N + col] = sum;
}
}cpp
// 调用方式（二维 Grid + 二维 Block）
void matrix_mul(float* A, float* B, float* C,
int M, int N, int K) {
// Block 维度：16×16 = 256 线程/block
dim3 threadsPerBlock(16, 16);
// Grid 维度：ceil(M/16) × ceil(N/16)
dim3 blocksPerGrid(
(N + threadsPerBlock.x - 1) / threadsPerBlock.x,
(M + threadsPerBlock.y - 1) / threadsPerBlock.y
);
matrix_mul_kernel<<<blocksPerGrid, threadsPerBlock>>>(
A, B, C, M, N, K);
}为什么一个线程算一个输出元素？
┌─────────────────────────────────────────────────────────────┐
│ 线程与输出的映射 │
├─────────────────────────────────────────────────────────────┤
│ │
│ C 是 M×N 的矩阵，有 M×N 个输出元素 │
│ 每个输出元素 C[i][j] 的计算是独立的！ │
│ │
│ ┌─────────────────────────────────────────────────────┐ │
│ │ Thread (block=0, thread=0) → 计算 C[0][0] │ │
│ │ Thread (block=0, thread=1) → 计算 C[0][1] │ │
│ │ Thread (block=0, thread=2) → 计算 C[0][2] │ │
│ │ ... │ │
│ │ Thread (block=1, thread=0) → 计算 C[1][0] │ │
│ │ ... │ │
│ │ Thread (block=M-1, thread=N-1) → 计算 C[M-1][N-1]│ │
│ └─────────────────────────────────────────────────────┘ │
│ │
│ M×N 个线程同时执行 = M×N 个输出同时计算 = 完美并行！ │
│ │
│ 这就是 GPU 编程的核心思想： │
│ 把计算拆分成 N 个独立的小任务，每个线程做一个 │
│ │
└─────────────────────────────────────────────────────────────┘更高效的共享内存版本（预告）
cpp
// 使用共享内存优化（大幅减少全局内存访问）
// 具体原理在第 5 章讲解，这里先看结构
__global__
void matrix_mul_shared_kernel(float* A, float* B, float* C,
int M, int N, int K) {
// 共享内存：同 block 内线程共享
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
float sum = 0.0f;
// 分块计算：每次加载一块到共享内存
for (int m = 0; m < (K + BLOCK_SIZE - 1) / BLOCK_SIZE; m++) {
// 线程协作加载到共享内存
if (row < M && (m * BLOCK_SIZE + threadIdx.x) < K)
As[threadIdx.y][threadIdx.x] = A[row * K + m * BLOCK_SIZE + threadIdx.x];
if (col < N && (m * BLOCK_SIZE + threadIdx.y) < K)
Bs[threadIdx.y][threadIdx.x] = B[(m * BLOCK_SIZE + threadIdx.y) * N + col];
// 等待所有线程加载完成
__syncthreads();
// 计算这一块
for (int k = 0; k < BLOCK_SIZE; k++)
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
// 等待所有线程计算完成，再加载下一块
__syncthreads();
}
if (row < M && col < N)
C[row * N + col] = sum;
}第5节：CUDA 错误处理——容易被忽略的坑
异步执行和错误检查
cpp
// ❌ 错误：cudaMalloc 的返回值不是错误码
float* d_a;
cudaError_t err = cudaMalloc(&d_a, N * sizeof(float));
if (err != cudaSuccess) {
printf("CUDA malloc failed: %s\n", cudaGetErrorString(err));
}
// ✅ 正确：kernel 启动是异步的！
// cudaGetLastError() 可以检查 kernel 启动是否有错
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
err = cudaGetLastError(); // kernel 启动后立即调用
if (err != cudaSuccess) {
printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}
// ✅ 正确：等待 GPU 完成后再检查结果
cudaDeviceSynchronize(); // 等待 GPU 所有任务完成
err = cudaGetLastError();
if (err != cudaSuccess) {
printf("Kernel execution failed: %s\n", cudaGetErrorString(err));
}完整错误检查封装
cpp
// 封装成宏，方便使用
#define CHECK_CUDA(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
printf("CUDA error at %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while (0)
// 使用
CHECK_CUDA(cudaMalloc(&d_a, N * sizeof(float)));
CHECK_CUDA(cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice));
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
CHECK_CUDA(cudaGetLastError()); // kernel 启动后立即检查
CHECK_CUDA(cudaDeviceSynchronize()); // 等待完成
CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));第6节：完整的项目结构示例
┌─────────────────────────────────────────────────────────────┐
│ 完整项目文件结构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ matrix_mul_project/ │
│ ├── CMakeLists.txt │
│ ├── include/ │
│ │ ├── matrix_mul.cuh ← kernel 声明 │
│ │ └── common.h ← 通用工具函数 │
│ ├── src/ │
│ │ ├── main.cpp ← Host 代码（C++） │
│ │ └── matrix_mul.cu ← Device 代码（CUDA） │
│ └── build/ │
│ └── (编译输出) │
│ │
└─────────────────────────────────────────────────────────────┘cmake
# CMakeLists.txt
cmake_minimum_required(VERSION 3.18)
project(matrix_mul LANGUAGES CXX CUDA)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_ARCHITECTURES 80) # A100
# CUDA Toolkit
find_package(CUDAToolkit REQUIRED)
# 头文件目录
include_directories(${CMAKE_SOURCE_DIR}/include)
# 可执行文件
add_executable(matrix_mul
src/main.cpp
src/matrix_mul.cu
)
target_link_libraries(matrix_mul PRIVATE CUDA::cudart)
# 启用 CUDA 扩展（CMake 自动处理 .cu 文件）
enable_language(CUDA)bash
# 编译
cd matrix_mul_project
mkdir build && cd build
cmake .. -DCMAKE_BUILD_TYPE=Release
cmake --build . --config Release
# 运行
./matrix_mul升华：CUDA 编程的核心心法
┌─────────────────────────────────────────────────────────────┐
│ CUDA 编程检查清单 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 1. 线程组织 │
│ ✅ 每个线程处理一个或多个数据元素 │
│ ✅ Block 维度设为 32 的倍数（warp 对齐） │
│ ✅ Grid 大小 >= GPU SM 数量（充分利用硬件） │
│ │
│ 2. 内存管理 │
│ ✅ cudaMalloc 分配 GPU 显存 │
│ ✅ cudaMemcpy 传输数据（注意方向 HostToDevice） │
│ ✅ cudaFree 释放显存 │
│ │
│ 3. 错误检查 │
│ ✅ 所有 CUDA 调用都要检查返回值 │
│ ✅ kernel 启动后立即检查 cudaGetLastError │
│ ✅ 验证结果前调用 cudaDeviceSynchronize │
│ │
│ 4. 性能意识 │
│ ✅ 线程数太少 → GPU 利用率低 │
│ ✅ 线程数太多 → 寄存器溢出，效率下降 │
│ ✅ 256 线程/block 是经验最优起点 │
│ │
└─────────────────────────────────────────────────────────────┘"AI 可查 vs 必须理解"清单
AI 可查：
✅ cudaMalloc/cudaMemcpy 的具体参数
✅ dim3 的具体用法
✅ CMake 中 CUDA 项目的具体配置
必须理解：
🔴 Host（CPU）负责调度，Device（GPU）负责计算
🔴 Grid / Block / Thread 的三层组织关系
🔴 线程索引计算：int i = blockIdx.x * blockDim.x + threadIdx.x
🔴 __global__ / __device__ / __host__ 的区别
🔴 kernel 启动是异步的，需要 cudaDeviceSynchronize 等待
🔴 为什么一个线程算一个输出元素 = 完美并行学习状态：🟡 开始学习