Maths-CS-AI Compendium 精读从 GPU 架构到 CUDA 内核编程实战【免费下载链接】maths-cs-ai-compendiumBecome a cracked AI/ML researcher/engineer with this unconventional textbook covering maths, computing, and ML with intuition.项目地址: https://gitcode.com/GitHub_Trending/mat/maths-cs-ai-compendium本文是 maths-cs-ai-compendium 开源教材 第 16 章 SIMD 与 GPU 编程 中 GPU Architecture and CUDA 一讲的深度展开。文章以该文档为主体结合仓库中硬件基础、Triton/TPU、量化与精度格式 等章节作交叉印证。读完你将掌握GPU 与 CPU 的根本设计差异、GPU 内存层次与瓶颈判断、CUDA 的网格-块-线程编程模型、warp/SIMT 执行语义以及内存合并、共享内存分块、流并发、kernel 融合等一套可落地的内核优化方法论并附三个可直接用nvcc编译验证的编码任务。为什么 GPU 主导了机器学习现代 NVIDIA GPU 拥有超过 10000 个 CUDA 核心而一颗 CPU 通常只有 4-128 个核心。100-1000 倍的核心数量优势正是 GPU 统治 ML 的直接原因训练一个 Transformer 需要数以万亿计的乘加操作GPU 能以 CPU 无法企及的规模并行处理这些运算。更关键的是即使你从不亲手编写 CUDA 内核理解 GPU 架构也能解释大量实践中的现象为什么 batch size 很重要——需要足够多的并行工作才能喂饱GPU饱和其吞吐能力为什么内存通常才是瓶颈而不是算力——大多数 ML 算子是内存受限的为什么某些操作scatter、条件分支在 GPU 上很慢——它们破坏了 GPU 的并行执行模型。这三点贯穿本文始终。值得注意的是本文档在仓库中属于第 16 章SIMD 与 GPU 编程的第 4 讲前接硬件基础roofline 模型、延迟 vs 吞吐后接 Triton、TPU 与 Pallas本文是整个从硬件到内核路线图的枢纽它讲述的是 CUDA C 这一层——比 C 内建函数第 2、3 讲抽象更高但比 Triton第 5 讲更贴近硬件。GPU vs CPU两种截然不同的设计哲学CPU 为延迟而设计最小化完成单任务的时间。它把大部分晶体管预算投入到缓存、分支预测器、乱序执行上——一切技巧都为了让一条线程跑得更快细节见硬件基础中对乱序执行、分支预测、投机执行的讲解。GPU 为吞吐而设计最大化每秒完成的任务数。它把大部分晶体管投入到执行单元ALU上。单个线程很慢但线程数以千计。| | CPU | GPU | |--|-----|-----| | 核心 | 4-128 个复杂、快速 | 1,000-20,000 个简单、较慢 | | 时钟频率 | 3-5 GHz | 1-2.5 GHz | | 缓存 | 大L3 可达 32 MB | 小以每 SM 共享内存为主 | | 分支预测 | 精密 | 无所有线程走同一路径 | | 擅长 | 低延迟、复杂控制流 | 高吞吐、数据并行工作负载 | | 典型算力FP32 | 1-5 TFLOPS | 30-80 TFLOPS | | 内存带宽 | 50-100 GB/s | 1-3 TB/s |用硬件基础中的比喻来说GPU 是公交车——单次操作延迟高但一次能载 50 人CPU 是出租车——延迟低但载客量有限。这也是为什么 GPU 适合训练需要吞吐每秒处理海量样本而 CPU 适合操作系统类任务需要低延迟立刻响应按键。GPU 的内存带宽优势10-30 倍往往比算力优势更重要。许多 ML 操作是内存受限的逐元素运算、归一化、注意力GPU 的高带宽能让数据以足够快的速度喂给核心。这正是 roofline 模型的核心结论可达到的 FLOPS min(峰值算力, 带宽 × 算力强度)。矩阵乘法算力强度高O(n³) 运算对 O(n²) 数据是计算受限的而 ReLU、逐元素加法算力强度低是内存受限的。GPU 内存层次理解瓶颈的第一步理解 GPU 内存至关重要因为内存访问是主要瓶颈而非计算。内存大小延迟带宽作用域寄存器每 SM 约 256 KB0 周期最高每线程私有共享内存每 SM 48-228 KB约 5 周期约 20 TB/s每线程块L1 缓存每 SM 128-256 KB约 30 周期每 SML2 缓存4-96 MB约 200 周期约 6 TB/s全局全局内存HBM24-192 GB约 400 周期1-3.3 TB/s全局三条关键结论寄存器最快也最有限。每个线程拥有私有寄存器组通常上限 255 个。单线程使用寄存器过多会降低占用率occupancy——能同时运行的线程变少进而削弱 GPU 用海量线程隐藏内存延迟的能力。共享内存是程序员管理的缓存由块内所有线程共享。它是写出快速 CUDA 内核的关键先从慢速全局内存把一块数据tile加载到快速共享内存再基于它计算。这就是主导 GPU 编程的tiling分块模式。全局内存HBM是 GPU 的主内存显存大但慢约 400 周期延迟。所有数据从这里开始、也最终回到这里。内核优化的目标就是最小化全局内存访问次数。CUDA 编程模型CUDACompute Unified Device Architecture统一计算设备架构是 NVIDIA 的 GPU 编程模型。你编写的是内核kernel在 GPU 上运行的函数由数千线程同时执行。层级网格、块、线程Grid一次启动 ├── Block (0,0) │ ├── Thread (0,0) │ ├── Thread (1,0) │ ├── Thread (2,0) │ └── ... 每块最多 1024 线程 ├── Block (1,0) │ ├── Thread (0,0) │ └── ... └── ... 可以有数百万个块Thread线程最小执行单元块内有唯一 IDthreadIdx.x。Block线程块一组可共享内存、可互相同步的线程。块 ID 为blockIdx.x块大小为blockDim.x最多 1024 线程。Grid网格单次内核启动的全部块。可以是 1D、2D 或 3D。每个线程用此式计算自己的全局索引int idx blockIdx.x * blockDim.x threadIdx.x;第一个 CUDA 内核// vector_add.cu — CUDA 源文件.cu 扩展名 #include stdio.h // __global__ 标记这是一个 GPU 内核从 CPU 调用在 GPU 上运行 __global__ void vector_add(const float* a, const 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() { int n 1 20; // 约 100 万元素 size_t bytes n * sizeof(float); // 分配主机CPU内存 float *h_a new float[n]; float *h_b new float[n]; float *h_c new float[n]; // 初始化 for (int i 0; i n; i) { h_a[i] 1.0f; h_b[i] 2.0f; } // 分配设备GPU内存 float *d_a, *d_b, *d_c; cudaMalloc(d_a, bytes); cudaMalloc(d_b, bytes); cudaMalloc(d_c, bytes); // 将数据从 CPU 拷贝到 GPU cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice); // 启动内核每块 256 线程块数足以覆盖 n 个元素 int block_size 256; int grid_size (n block_size - 1) / block_size; // 向上取整除法 vector_addgrid_size, block_size(d_a, d_b, d_c, n); // 将结果从 GPU 拷贝回 CPU cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost); // 验证 printf(c[0] %f (expected 3.0)\n, h_c[0]); // 释放内存 cudaFree(d_a); cudaFree(d_b); cudaFree(d_c); delete[] h_a; delete[] h_b; delete[] h_c; return 0; }说明原文此处存在一处笔误——结果回拷应使用d_c而非d_a否则拷贝的是输入数组 a上文已修正。这也是一个很好的自查点验证输出c[0]是否为 3.0 能立刻暴露这类错误。编译与运行需要 NVIDIA 编译器 nvcc且主机装有支持 CUDA 的 NVIDIA GPU 驱动# 用 NVIDIA 编译器编译 nvcc -O3 -o vector_add vector_add.cu ./vector_addCUDA 关键 C 概念__global__CUDA 关键字标记内核函数。从 CPUhost调用在 GPUdevice上运行。grid_size, block_size内核启动语法指定使用多少块、每块多少线程。cudaMalloc/cudaFree分配/释放 GPU 内存相当于 GPU 版的new/delete。cudaMemcpy在 CPU 与 GPU 之间拷贝数据。这往往是最主要的瓶颈——PCIe 带宽约 32 GB/s而 GPU 内存带宽约 3 TB/s相差近两个数量级。从第 0 讲ML 框架是如何工作的我们知道PyTorch 的torch.matmul调用链是Python → ATenC→ cuBLASNVIDIA 的 GPU 优化 BLAS。cuBLAS 中的每个算子本质上就是这样手写优化的 CUDA 内核本文要写的 kernel正是 PyTorch 底层真实发生的事情。Warp 与 SIMTGPU 如何执行指令GPU 以 32 个线程为一组执行这一组被称为warp经线。一个 warp 内的全部 32 个线程同时执行同一条指令Single Instruction, Multiple Threads——SIMT。这是 SIMD 在 GPU 上的对应物但发生在线程层面SIMD 与 AVX-512 掩码寄存器的对比见 x86 与 AVX 讲。Warp 发散warp divergence当同一 warp 内线程走到if的不同分支时GPU 无法在一个 warp 内同时执行两条不同指令于是它串行执行两个分支屏蔽掉不该参与的分支上的线程。这会让性能减半甚至更差。// 糟糕warp 发散同一 warp 内线程走不同路径 if (threadIdx.x % 2 0) { c[idx] a[idx] b[idx]; // 偶数线程执行这个 } else { c[idx] a[idx] - b[idx]; // 奇数线程执行这个同一 warp串行化 } // 更好无分支无发散 float sign (threadIdx.x % 2 0) ? 1.0f : -1.0f; c[idx] a[idx] sign * b[idx]; // 所有线程执行同一条指令注意由于 warp 大小为 32而threadIdx.x % 2的奇偶模式在 0-31 号线程中恰好各占一半因此该例会把 warp 切成两个串行路径。通常发散惩罚的程度取决于分支的线程分布。内存合并第一个性能关口合并访问coalesced access当连续的线程访问连续的内存地址时GPU 把它们合并为一次内存事务。这是性能的关键。// 好合并——线程 0 读 a[0]线程 1 读 a[1]…… c[idx] a[idx] b[idx]; // 差跨步——线程 0 读 a[0]线程 1 读 a[stride]…… c[idx] a[idx * stride] b[idx * stride]; // stride 1 浪费带宽量化地看一个 warp 的 32 个线程合并访问只需加载 128 字节32 × 4 字节的 float32即一次事务跨步访问则需要多次事务每次仍加载 128 字节却只用其中一小部分。stride 为 32 是最坏情况每次事务加载 128 字节但只有一个线程用到 4 字节约 3% 利用率。共享内存与分块Tiling最重要的优化模式分块模式是 GPU 编程中最重要的优化技术。思路把数据块从慢速全局内存加载到快速共享内存在共享内存上计算再把结果写回全局内存。// 使用共享内存分块的矩阵乘法简化版 __global__ void matmul_tiled(const float* A, const float* B, float* C, int M, int N, int K) { // A 和 B 各一块 tile 的共享内存 __shared__ float tile_A[TILE_SIZE][TILE_SIZE]; __shared__ float tile_B[TILE_SIZE][TILE_SIZE]; int row blockIdx.y * TILE_SIZE threadIdx.y; int col blockIdx.x * TILE_SIZE threadIdx.x; float sum 0.0f; // 按 tile 循环 for (int t 0; t (K TILE_SIZE - 1) / TILE_SIZE; t) { // 将 A、B 各一个 tile 加载进共享内存 if (row M t * TILE_SIZE threadIdx.x K) tile_A[threadIdx.y][threadIdx.x] A[row * K t * TILE_SIZE threadIdx.x]; else tile_A[threadIdx.y][threadIdx.x] 0.0f; if (col N t * TILE_SIZE threadIdx.y K) tile_B[threadIdx.y][threadIdx.x] B[(t * TILE_SIZE threadIdx.y) * N col]; else tile_B[threadIdx.y][threadIdx.x] 0.0f; __syncthreads(); // 等待所有线程加载完毕 // 基于本 tile 计算部分点积 for (int k 0; k TILE_SIZE; k) { sum tile_A[threadIdx.y][k] * tile_B[k][threadIdx.x]; } __syncthreads(); // 加载下一 tile 之前再同步 } if (row M col N) C[row * N col] sum; }两个核心原语__shared__声明块内所有线程共享的内存快速、片上。__syncthreads()屏障等待块内所有线程到达此处。在写入共享内存与读取它之间必须调用否则部分线程会读到过期数据。为什么分块有效不分块时每个线程的每次乘法都要从全局内存取数分块后一个 TILE_SIZE × TILE_SIZE 的数据块只从全局内存加载一次被块内所有线程复用。复用因子就是 TILE_SIZE全局内存流量因此降低 TILE_SIZE 倍。矩阵乘法的内层循环向量点积还天然映射到 SIMD 乘加指令——这正是 BLAS 库cuBLAS、MKL、OpenBLAS为何被重度 SIMD/GPU 优化的根本原因。流Streams与并发让 GPU 始终忙碌默认情况下CUDA 操作是串行的CPU 启动一个内核等它结束后再启动下一个。**流stream**让不同流的操作可以重叠cudaStream_t stream1, stream2; cudaStreamCreate(stream1); cudaStreamCreate(stream2); // 这些操作可以重叠不同流中的操作并发执行 cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1); cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2); kernel1grid, block, 0, stream1(d_a, d_c); kernel2grid, block, 0, stream2(d_b, d_d);流能把数据传输与计算重叠一个流的内核在跑时另一个流正在拷贝数据。这掩盖了 PCIe 传输延迟让 GPU 保持忙碌。这正是第 0 讲所述 ML 框架调度层torch.compile、XLA 等在底层反复优化的模式之一。剖析 CUDA 代码# NVIDIA Nsight Compute内核级剖析 ncu --set full ./my_program # NVIDIA Nsight Systems系统级时间线 nsys profile ./my_program # 快速指标 ncu --metrics sm__throughput,dram__throughput ./my_program看什么占用率occupancySM 容量被使用的比例。低于 50% 意味着线程太少不足以隐藏内存延迟。常见诱因单线程寄存器过多、每块共享内存过多。内存吞吐与峰值带宽对比。若不足峰值的 50%说明访问模式低效未合并访问、bank 冲突。计算吞吐与峰值 FLOPS 对比。若内存与计算吞吐都很低内核是延迟受限的并行度不足。结合 roofline 模型先用算力强度判定内核是内存受限还是计算受限再对症下药——不要盲目优化。高级优化技术在合并访问与共享内存分块之外高性能 GPU以及 CPU代码还常用以下技巧。数据布局AoS vs SoAAoSArray of Structures结构体数组每个元素的所有字段紧挨着存放。[{x,y,z}, {x,y,z}, {x,y,z}]。SoAStructure of Arrays数组结构体每个字段存放在各自连续的数组中。{[x,x,x], [y,y,y], [z,z,z]}。// AoS对 SIMD/GPU 不利访问所有 x 值会触碰不连续内存 struct Particle { float x, y, z, mass; }; Particle particles[N]; // particles[0].x 与 particles[1].x 相距 16 字节 // SoA对 SIMD/GPU 有利所有 x 值连续 struct Particles { float x[N], y[N], z[N], mass[N]; }; // x[0] 与 x[1] 相距 4 字节——完美适合合并访问与 SIMD对数据并行工作负载SIMD、GPUSoA 几乎总是更快AoS 仅在总是同时访问某元素全部字段时更优数值代码中很少见。PyTorch 张量天生就是 SoA每个特征是一个连续维度。软件预取Software Prefetching可以主动告诉 CPU 提前加载数据从而隐藏内存延迟#include xmmintrin.h // 提供 _mm_prefetch for (int i 0; i n; i 4) { _mm_prefetch((char*)(a i 64), _MM_HINT_T0); // 预取 64 个元素之后的数据 // 用 SIMD 处理 a[i:i4] __m128 va _mm_load_ps(a i); // ... }预取指令只是提示数据若已在缓存中则是 no-op若不在CPU 会在执行其他指令的同时于后台拉取。预取距离此例为 64 个元素之后需要根据内存延迟与循环迭代时间调参。这与第 0 讲中何时写自定义 C 内核的第 2 条融合操作提升性能是同一套性能思维的延伸。内核融合Kernel Fusion内核融合把多个操作合并进单个内核避免把中间结果写回内存。这是对 ML 影响最大的单个 GPU 优化// 未融合3 次内核启动、3 次全局内存往返 y matmul(x, W) // 将 y 写入全局内存 z y bias // 读 y写 z out relu(z) // 读 z写 out // 融合1 次内核启动、1 次全局内存写入 out fused_matmul_bias_relu(x, W, bias) // y 和 z 从不离开 SRAM对内存受限操作bias 加、ReLU、LayerNorm内存流量主导执行时间融合则彻底消除这部分流量。PyTorch 的torch.compile与 Triton 可以自动或几乎零成本地完成融合——Triton 的 softmax 示例第 5 讲展示的正是这一点PyTorch 内置F.softmax启动 3 个内核max、exp-and-sum、divide而 Triton 版一个内核在寄存器/SRAM 内完成全部工作。这也是 roofline 模型中融合把三个内存受限操作变成一个计算受限操作的具体落地。混合精度内核Mixed-Precision Kernels用低精度FP16、BF16、INT8计算、高精度FP32累加可以兼得速度与精度// Tensor CoreFP16 矩阵相乘、FP32 累加 // 每条 Tensor Core 指令D (FP32) A (FP16) × B (FP16) C (FP32) nvcuda::wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);FP16 只有 FP32 一半大小因此内存带宽翻倍带宽通常是瓶颈缓存能容纳 2 倍数据。Tensor Core 处理 FP16 的速度是 FP32 CUDA 核心的 8-16 倍。这就是混合精度训练对应第 6 章深度学习内容能带来 2-3 倍加速且精度损失极小的原因。仓库 第 17 章量化文档给出了完整的精度格式谱系可与之对照TF32A100 训练、FP16混合精度训练、BF16与 FP32 同指数范围训练更安全、FP8 E4M3前向、FP8 E5M2梯度。H100 在 FP8 下 989 TFLOPS而 FP32 只有 67 TFLOPS相差约 15 倍——这就是低精度内核收益的数量级来源。内存池分配器Memory Pool AllocatorscudaMalloc很慢每次调用约 1 ms因为它要与 GPU 同步。训练循环每轮迭代都要分配临时缓冲区累积开销可观。内存池PyTorch 的缓存分配器、CUDA memory pools预先分配一大块 GPU 内存从中二次分配避免系统调用# PyTorch 自动做这件事——但理解其原理很重要 # 每次 torch.empty() 都从池中复用内存不触发 cudaMalloc temp torch.empty(1024, 1024, devicecuda) # 微秒级而非毫秒级这也解释了为何 PyTorch 的torch.cuda.memory_allocated()与torch.cuda.max_memory_allocated()数值不同allocated 是当前使用量max 是峰值池可能持有比当前使用更多的内存。基于剖析的优化Profile-Guided Optimisation不要盲目优化。先剖析定位瓶颈优化它再重新剖析。用 roofline 模型判断瓶颈在内存还是计算内存受限算力强度低优化数据布局SoA、融合内核、降低精度、预取。计算受限算力强度高用 Tensor Core、提高并行度、用更快的指令FMA。延迟受限并行度不足提高占用率、减少寄存器使用、启动更多线程。大多数 ML 工作负载是内存受限的。由此得出一个反直觉的结论更强的 GPU更多 FLOPS往往无济于事更快的内存HBM3 vs HBM2e帮助更大。这正是 A100→H100 升级的意义不止于 FLOPS 的原因——H100 的内存带宽也是 2 倍。NVIDIA GPU 代际一览代际年份关键创新AI 相关性Pascal (P100)2016HBM2、NVLink第一代严肃的深度学习 GPUVolta (V100)2017Tensor Core混合精度矩阵乘开启 FP16 训练125 TFLOPS TF32Ampere (A100)2020TF32、稀疏性、第三代 Tensor Core312 TFLOPS TF322:4 结构化稀疏Hopper (H100)2022Transformer EngineFP8、HBM3989 TFLOPS FP8动态精度切换Blackwell (B200)2024第二代 Transformer Engine、NVLink 52.5 PFLOPS FP4多晶粒设计以上性能数字为文档所述的设计指标实际可达性能受散热、供电与工作负载形态影响。Tensor Core是专用矩阵乘单元。单条 Tensor Core 指令在一个周期内完成 4×4 矩阵乘D A×B C而普通 CUDA 核心需要 64 次 FMA。Tensor Core 正是混合精度训练FP16 计算、FP32 累加如此之快的硬件基础。Transformer EngineHopper 及以后在单个层内部动态切换 FP8 与 FP16只在需要处使用更高精度。它在不牺牲模型质量的前提下最大化吞吐是专为 Transformer 架构attention MLP当代 AI 的主流形态设计的。第 17 章量化文档详细说明了 FP8 前向/梯度分工E4M3/E5M2与动态损失缩放如何支撑 FP8 训练。编码任务用 nvcc 编译下面三个任务是原文档的实战部分依次覆盖内核编写与数据传输、共享内存分块、warp 发散三个主题。全部可直接保存为.cu文件编译运行。任务 1CUDA 版 ReLU计时包含数据传输该任务训练内核编写、cudaMalloc/cudaMemcpy并直观感受 host↔device 传输瓶颈。// task1_relu.cu // Compile: nvcc -O3 -o task1_relu task1_relu.cu #include stdio.h #include stdlib.h #include cuda_runtime.h __global__ void relu_kernel(const float* input, float* output, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { output[idx] input[idx] 0.0f ? input[idx] : 0.0f; } } int main() { const int N 1 24; // 约 1600 万元素 size_t bytes N * sizeof(float); // 分配主机内存 float* h_input (float*)malloc(bytes); float* h_output (float*)malloc(bytes); for (int i 0; i N; i) { h_input[i] (float)(i % 100) - 50.0f; // 正负混合 } // 分配设备内存 float *d_input, *d_output; cudaMalloc(d_input, bytes); cudaMalloc(d_output, bytes); // 为完整流水线计时拷入 GPU、计算、拷回 cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); cudaMemcpy(d_input, h_input, bytes, cudaMemcpyHostToDevice); int block_size 256; int grid_size (N block_size - 1) / block_size; relu_kernelgrid_size, block_size(d_input, d_output, N); cudaMemcpy(h_output, d_output, bytes, cudaMemcpyDeviceToHost); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0; cudaEventElapsedTime(ms, start, stop); // 验证 int errors 0; for (int i 0; i N; i) { float expected h_input[i] 0.0f ? h_input[i] : 0.0f; if (h_output[i] ! expected) errors; } printf(Time (including transfers): %.2f ms\n, ms); printf(Bandwidth: %.1f GB/s\n, 2.0 * bytes / ms / 1e6); // 读 写 printf(Errors: %d / %d\n, errors, N); cudaFree(d_input); cudaFree(d_output); free(h_input); free(h_output); return 0; }要点cudaEvent用于在 GPU 侧精确计时带宽按读一次 写一次共 2×bytes 计算。ReLU 是典型的内存受限算子——实测带宽会接近硬件极限而算力利用率极低这印证了roofline 模型的判断。任务 2共享内存分块矩阵乘对比朴素版该任务训练共享内存、__syncthreads并实证分块为何重要。// task2_matmul.cu // Compile: nvcc -O3 -o task2_matmul task2_matmul.cu #include stdio.h #include cuda_runtime.h #define TILE 16 // 朴素矩阵乘每个线程计算 C 的一个元素 __global__ void matmul_naive(const float* A, const float* B, float* C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row N col N) { float sum 0.0f; for (int k 0; k N; k) { sum A[row * N k] * B[k * N col]; } C[row * N col] sum; } } // 分块矩阵乘用共享内存减少全局内存访问 __global__ void matmul_tiled(const float* A, const float* B, float* C, int N) { __shared__ float sA[TILE][TILE]; __shared__ float sB[TILE][TILE]; int row blockIdx.y * TILE threadIdx.y; int col blockIdx.x * TILE threadIdx.x; float sum 0.0f; for (int t 0; t (N TILE - 1) / TILE; t) { sA[threadIdx.y][threadIdx.x] (row N t*TILEthreadIdx.x N) ? A[row * N t*TILE threadIdx.x] : 0.0f; sB[threadIdx.y][threadIdx.x] (t*TILEthreadIdx.y N col N) ? B[(t*TILE threadIdx.y) * N col] : 0.0f; __syncthreads(); for (int k 0; k TILE; k) sum sA[threadIdx.y][k] * sB[k][threadIdx.x]; __syncthreads(); } if (row N col N) C[row * N col] sum; } int main() { const int N 1024; size_t bytes N * N * sizeof(float); float *d_A, *d_B, *d_C; cudaMalloc(d_A, bytes); cudaMalloc(d_B, bytes); cudaMalloc(d_C, bytes); // 全 1 初始化便于验证C 应全为 N float* h_A new float[N*N]; for (int i 0; i N*N; i) h_A[i] 1.0f; cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_A, bytes, cudaMemcpyHostToDevice); dim3 block(TILE, TILE); dim3 grid((NTILE-1)/TILE, (NTILE-1)/TILE); // 基准测试朴素版 cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); for (int i 0; i 10; i) matmul_naivegrid, block(d_A, d_B, d_C, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float naive_ms; cudaEventElapsedTime(naive_ms, start, stop); // 基准测试分块版 cudaEventRecord(start); for (int i 0; i 10; i) matmul_tiledgrid, block(d_A, d_B, d_C, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float tiled_ms; cudaEventElapsedTime(tiled_ms, start, stop); double gflops_naive 2.0 * N * N * N * 10 / naive_ms / 1e6; double gflops_tiled 2.0 * N * N * N * 10 / tiled_ms / 1e6; printf(Naive: %.2f ms, %.1f GFLOPS\n, naive_ms/10, gflops_naive); printf(Tiled: %.2f ms, %.1f GFLOPS\n, tiled_ms/10, gflops_tiled); printf(Speedup: %.1fx\n, naive_ms / tiled_ms); cudaFree(d_A); cudaFree(d_B); cudaFree(d_C); delete[] h_A; return 0; }要点dim3允许二维网格/块对应blockIdx.y、threadIdx.y两版用同一 grid 与 block 配置对比公平迭代 10 次取均值以摊薄启动开销。运行后会看到分块版对朴素版的明显加速——这正是矩阵乘法算力强度高、值得投入共享内存优化的依据。任务 3演示 warp 发散编写一个让同一 warp 内线程走不同分支的内核与无分支版本对比。// task3_divergence.cu // Compile: nvcc -O3 -o task3_diverge task3_divergence.cu #include stdio.h #include cuda_runtime.h // 差warp 发散——奇偶线程走不同路径 __global__ void divergent_kernel(float* data, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { if (idx % 2 0) { data[idx] data[idx] * 2.0f 1.0f; } else { data[idx] data[idx] * 0.5f - 1.0f; } } } // 好无分支——所有线程执行同一条指令 __global__ void branchless_kernel(float* data, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { float scale (idx % 2 0) ? 2.0f : 0.5f; float offset (idx % 2 0) ? 1.0f : -1.0f; data[idx] data[idx] * scale offset; } } int main() { const int N 1 24; float* d_data; cudaMalloc(d_data, N * sizeof(float)); cudaMemset(d_data, 0, N * sizeof(float)); int block 256, grid (N block - 1) / block; cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); // 发散版 cudaEventRecord(start); for (int i 0; i 100; i) divergent_kernelgrid, block(d_data, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float div_ms; cudaEventElapsedTime(div_ms, start, stop); // 无分支版 cudaEventRecord(start); for (int i 0; i 100; i) branchless_kernelgrid, block(d_data, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float nodiv_ms; cudaEventElapsedTime(nodiv_ms, start, stop); printf(Divergent: %.2f ms\n, div_ms / 100); printf(Branchless: %.2f ms\n, nodiv_ms / 100); printf(Speedup: %.2fx\n, div_ms / nodiv_ms); cudaFree(d_data); return 0; }要点发散代价来自 SIMT 串行化而条件运算可重写为比例 × 值 偏移的统一公式。现代编译器有时会自动做这类转换但手工重写更可控。实测的速度差异会直观展示在 GPU 上控制流重写为数据流往往比分支更快。从 CUDA 出发本章的抽象阶梯与后续延伸写完这三个任务后你已经站在了能写、能调 CUDA 内核的起点。整个第 16 章的抽象阶梯在此收拢C 内建函数第 2-3 讲NEON/AVX最低层、最可控 ↓ CUDA C本文GPU 专用 ↓ Triton / Pallas第 5 讲Python 编写、自动编译为 PTX ↓ JAX / PyTorch最高层、全自动每一层都以控制权换取便利性。Triton 与 TPU 讲明确指出Triton 用约 20% 的精力获得 CUDA 约 80% 的性能——因为它自动处理线程映射、内存合并、共享内存管理但理解本文的 warp、合并、分块、占用率概念是你在 Triton 调优如triton.autotune的 BLOCK_SIZE 选择时能做出正确决策的前提。无论你最终停留在哪一层本文的底层视角都会让你成为更高效的上层使用者——这正是本仓库教材的核心理念。【免费下载链接】maths-cs-ai-compendiumBecome a cracked AI/ML researcher/engineer with this unconventional textbook covering maths, computing, and ML with intuition.项目地址: https://gitcode.com/GitHub_Trending/mat/maths-cs-ai-compendium创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考