编程模型与API：CUDA编程模型回顾

各位同学好，今天我们来聊聊张量加速核性能调优的基础——编程模型与API。这部分内容，说白了就是你在GPU上写代码时，到底怎么组织你的计算任务。我刚开始接触CUDA时，也被Grid、Block、Thread这些概念绕得晕头转向，后来调了几个项目才真正摸透。

Grid-Block-Thread：GPU的“三体”结构

CUDA的编程模型，核心就是三层结构：Grid、Block、Thread。你可以这么理解：

• Thread：最基础的执行单元，每个线程执行一个任务实例。
• Block：一组线程的集合，它们可以共享显存（Shared Memory），也能通过同步指令（__syncthreads()）协同工作。
• Grid：一组Block的集合，一个Kernel函数对应一个Grid。

我个人的习惯是，先想清楚数据怎么分块，再决定Block和Grid的尺寸。举个例子，你要处理一个1024x1024的矩阵，每个线程处理一个元素，那你可以用32x32的Block，这样Grid就是32x32个Block。

为什么会这样？因为GPU的调度单位是Warp（32个线程一组），Block内线程数不是32的倍数，就会有空闲的CUDA Core，说白了就是浪费算力。

WMMA API：wmma::fragment 实战

接下来是WMMA（Warp Matrix Multiply-Accumulate）API。这是NVIDIA在Volta架构引入的，专门用来加速矩阵乘法的。你想想看，传统的矩阵乘法需要手动管理Shared Memory、Tile切分，代码写起来又臭又长。WMMA直接帮你封装好了。

核心就是 wmma::fragment 这个类。它代表一个矩阵分片，可以存储在寄存器中。用法很简单：

#include <mma.h>
using namespace nvcuda;

// 定义分片
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

// 加载数据
wmma::load_matrix_sync(a_frag, d_a, 16);
wmma::load_matrix_sync(b_frag, d_b, 16);
wmma::fill_fragment(c_frag, 0.0f);

// 执行矩阵乘加
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

// 存储结果
wmma::store_matrix_sync(d_c, c_frag, 16, wmma::mem_row_major);

我在项目中遇到过一个问题：WMMA的fragment尺寸是固定的（16x16或8x8或32x8），如果你的矩阵维度不是这些值的倍数，就得做边界处理。嗯，这里要注意，WMMA不支持动态尺寸，所以你得自己补齐。

cuBLAS/cuTLASS库简介

cuBLAS是NVIDIA官方的BLAS库，说白了就是帮你把矩阵运算做到极致优化。你直接调用 cublasSgemm 或 cublasHgemm 就行，不用自己手写Kernel。

但cuBLAS有个问题：它是个黑盒。你没法控制它内部的Tile大小、流水线策略。这时候cuTLASS就派上用场了。

cuTLASS是NVIDIA开源的模板库，它把矩阵乘法拆成了多个可组合的组件：

• Epilogue：后处理（比如加偏置、ReLU）
• Mainloop：核心计算循环
• TiledMMA：基于WMMA的Tile级矩阵乘

我建议你从cuTLASS的示例代码入手，比如 gemm_64x64x64 这个例子。它展示了如何用cuTLASS实现一个高效的矩阵乘法Kernel。

PTX指令级编程入门

最后聊聊PTX。PTX是CUDA的中间表示，介于高级语言和机器码之间。你写CUDA C++代码，nvcc编译器会先编译成PTX，再转成SASS（机器码）。

为什么要学PTX？因为有些优化，高级语言做不到。比如：

• 手动控制寄存器分配
• 使用特殊指令（如 mma.sync）
• 精细化的内存访问模式

一个简单的PTX例子：

// 加载全局内存
ld.global.f32 %f1, [%rd1];
// 执行FMA（乘加）
fma.rn.f32 %f2, %f1, %f3, %f4;
// 存储到共享内存
st.shared.f32 [%rd2], %f2;

我个人习惯是，先用CUDA C++写出正确版本，然后用 --ptxas-options=-v 查看生成的PTX，看看编译器有没有做傻事。比如，我遇到过编译器把循环展开得太狠，导致寄存器溢出，性能反而下降。这时候手动写PTX就能控制展开因子。

知识体系总览

下面这张图展示了本章的核心逻辑：从编程模型到API，再到底层指令，层层递进。

这张图展示了本章的脉络：从顶层的CUDA编程模型，到中间的WMMA和cuBLAS/cuTLASS API，到底层的PTX指令。你想想看，调优的过程其实就是在这三层之间反复横跳——先用高级API快速实现，再根据性能瓶颈下探到更低层次做优化。

好了，这一章的内容就到这里。记住，编程模型是骨架，API是肌肉，PTX是神经末梢。三者结合，才能写出高性能的张量加速核。