前言

在高性能矩阵计算（GEMM）领域，模板化算子库已成为 AI 框架与底层硬件之间的关键桥梁。NVIDIA 生态中的 CUTLASS（CUDA Templates for Linear Algebra Subroutines）长期作为行业标杆，而 CANN 开源项目推出的 catlass（CANN Templates for Linear Algebra Subroutines）则代表了面向新型 AI 加速器的另一条技术路径。尽管二者名称相似、目标一致——提供可组合、高性能的 GEMM 模板——但其设计理念、抽象层次与硬件适配策略存在根本性差异。

1. 设计哲学：通用 CUDA vs. 硬件亲和模板

1.1 CUTLASS：基于 CUDA 的通用模板库

CUTLASS 本质上是一个 高度参数化的 CUDA C++ 模板库。它假设底层硬件为 NVIDIA GPU，具备以下特性：

• 统一的 SIMT（Single Instruction, Multiple Thread）执行模型；
• 共享内存（Shared Memory）作为软件管理的缓存；
• Tensor Core 提供固定形状（如 16×8×16）的 MMA（Matrix Multiply-Accumulate）指令。

因此，CUTLASS 的核心抽象围绕 Warp-level GEMM 构建，通过 Mma、Epilogue、ThreadblockSwizzle 等组件组合，生成高效的 CUDA Kernel。其优势在于灵活性——同一套模板可适配多种 NVIDIA 架构（Volta/Ampere/Hopper）。

// CUTLASS 示例：定义一个 FP16 GEMM
using Gemm = cutlass::gemm::device::Gemm<
    cutlass::half_t, cutlass::layout::RowMajor,   // ElementA, LayoutA
    cutlass::half_t, cutlass::layout::ColumnMajor, // ElementB, LayoutB
    float,           cutlass::layout::RowMajor,   // ElementC, LayoutC
    cutlass::arch::OpClassTensorOp,               // 使用 Tensor Core
    cutlass::arch::Sm80                           // 目标 SM 版本
>;

✅ 特点：依赖 CUDA 编译器（nvcc）进行模板实例化与优化，硬件细节通过 arch 命名空间封装。

1.2 catlass：面向专用 AI 加速器的分层模板

相比之下，catlass 并非基于通用并行编程模型（如 CUDA/OpenCL）。根据其 README.md，catlass 的设计目标是：

“通过抽象分层的方式将矩阵类算子代码模板化，从而实现算子计算逻辑的白盒化组装，让算子代码可复用，可替换，可局部修改。”

其核心理念是 “硬件亲和” 与 “白盒化”。catlass 不假设通用线程模型，而是直接映射到目标 AI 加速器的计算单元、片上存储、DMA 引擎与专用指令集。为此，catlass 定义了自己的四层抽象架构：

• Algorithm Layer：定义 GEMM、Conv 等高层算法；
• Tile Policy：指定分块大小（如 M/N/K）、流水级数（Stage）、数据布局；
• Kernel Layer：组织计算与访存流水；
• Hardware Primitive：封装底层指令（如 mma, load, store）。

💡 关键差异：catlass 的每一层均可被开发者显式替换或修改，实现“局部定制”，而 CUTLASS 的底层（如 Warp MMA）对用户基本透明。

---

2. 编程模型与内存访问：显式控制 vs. 隐式优化

2.1 内存层级与数据搬运

CUTLASS：依赖 Shared Memory + L2 Cache

CUTLASS 利用 CUDA 的两级内存：

• Global Memory → Shared Memory（由程序员显式管理）
• Shared Memory → Registers（由编译器自动调度）

数据搬运通过 Iterator 和 PredicatedTileIterator 实现，编译器负责 coalescing。

catlass：显式 DMA 与片上 Buffer 控制

catlass 面向的硬件通常具备独立的全局内存与片上存储（On-Chip Buffer），且无统一地址空间。因此，catlass 显式建模 DMA 传输：

// catlass/include/catlass/gemm/kernel/gemm_kernel.hpp
template<typename TileScheduler>
__global__ void GemmKernel(...) {
    // 1. 声明片上 Buffer
    __shared__ typename TileScheduler::SmemLayoutA smem_a;
    __shared__ typename TileScheduler::SmemLayoutB smem_b;

    // 2. 启动 DMA 引擎预取第一块数据
    dma_engine.load_tile(smem_a, global_a_ptr, ...);

    // 3. 流水循环
    for (int k = 0; k < K_tiles; ++k) {
        // 等待 DMA 完成
        dma_engine.wait();

        // 启动下一块 DMA（重叠计算与传输）
        if (k + 1 < K_tiles)
            dma_engine.load_tile(smem_a_next, global_a_ptr + offset, ...);

        // 执行 MMA 计算
        mma_compute(smem_a, smem_b, reg_c);
    }
}

📌 关键点：dma_engine 是 catlass 自定义的硬件抽象，对应 driver 层暴露的 DMA 控制接口。开发者需手动插入同步点（wait()），确保数据就绪。

2.2 数据布局与向量化

CUTLASS 支持多种布局（RowMajor/ColumnMajor），但向量化由编译器自动推导。

catlass 则强制要求开发者指定向量化宽度，以匹配硬件 DMA 引擎的位宽：

// examples/00_basic_matmul/basic_matmul.cpp
using ElementA = half;
using LayoutA = cutlass::layout::RowMajor;
using VecWidthA = cutlass::AlignedVector<half, 8>; // 128-bit 向量

// 在 TilePolicy 中指定
struct MyTilePolicy {
    static constexpr int kBlockM = 128;
    static constexpr int kBlockN = 128;
    static constexpr int kBlockK = 64;
    using VecA = VecWidthA; // 显式绑定向量化
};

这种设计虽增加开发复杂度，但确保内存访问完全对齐，避免硬件性能惩罚。

---

3. 融合能力与扩展性：随路量化 vs. Epilogue

3.1 CUTLASS 的 Epilogue 机制

CUTLASS 通过 Epilogue 组件支持后处理融合（如 BiasAdd、ReLU、Scale）：

using EpilogueOp = cutlass::epilogue::thread::LinearCombinationRelu<float, ...>;

Epilogue 在 GEMM 结果写回 Global Memory 前执行，但仅限于简单逐元素操作。

3.2 catlass 的“随路量化”与多阶段融合

catlass 更进一步，支持 “随路量化”（On-the-fly Quantization）——在计算流水过程中插入量化逻辑，而非仅在结尾。

例如，在 examples/32_w4a8_matmul/ 中，INT4 反量化与 GEMM 融合：

// include/catlass/gemm/tile/tile_copy.hpp (v1.3.0)
template<typename QuantPolicy>
__device__ void TileCopyWithDequant(
    FragmentAccumulator &accum,
    FragmentQuant &quant_frag,
    float scale
) {
    // 1. 从 INT4 解包为 FP16
    auto dequant_frag = unpack_int4(quant_frag);
    
    // 2. 应用缩放因子
    auto fp16_frag = dequant_frag * scale;
    
    // 3. 与累加器融合（可能触发 MMA）
    mma_pipe.fuse_with_accumulator(fp16_frag, accum);
}

✅ 优势：避免中间结果写回 Global Memory，节省带宽；支持更复杂的融合（如 Per-Token Scale）。

此外，catlass 的 Kernel Layer 可被整体替换，实现 FlashAttention、Sparse MatMul 等定制算子，而 CUTLASS 需通过 Custom Epilogue 或完全重写 Kernel。

---

4. 硬件指令集利用：模板特化 vs. 内联汇编

4.1 CUTLASS：PTX 指令封装

CUTLASS 通过 mma.sync PTX 指令调用 Tensor Core，由 nvcc 编译为 SASS。

// CUTLASS 内部使用
asm("mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 {...}");

用户无法直接控制指令发射。

4.2 catlass：硬件原语模板特化

catlass 将硬件指令封装为 可特化的模板函数，位于 include/tla/（Tensor Linear Algebra）目录：

// include/tla/mma/mma.h
template<typename ArchTag, typename ElementA, typename ElementB, typename ElementC>
struct MmaTraits;

// 特化：针对特定硬件架构的 INT8 MMA
template<>
struct MmaTraits<ArchXYZ, int8_t, int8_t, int32_t> {
    static __device__ void mma(
        const FragmentA &a,
        const FragmentB &b,
        FragmentC &c
    ) {
        // 内联汇编或 intrinsic
        asm volatile("mma.i8.i8.o32 %0, %1, %2, %3"
                     : "=r"(c) : "r"(a), "r"(b), "r"(c));
    }
};

开发者可通过继承或特化 MmaTraits，适配新硬件指令，实现指令级定制。

---

5. 工具链与调试支持

5.1 CUTLASS：依赖 nsight-compute

性能分析依赖 NVIDIA Profiler，调试困难。

5.2 catlass：集成 msDebug 与 Tiling Tuner

catlass 仓库提供 tools/tuner 目录下的 Tiling 自动寻优工具，可搜索最优分块参数：

python tools/tuner/tune_gemm.py \
    --dtype w4a8 \
    --shapes "M:1024,2048 N:1024,4096 K:4096" \
    --output best_config.json

同时，msDebug 工具（需驱动支持）可单步调试 Kernel，查看寄存器与片上内存状态——这是 CUTLASS 无法提供的能力。

⚠️ 注意：使用 msDebug 前需检查驱动版本，详见 docs/tools/msdebug.md。

---

结语

CUTLASS 与 catlass 代表了两种不同的高性能计算范式：

• CUTLASS 是 通用 GPU 编程模型的极致优化，依赖强大编译器与硬件一致性；
• catlass 是 专用 AI 加速器的白盒化模板库，强调硬件亲和、显式控制与可定制性。

对于 CANN 生态开发者，catlass 提供了前所未有的算子开发自由度——从 DMA 调度到指令发射均可干预，从而在特定 shape 下达到 0.98~1.2 倍的标杆性能。而 CUTLASS 则更适合快速原型与跨代 GPU 部署。

随着大模型对定制算子需求激增，catlass 的“分层可替换”设计理念，或将成为下一代 AI 编译器的重要参考。

🔗 相关链接：

• CANN 组织主页：https://atomgit.com/cann
• catlass 仓库地址：https://atomgit.com/cann/catlass