企业官网建设流程全解析

如果你已经在NVIDIA卡上做过几轮LLM推理性能优化，大概率会有一种感觉：瓶颈从来不在算力，而在怎么把数据喂给Tensor Core。无论是MHA里的QKV投影，还是FFN的两个大GEMM，又或者是KV Cache相关的特殊算子，底层都绕不开一个核心问题：如何写出一个在特定GPU上跑满带宽和算力的矩阵乘法。

CUTLASS这套模板库解决的就是这个问题。它不是一个pip install就完事的黑盒，而是一套让GEMM、Conv、Attention算子作者能精确控制每一笔访存的C++模板体系。这篇文章我会从源码结构、CuTe DSL、Collective编程模型、工程能力和AI推理落地几个维度做一次深度拆解。内容稍微长，但看到最后，你应该能自己判断：什么时候该直接用cuBLAS，什么时候该拿起CUTLASS手写。

1. 源码地图：CUTLASS 3.x的分层设计与目录导航

1.1 为什么我不建议你从gemm.hpp开始读源码

很多第一次接触CUTLASS的人，会本能地打开include/cutlass/gemm/device/gemm.h，然后在一千多行模板代码里迷失。这不是你的问题——CUTLASS 2.x的gemm.hpp本来就是面向模板特化设计的，大量逻辑隐藏在platform/和thread/路径下，读起来极其痛苦。

CUTLASS 3.x转型之后，源码结构已经清晰得多。顶层设计上，整个库拆成三大块：

• cute/：CuTe DSL，负责描述张量布局、线程映射和数据搬运，是整个3.x的核心抽象；
• cutlass/gemm/：GEMM相关的高层组合，包括kernel/、collective/、tile_scheduler.hpp等；
• cutlass/epilogue/：输出阶段的collective实现，处理累加器的缩放、偏置、激活函数和写入。

如果你现在拿到3.7的主线代码，我建议的阅读顺序是：先看cute/layout.hpp和cute/tensor.hpp，把Layout和Tensor这两个概念吃透；再去读cutlass/gemm/collective/下面针对sm90的几个Collective实现；最后才回头看gemm_universal.hpp，因为那个文件只是把不同Collective粘到一起的"壳"。

1.2 三个核心层：Kernel、Collective、CuTe

CUTLASS 3.x的架构本质上是三层结构：

第一层是CuTe，它定义了Lambda、Layout、Tensor、MMA_Atom、Copy_Atom这些基础类型。你可以把CuTe理解成一套"数据排布的数学语言"，它不关心你是做GEMM还是做Attention，只负责回答一个问题：给定一个逻辑坐标系，如何映射到全局内存、共享内存或寄存器文件的物理坐标。

第二层是Collective，它在CuTe的基础上封装了具体的计算和搬运策略。比如CollectiveMma处理主循环里的矩阵乘法，CollectiveEpilogue处理累加器写回。这一层会用到sm90的TMA、wgmma，也会用到sm80的cp.async和mma.sync。

第三层是Kernel层。它只做两件事：决定Tile的调度方式（比如Simple、Persistent还是StreamK），然后调用Collective完成实际计算。gemm_universal.hpp里的CUTLASS_DEVICE入口函数，基本就是把这几个组件按顺序排起来。

这种分层最大的好处是：你不需要为了测试一个新的Mainloop去重写整个GEMM。只要CuTe里描述的布局对，Collective能编译过，Kernel层基本不需要动。

1.3 读懂源码前先理解的三个硬件事实

直接读CUTLASS源码之前，必须先确认自己理解下面三件事，不然代码看起来就是天书：

第一，Tensor Core是有固定 shape 的。以sm80的mma.sync.aligned.m16n8k8为例，一条指令一次算16行8列8深，数据必须按特定布局放在线程的寄存器里。CUTLASS里的MMA_Atom就是把这些指令包成一个个"原子"，CuTe的make_tiled_mma再把原子在warp内重复排布，覆盖更大的块。

第二，sm90的TMA（Tensor Memory Accelerator）是一台独立于线程的异步拷贝引擎。它可以通过一个描述符把全局内存里的数据批量搬进共享内存，完全不需要线程一条条访问。CUTLASS 3.x里大量实现在MainloopSm90TmaGmmaWarpSpecialized这样的名字中，核心配置就是"由warp specialized分工，Producer warp负责TMA搬运，Consumer warp负责wgmma计算"。

第三，sm90引入了线程块簇（Cluster），多个CTA可以共享TMEM和DSMEM空间。CUTLASS的ClusterShape模板参数就是在描述这个协作粒度。看懂这三个硬件特性，你再看CUTLASS源码时，就不会总问"为什么要写得这么复杂"了。

2. CuTe DSL原理：Layout代数把"数据排布"变成了可计算对象

2.1 Layout不是数组，是函数：Shape + Stride的世界

大多数从PyTorch或NumPy过来的开发者，对张量的直觉是"多维数组"，而CuTe的Layout完全换了一个心智模型：Layout本质上是一个从逻辑坐标到内存偏移的纯函数。

using namespace cute; // 一个128行64列的row-major矩阵 auto layout = Layout<Shape<_128, _64>, Stride<_64, _1>>{}; // 把ptr包装成一个Tensor，逻辑形状和物理布局绑定 float* ptr = ...; auto tensor = make_tensor(ptr, layout);

Shape<_128,_64>表示逻辑维度是128和64，Stride<_64,_1>表示第0维步长64、第1维步长1。坐标(i, j)对应的内存偏移就是i * 64 + j * 1。这个模型看起来朴素，但它的强大之处在于可以表达各种奇怪的排布：交错数据、pad后的共享内存、按转置读取等。

在CUTLASS里的一切计算，本质上都是Layout的组合与变换。Shared Memory里的一块张量，其实就是一个指针加上一个Layout；寄存器里的一个Tensor，也就是一个array<float, N>绑上一个Layout。把这个想法贯彻到底，整个GEMM就变成了"几个张量之间的搬运和计算"，而CuTe的算法会帮你自动推导每个线程在哪个寄存器上操作。

2.2 布局操作与TiledMMA：怎么把一个"大GEMM"切给warp

CuTe最核心的价值，是把"MMA如何在warp内展开"这个难题，变成了几个Layout操作。

先看一个典型例子。假设你的Tile是128x128，每个warp要算其中一小块。你首先需要local_tile从大Tensor里切出一个子块；然后用make_tiled_mma把一个MMA_Atom和一份warp映射组合起来，得到一个"TiledMMA"。

// sm80上的一条fp16 MMA原子 using MMA_Atom = MMA_Atom<SM80_16x8x8_F16F16F32F32_TN>; // 把原子在warp维度重复8次，覆盖更大的形状 auto tiled_mma = make_tiled_mma( MMA_Atom{}, Layout<Shape<_8, _2>, Stride<_2, _1>>{} // 当前示例中的线程映射 );

这个Layout<Shape<_8,_2>, Stride<_2,_1>>就是CuTe里最神奇的部分：它决定了32个线程如何映射到MMA原子内部的Tile坐标。tiled_mma随后可以通过partition_A、partition_B和partition_C，自动把共享内存里的A、B矩阵和累加器切分成每个线程负责的寄存器片段。

你需要做的事，只剩三步：第一，用local_tile切出共享内存中的Tile；第二，用tiled_mma.partition_*取出当前线程的寄存器张量；第三，调用mma.callback或gemm循环完成累加。背后的坐标推导和线程对应关系全部由CuTe的Layout代数完成。

这也是为什么CUTLASS 3.x的代码看起来比2.x更"数学"：它不在代码里到处写线程ID与行列坐标的if判断，而是通过Layout的计算自动完成映射，几乎找不到需要手工取模的代码。

2.3 进阶：组合、补全、互逆与coalesce这几个操作解决了什么问题

CuTe里有一组非常值得深入研究的Layout操作，读懂了它们，你才算真正理解CuTe DSL。

第一个是composition（布局组合）。它的作用是把两个布局串成一个：先经过第一个布局做一个坐标变换，再经过第二个布局做内存映射。在partition_A等操作里，composition是底层的主力，它能把"从大Tile中取局部块"和"按线程摊位映射到寄存器"这两个步骤无缝衔接。

第二个是complement（补布局）。补布局是生成一个"双射"的关键手段。假设你已经定义了一个从逻辑坐标到寄存器坐标的映射，你还需要一个"补齐"的布局来描述剩余那些线程或内存应该如何排列。很多看起来很绕的源码，其实都是在算某个Layout的complement。

第三个是coalesce（合并维度）。它把相邻的连续维度合并成更大的步长，在分析内存访问是否连续时非常有用。做性能优化时，经常需要检查一个全局内存Tensor经过多次partition之后，在内存地址上到底有没有保持"相邻线程访问相邻地址"，用coalesce一眼就能看出来。

我之前在Jetson Orin上调一个自定义算子时，就因为忽视coalesce，导致全局内存访存路径呈"跳变"状态，带宽利用率只剩三成。后来用CuTe把布局打出来一看，Stride完全不连续，改完Layout定义后性能直接翻倍。

2.4 CuTe不止做GEMM：它是"张量计算的编译器基础设施"

很多读者容易把CuTe局限在"GEMM的辅助工具"这个印象里，这是对CuTe DSL最大的误解。

CuTe实际是一个通用的"张量映射算法库"。它并不在意你的计算到底是GEMM还是FlashAttention，也不在乎你最后是调用wgmma还是普通的ldmatrix。只要能描述成"从全局内存搬数据到共享内存，再从共享内存搬到寄存器，再算一次"，CuTe的原子和布局操作就都可以用。

官方仓库里的cute/atom/mma_traits_sm90.hpp、cute/atom/copy_traits_sm90.hpp，分别封装了sm90的MMA指令和TMA拷贝指令。你在自研Attention、MOE、甚至非矩阵类算子时，依然可以复用CuTe来做数据排布，只是把计算替换为你自己的CUDA C++逻辑。

这也是CUTLASS 3.x和2.x一个很大的区别：3.x不再是一个只会做GEMM的库，而更像一个面向GPU算子作者的领域特定语言和基础设施。理解了这一点，你在做AI推理自定义内核时，就能把CuTe用得更顺手。

3. Collective编排层：Mainloop、Epilogue与TileScheduler怎么配合

3.1 从collective_mma.hpp看主循环的四个阶段

在CUTLASS 3.x里，一个GEMM主循环（Mainloop）的完整执行，可以被拆成四个阶段。

第一阶段是全局内存到共享内存的搬运。对sm90来说，这一步通常由TMA发起，只需要几行代码构造一个copy操作。TMA描述符在kernel开头创建一次，后续每一轮迭代只需要更新偏移量，数据就到了共享内存。

第二阶段是共享内存到寄存器的加载。在非sm90路径或特殊布局下，这一步由ldmatrix或普通ld.shared完成；在sm90的wgmma路径下，共享内存数据甚至可以直接被wgmma指令消费，减少了中间寄存器拷贝。

第三阶段是MMA计算。CUTLASS把连续的若干次MMA组合成一个"GMMA"操作，在warpgroup级别执行。CollectiveMma会维护一个寄存器累加器数组，每次迭代把共享内存中的A、B分块乘进去。

第四阶段是累加器的更新与传递。累加器经过多轮K迭代后，最终需要交给Epilogue。这里有一个非常容易被忽略的问题：Mainloop和Epilogue之间，累加器到底是留在寄存器里，还是先写回共享内存，还是通过TMEM中转——这个决策直接决定了性能上限，也是MainloopSm90TmaGmmaWarpSpecialized这类实现中大量barrier和warp specialization代码存在的意义。

理解这四个阶段后，再读collective_mma.hpp就会轻松很多，因为代码的注释和函数命名基本就是围绕这条流水线写的。

3.2 Epilogue不是"写回Global Memory"这么简单

很多人在研究CUTLASS时，会主攻Mainloop，却轻视Epilogue。但对于AI推理来说，Epilogue往往才是最耗时、最需要灵活性的部分。

现代推理算子几乎不会直接输出累加器的原始int32或fp32结果。要么需要乘以一个scale再转成fp16/bf16，要么需要加上bias再做SiLU或ReLU，要么需要做GELU、Softmax之外的更复杂融合。CUTLASS 3.x的CollectiveEpilogue提供了统一的Visitor机制，允许你在输出前插入多个变换步骤。

我见过最典型的一个优化案例是：一个模型的前向里有GEMM加LayerNorm的操作。传统做法是GEMM算完把结果写回全局内存，再读出来做LayerNorm。用CUTLASS的Epilogue则可以直接在累加器上完成部分归约和缩放，然后把结果写回。表面上看只是省了一次全局读写，实测在A100上端到端算子延迟能降20%以上。

所以，在做推理落地时，不要只盯着GEMM本身，认真想想哪些算子可以"融进Epilogue"。这是CUTLASS带来的额外优化空间。

3.3 TileScheduler：StreamK和Persistent调度决定性能下限与上限

GEMM性能不仅取决于单个Tile算得快不快，还取决于整个GPU上的Tile调度得是否均匀。

CUTLASS 2.x的经典实现是"CTA按M、N格子切分，每个CTA算一个Tile"，遇到M很小、K很大的形状时，会有大量CTA结束后SM空转。CUTLASS 3.x引入了tile_scheduler.hpp，里面至少有三类调度器值得关注：

• 普通Tile调度：每个CTA只负责一个固定的输出Tile，实现简单，适合形状规整、M足够大的场景；
• Persistent流水线调度：CTA以Persistent方式驻留在SM上，通过原子索引动态领取下一个Tile，减少了CTA启动和回收的开销；
• StreamK调度：在K维度上把工作切碎，多个CTA并行处理同一输出Tile的部分K，通过原子加做归约。这种调度特别适合小M、大K的Decode场景。

实际推理落地时，我建议不要一开始就追求StreamK。它虽然能解决负载均衡问题，但原子归约和最后阶段的累加会带来额外的精度和调度风险。先用普通调度把基线跑通，再对照Profiling结果决定要不要上StreamK。

选调度器时还有一个很容易忽略的点：TileScheduler的模板参数会改变kernel的全局状态布局。同一个CollectiveMma，换调度器后，共享内存占用和寄存器用量都可能变。所以调度和Tile尺寸必须一起sweep，不要单独调某一个。

4. 工程能力评测：CUTLASS真正拉开差距的地方

4.1 构建系统、examples和测试框架的组织方式

CUTLASS的工程能力在开源算子库里是第一梯队。这不光是"代码能跑"层面的工程能力，而是"大规模模板库可持续维护"的工程能力。

首先是构建系统。CUTLASS使用CMake，但与普通项目不同，它编译的不是一个可执行文件，而是大量测试和example。官方提供了一套CMakeLists.txt和tools/library/scripts，用来生成不同架构、不同数据类型的GEMM配置组合，最后封装成libcutlass.so或libcutlass.a。你甚至可以只生成自己关心的几个kernel，节省大量编译时间。

其次是examples。examples/目录下每个子目录都是一个小而完整的独立示例，从最简单的00_basic_gemm到sm90的50_hopper系列。我第一次接触CUTLASS 3.x时，就是照着examples/47_hopper_gemm_with_collective_builder跑通的。这些example不是玩具，它们会直接调用CUTLASS不同层次的API，告诉你完整的调用链长什么样。

测试框架方面，CUTLASS维护了test/目录，单元测试覆盖了CuTe的布局运算、各个架构下的MMA指令、Collective和Epilogue的组合。更关键的是，CUTLASS的测试里有一个"与cuBLAS对比"的reference模块，你在自研kernel时可以直接复用它的验证逻辑。我建议你把test/unit/cute和test/unit/gemm下的代码当作额外的文档来读，很多API的边界条件写得很清楚。

4.2 Profiling工具链：sweep、JSON、与Nsight Compute的配合

CUTLASS在tools/profiler下自带一个cutlass_profiler，这是我在调GEMM时最常用的工具。它能自动遍历你指定的多个配置组合，跑一轮并输出Kernel时间、TOPS、带宽利用率等指标。例如：

./cutlass_profiler --kernels=sm90_xmma_gemm_f16f16_bf16_bf16_f32_tn_f16_tensor_op_f32 \ --m=4096 --n=4096 --k=4096 \ --m-shape=128,128,256 --n-shape=128,256,512

这个Profiler的价值在于，它能快速帮你建立"不同TileShape、不同ClusterShape下性能如何变化"的全局视图，而不是靠猜。

拿到初筛结果后，再用Nsight Compute（ncu）做单kernel深挖。重点看四类指标：

• 访存指标：SMEM bank conflict、global load命中率、TMA带宽利用率；
• 计算指标：Tensor Core pipe利用率、MMA每周期发射数；
• 调度指标：warp stall原因、barrier等待时间、Occupancy是否被SMEM限制；
• 发射指标：issue slot利用率、Control flow overhead。

CUTLASS的kernel在ncu下有非常好的symbol名和源码行号映射，基本能做到"哪一行代码导致哪个stall"的定位。这一点比很多闭源库或手写kernel强太多。

4.3 源码质量与迭代速度：什么样的社区生态在支撑它

我长期跟踪CUTLASS的commit历史后，一个直观感受是：它其实是在"为NVIDIA硬件做算子落地的工程验证"。每次新卡发布前，CUTLASS都会提前拿到新指令集的支持，并在主线里更新。

3.x主线的源码质量，可以概括为"注释密度高、模板边界清晰、但阅读曲线陡峭"。很多核心模板的注释甚至比代码还长，重点解释了这个模板参数组合对应的是哪个硬件功能。尤其要提的是，CUTLASS对sm90和Blackwell的代码路径会明确标注"这是为哪一代卡优化"，不会默认所有架构走同一份代码。

不过也要说句公道话：CUTLASS的代码风格并不是所有人都会喜欢。因为它大量使用CRTP、tag dispatch、ndexpr和折叠表达式，新手阅读时会有很大的心理负担。但一旦你习惯了模板元编程的表达方式，你会觉得这套体系远比密密麻麻的CUDA C代码更易于维护。

对于AI推理团队来说，与其把CUTLASS看成一个"拿来即用的算子库"，不如把它看成一个"拥有NVIDIA官方维护的、可持久跟随硬件更新的算子开发框架"。这意味着你的自定义内核在下一代GPU出来时，有更大概率通过升级CUTLASS版本而不是重写来获得支持。

5. AI推理落地：从LLM算子到生产后端的完整链路

5.1 PreFill与Decode的两种GEMM特征差异

LLM推理环境下的GEMM，和你在Profile里看到的大矩阵乘法是两回事。如果你直接拿固定形状的大矩阵测CUTLASS性能，实战时大概率会失望，因为真实推理的GEMM形状分布极其不均匀。

PreFill阶段的特点是：M等于序列长度和batch的乘积，可能从几百到几千，K和N对应隐藏层宽度，通常在4096或8192。这个阶段单个GEMM的形状仍然比较大，Tensor Core利用率能撑起来。优化重点放在TMA流水线的stage数量、TileShape与M的匹配，以及如何在多buffer之间隐藏TMA延迟。

Decode阶段则是完全不同的世界。M通常只有1到几十，实际计算退化成GEMV。这种情况下，直接跑一个标准GEMM kernel是极度浪费的。一样的Tile，大部分位置都在搬运空气。这时候你要么修改GEMM的调度，让一个CTA处理多个输出行，要么在K维度上做split和reduce，要么把逻辑转成"批量GEMV"来减少并行浪费。

我在实际项目里，PreFill阶段能稳定拿到cuBLAS的85%到95%性能，但Decode阶段如果不做定制，往往只能拿到50%左右。所以如果你要自研推理算子，第一个要问自己的问题就是：我的场景是PreFill多还是Decode多？两者对kernel的要求差得很远。

5.2 量化、窄类型与现代TC：INT8/FP8在CUTLASS里的用法

AI推理落地绕不开量化。它在CUTLASS里的体现，主要是"窄类型GEMM + 对应缩放因子"的组合。

CUTLASS 3.x对sm90的FP8支持很有代表性。它分别提供E4M3和E5M2两种FP8格式的MMA原子，并在Epilogue里集成了从FP32累加器转成FP8输出的逻辑。实际使用中，你需要额外处理scale对齐，因为FP8的指数范围有限，A和B的scale通常要单独计算，再转化到全局的scale里。

INT8方面，CUTLASS同样保持了对sm80以来Per-Column和Per-Output量化的支持。很多推理框架里的Weight-Only量化（比如AWQ和GPTQ）后端，都依赖这类能力：权重被预先量化和重排，激活在运行时使用，用很小的额外开销完成反量化。

我在整合INT8 GEMM到推理管线时，一个重要的教训是：不要只盯着kernel的TOPS，还要看两头的转换开销。权重量化后如果不在kernel启动前做re-order，kernel内部会频繁命中bank conflict或需要额外的通道shuffle。CUTLASS的CutlassLayout提供了预置的卷积权重重排工具，但很多人根本不知道要用它。

量化模型的精度验证又是另一个坑。建议在集成前，先用一组固定的输入跑一遍reference fp16结果，再跑量化后的结果，拉出每个输出通道的最大绝对误差。如果误差集中某几个通道，大概率是scale选择有问题，而不是GEMM kernel本身的问题。

5.3 决策边界：何时直接用cuBLAS/TensorRT，何时必须自研

很多新读者热情高涨，觉得有了CUTLASS就能自己写GEMM超过cuBLAS。实际上，大多数场景直接调cuBLAS或TensorRT是更理性的方案。我把自己的决策逻辑总结成四个问题：

• 你的算子形状是否足够"非常规"？比如变长序列、稀疏模式、特殊融合逻辑，cuBLAS没有一个直接对应的API。这时候值得自研。
• 你是否需要对多核/多卡做算法级协同？比如跨苏格块的Attention、多查询聚合的MOE路由，这种"算子之上的控制流"用CUTLASS写更自然。
• 你是否需要和已有的推理引擎深度绑定？比如TensorRT-LLM的自定义plugin，你需要在它的框架约束下写算子，这时候自研内核往往比塞进cuBLAS更灵活。
• 你的团队是否有持续的Tuning时间？CUTLASS不是一次写完就完事的，你需要针对不同GPU型号做TileShape和调度策略的适配。如果只做一次性优化，用cuBLAS性价比反而更高。

对我而言，CUTLASS最适合的位置是：为那些cuBLAS覆盖不到、或者融合需求极强的算子提供一个"可控的基础设施"，而不是替代cuBLAS本身。同样，如果你的后端框架已经用了TensorRT-LLM或NVIDIA NIM，它底层可能已经调了CUTLASS或cuBLAS，你再去重写反而会失去框架层面的优化联动。

5.4 部署实战清单：环境检查、精度验证、性能三步

最后给你一份我在生产环境做CUTLASS算子落地时的检查清单，每条都是踩过坑之后总结的。

环境检查这一关，很多人会忽略。编译器版本、CUDA版本和GPU驱动版本必须匹配。比如sm90的某些TMA指令，需要CUDA 12.x以上和对应驱动版本；如果你在Jetson Orin这类嵌入式设备上构建，还要确认JetPack SDK自带哪一版CUDA，是否支持你需要的架构。曾有人问为什么nvidia-smi都报错显示驱动通信失败，这种环境问题解决之前，任何CUTLASS内核都不用谈。

然后是精度验证。不要只看"平均误差小"。在GEMM场景下，最危险的错误是某些K维度分块的累加顺序改变，导致个别输出通道出现无法接受的偏差。我建议用三组数据验证：全零输入、全一输入、随机正态分布输入。全零和全一能快速暴露Index或Layout的错位，随机输入则测试数值范围。

性能验证方面，做两次就够了，但每次都要和正确的基线比。第一次和cuBLAS比，看是否达到期望的量级；第二次和相同CUTLASS内核但不同Tile配置比，做一轮小规模sweep。注意sweep时要固定L2 cache影响，多次跑同一kernel取中位数或最小值，避免被随机噪声误导。

最后是集成阶段。CUTLASS内核一旦被放进推理服务，它就和普通CUDA kernel一样，需要处理stream、多卡、内存池和跨kernel间的依赖。一个容易被忽略的点是：如果多个推理请求并发，kernel的SMEM占用会直接影响同一SM上能并发几个block，这会反过来影响端到端吞吐。所以最终部署前，一定要在满负载并发场景下重新做一次性能验证，不要只信单kernel的Profile报告。

我在实际项目中见过太多次这样的场景：单测一切正常，Profiling数据非常漂亮，一上线并发负载就崩。原因就是TileSize和共享内存分配只考虑了单kernel工况，没有考虑多个block在SM上的并发竞争。

最后说点个人体会

做算子优化的这些年，CUTLASS给我的最大帮助，不是"直接拿来用的GEMM"，而是让我的脑子里有了一套"如何把计算映射到硬件"的思维框架。

CuTe的Layout代数，让我在写任何高性能kernel之前，都会先思考数据的物理排布；Collective分层，让我把"搬运"和"计算"当成两个独立问题去设计；StreamK和Persistent调度，让我明白算法层面的负载均衡和硬件层面的指令流水一样重要。

如果你刚开始接触CUTLASS，不必急着追新鲜的内核实现。先把CuTe的Layout概念吃透，再对照examples/00_basic_gemm和examples/47_hopper_gemm_with_collective_builder把流程跑通。最后再用cutlass_profiler做一轮自己的sweep，对比一下官方默认配置和cuBLAS的差异。

这不只是一次源码阅读，更像是一次对"GEMM到底是怎么在GPU上跑"的完整祛魅。等你丢下gemm.hpp，能用自己的话解释L2 Cache residency、TMA pipeline和warp specialization之间的关系，恭喜你，你已经真正走出"调用框架"的阶段，进入了"设计算子"的阶段。