这篇不聊概念直接聊我在实际做“DeepGEMM”这类项目时的完整思路和踩坑记录。如果你正打算搞一个面向深度学习场景的高性能GEMMGeneral Matrix Multiplication通用矩阵乘法算子或者你在训练框架里被手写算子折磨过这篇文章应该能帮你少走不少弯路。先给刚接触的朋友一个定位。所谓“DeepGEMM”本质上是把深度学习编译器比如类TVM、类Triton的栈和底层GEMM优化结合起来做成一个“用DSL描述矩阵乘、自动生成高性能GPU内核”的工具。它解决的问题很具体训练和推理框架里标准的cuBLAS类库覆盖不了所有shape和融合算子手写CUDA内核又费时费力而且换一个GPU架构就得重调。DeepGEMM的思路是——把GEMM的“配方”写进编译器让编译器自动帮你生成并调优内核。下面我从项目拆解、核心技术、最小实现到排查经验完整过一遍。1. 先拆项目“DeepGEMM”到底想解决什么问题1.1 名字里藏着的信息量“Deep”指的是深度学习场景不仅仅是一个简单的“deep learning”标签它还隐含了几个硬需求要支持Tensor Core、要能和激活函数/归一化等算子融合、要适配不同精度的训练和推理FP16、BF16、TF32。而GEMM则是整个深度学习计算图的“骨架”操作卷积可以拆成GEMM、Transformer里的Attention QKV投影和线性层全是GEMM、推荐系统里的Dense层还是GEMM。把这两个词组合成“DeepGEMM”通常意味着项目不再满足于“写一个能跑的GEMM”而是要打造一个“面向深度学习场景的GEMM内核生成器/算子库”。我见过不少团队在自制训推框架时卡在这一层上层图优化做得再好下层GEMM打不满GPU延迟和吞吐直接崩盘。DeepGEMM这类项目就是来解决这个断层的。1.2 方案选型为什么不能让框架直接调库有人可能会问cuBLAS已经很快了为什么不直接用这个问题的答案做过算子融合的人都会苦笑。其一cuBLAS的API粒度太粗它不知道你的后续算子是什么你只能在H2D/D2H或者kernel之间来回切换平白增加显存读写。其二cuBLAS不开放某些底层的tile调度策略遇到非标准shape比如batch很小但seq很长、或者head维度和GEMM的K维纠缠在一起时它未必给你最优的kernel配置。所以我倾向于把DeepGEMM设计成一个“半编译半运行时”的组件上层用编译器或者接近编译器的DSL解析层描述GEMM的tiling、流水、精度策略下层生成CUDA C代码再由NVCC编译成SASS。这样既保留了对底层的完全控制又能利用编译器的自动调优autotuning去搜索block size、pipeline stage、swizzle模式这些超参数。在真正动手前我习惯把三套方案摆在一起做取舍像下表这样方案灵活性性能上限开发成本适用场景直接封装 cuBLAS低中受库版本和API限制低快速上线、shape 规整的模型CUTLASS 手写模板高高高需要极致性能、团队有底层能力DeepGEMM 式的编译器生成高高中高融合算子多、多架构适配需求强实测下来方案三在“灵活性和性能上限”上最均衡。它最核心的思路是把矩阵乘法优化专家的知识tiling、片段切分、共享内存缓存、流水线重排固化成编译器后端规则而不是固化在一个不可变的内核文件里。1.3 项目边界不是所有GEMM都要自己写一开始容易犯的错是野心太大什么shape都想覆盖。我建议把边界划清楚第一优先是训练推理中高频出现的“规则GEMM”比如M1~128推理场景、M512~8192训练场景、N固定为模型隐藏层宽度、K可变。第二优先是融合场景比如GEMMGELU、GEMMResidualLayerNorm。第三优先才是一些稀疏或者特殊布局的GEMM。把边界定下来编译器生成规则的复杂度就骤降。我最初连“M3”这种秃头shape都想高效支持结果自找麻烦。后来学乖了规则范围内的shape走DeepGEMM极端shape直接fallback到库函数绝不让编译器在这类case上做无谓搜索。2. 核心技术点搞懂这四块你就能改内核2.1 先理解GEMM为什么是“算术强度之王”GEMM的计算量是 2MNK访存量是输入输出数据量MK NK M*N。当M、N、K都比较大时计算密度远大于访存密度。用一句话说GEMM是一个计算受限compute-bound的算子。这个看似简单的结论决定了DeepGEMM内核设计的方向要把数据从全局显存里搬一次然后尽可能在片上算多笔。所以内核里的核心动作不是“计算”而是“调度搬运”。以典型的Transformer形状为例假设M4096、N4096、K4096FP16输入。总计算量约137 GFLOPs输入输出显存总量不考虑写回约100 MB。现代GPU的L2带宽通常在几TB/s量级算下来如果只访存一次理论最短访存时间是几十微秒计算时间在满频Tensor Core下也是几十微秒量级。这意味着只有在“访存和计算完全重叠”的前提下才能逼近硬件极限。这也是为什么DeepGEMM的生成内核里我特别看重一个设计——double buffering双缓冲。用共享内存和寄存器分别做“当前计算”和“下一块数据加载”的流水线让数据的搬运和矩阵乘的算子在时间上完全错开。如果编译器生成的代码没有这一步大概率跑不到峰值的50%。2.2 Tensor Core现代GEMM的性能引擎如果还是只靠CUDA Core的FFMA指令去算GEMM那即使调度做到极致也吃不满当代GPU的算力。DeepGEMM这类项目一定绕不开Tensor Core的使用。有一件事我特别想说清楚Tensor Core上的“矩阵乘”和我们普通写的“矩阵乘”并不完全等价。它是以“warp级矩阵分片”为单位操作的一个warp32线程一起负责一个大的矩阵片比如16x16x16每个线程只持有这个矩阵片的一部分数据。用CUDA的WMMA API写起来大概是#include mma.h using namespace nvcuda; // warp-level GEMM: D[m16n16k16] A[m16k16] * B[k16n16] wmma::fragmentwmma::matrix_a, 16, 16, 16, half, wmma::row_major a_frag; wmma::fragmentwmma::matrix_b, 16, 16, 16, half, wmma::col_major b_frag; wmma::fragmentwmma::accumulator, 16, 16, 16, float c_frag; wmma::load_matrix_sync(a_frag, ptr_a, lda); wmma::load_matrix_sync(b_frag, ptr_b, ldb); wmma::mfma_sync(c_frag, a_frag, b_frag, c_frag); wmma::store_matrix_sync(ptr_c, c_frag, ldc, wmma::mem_row_major);这里有几个坑。第一fragment的数据布局不是直观的“线程0拿第0行”而是按硬件的load/store指令排好的你只需要保证load和store时给对指针和leading dimension即矩阵的“一行”有多少元素中间的计算交给硬件。第二累加器。Tensor Core的累加通常是FP32哪怕输入是FP16这也避免了累加过程中的精度漂移。第三如果你要手动调优还是要看一眼PTX/SASS里是否真的生成了HMMA指令而不是退化成FFMA。在出手写CUDA内核之前我强烈建议先用WMMA API跑通功能再去考虑ldmatrix这种更高阶的指令级优化。DeepGEMM项目的第一阶段完全可以用WMMA API作为生成器的后端。等你想榨干最后一点性能时再换成直接内联PTX的ldmatrix mma指令。2.3 Block级调度怎么把大矩阵切成块、塞进共享内存GPU调度GEMM的基本套路是一个Block负责输出矩阵的一个矩形象限例如128x128的输出块所有Block按照某种启发式顺序遍历整个输出矩阵。决定输出块大小的逻辑很简单——让共享内存装下对应的A块和B块。共享内存总容量有限通常几十KB到一百多KBA块和B块加在一起不能超过上限。输出块越小共享内存压力越小但全局访存的复用率越低输出块越大复用好但可能装不下或者占用过高导致低占用率。实际做DeepGEMM生成器时我一般会设置几个候选组合比如128x128输出块k块32双缓冲64x64输出块k块64单缓冲128x256输出块k块16双缓冲然后让编译器跑一遍自动调优选最优。不建议写死一个尺寸因为不同shape下最优配置不一样。比如K很小的时候k块64就不会比k块16好反而多占共享内存。还有一个细节——共享内存里的bank冲突。共享内存在硬件上被分成32个bank每次访存宽度是4字节。如果多个线程同时访问同一个bank内的不同地址就会触发多路冲突大大降低带宽利用率。经典解法是“padding”也就是在矩阵每一行末尾多分配一个元素让不同线程的地址错开。DeepGEMM的代码生成器里我会专门加一个“swizzle模式”的参数自动生成带padding或者按特定pattern打乱数据排布的加载代码。这块如果直接手写每次都要算半天而交给编译器生成的收益就非常明显。2.4 编译器侧用IR和自动调优来生成内核有了上面的技术点接着就要把它们“编进化”。在DeepGEMM项目里我习惯把内核描述拆成三层算法层描述GEMM的tiling结构、数据类型、累加方式、是否有融合算子。这里可以用非常简单直接的DSL比如gemm(128, 128, 32, dtypehalf, accumfloat, fusegelu)调度层描述共享内存的swizzle策略、pipeline的stage数量、是否用cp.async异步拷贝、warp布局。这一层是自动调优重点搜索的空间。生成层把这些配置翻译成CUDA C再交给NVCC。自动调优的搜索空间设计也很有讲究。我通常先固定算法层的tiling只搜索调度层的关键参数否则搜索空间爆炸。比如先跑一个baseline再开一个2-3轮“坐标下降”式的搜索先搜block大小再搜pipeline stage再搜swizzle。每次改动只调一个维度记录performance下一轮在这个基础上继续。这种搜索策略不一定找到全局最优但能在可接受时间内稳定拿到接近最优的结果。还有一个容易忽略点编译器优化选项。生成出来的CUDA代码要确保NVCC开启 -O3、--use_fast_math注意精度损失、以及针对目标架构的设置。我在项目里会把“编译器参数”也纳入配置管理不然你在本地跑得好好的换个机器忘了加flag性能能掉30%以上。3. 一起动手从零写一个最简DeepGEMM内核3.1 环境准备与最小工程结构我推荐从CUDA 12.x开始直接支持WMMA和cp.async特性。你的机器上需要NVIDIA驱动、CUDA Toolkit、以及nsight compute后面性能分析会用到。工程结构按下面这样搭一个最小的deepgemm_min/ ├── CMakeLists.txt ├── src/ │ ├── main.cpp │ └── gemm_kernel.cu ├── tests/ │ └── test_gemm.py └── scripts/ └── tuning.py一个最小的CMakeLists.txt大概是cmake_minimum_required(VERSION 3.18) project(deepgemm_min LANGUAGES CXX CUDA) set(CMAKE_CUDA_ARCHITECTURES 80;90) set(CMAKE_CUDA_STANDARD 17) set(CMAKE_CXX_STANDARD 17) cuda_add_executable(gemm_demo src/main.cpp src/gemm_kernel.cu) target_link_libraries(gemm_demo PRIVATE cublas)架构设置一定要带上你实际用的卡比如A100是80H100是90不要只写一个native否则换卡就得重新编译。3.2 用WMMA写的第一个可用内核我们先不考虑性能极致优化先让DeepGEMM“能跑起来”验证正确性。下面的内核是一个benchmark级的GEMM能正确输出结果也能看到WMMA API的用法。// gemm_kernel.cu #include mma.h #include cuda_runtime.h #include cstdio using namespace nvcuda; #define TILE_M 16 #define TILE_N 16 #define TILE_K 16 __global__ void wmma_gemm_kernel(const half* A, const half* B, float* C, int M, int N, int K, float beta) { // 这里把每个warp负责一个 16x16 输出块简单起见不做block级tiling int warp_m (blockIdx.x * blockDim.x threadIdx.x) / 32; int warp_n (blockIdx.y * blockDim.y threadIdx.y) / 32; // 实际应该用更规整的索引方式为了演示先写得直观些 int m0 blockIdx.x * TILE_M; int n0 blockIdx.y * TILE_N; wmma::fragmentwmma::matrix_a, 16, 16, 16, half, wmma::row_major a_frag; wmma::fragmentwmma::matrix_b, 16, 16, 16, half, wmma::col_major b_frag; wmma::fragmentwmma::accumulator, 16, 16, 16, float c_frag; wmma::fill_fragment(c_frag, 0.0f); for (int k 0; k K; k TILE_K) { const half* a_ptr A m0 * K k; const half* b_ptr B n0 * K k; // B列主序所以从C视角看需要按列取 wmma::load_matrix_sync(a_frag, a_ptr, K); wmma::load_matrix_sync(b_frag, b_ptr, K); wmma::mfma_sync(c_frag, a_frag, b_frag, c_frag); } float* c_ptr C m0 * N n0; wmma::store_matrix_sync(c_ptr, c_frag, N, wmma::mem_row_major); }这个内核的最大问题是它没有block级共享内存缓存每个warp都直接从全局内存取A和B数据访存会非常糟糕。但它适合作为正确性基准——如果连这个都跑不对说明你的fragment布局或者索引写错了。3.3 增加Block级共享内存向实用内核迈进真实可用的内核至少要有share memory这一层。我把一个Block设为128个线程负责一个64x64的输出块。每个Block先把A的64x64块和B的64x64块搬进shared memory然后由每个warp处理其中的16x16子块。当一个Block协同工作的时候真正的难点在于“谁是搬运工谁是计算工”。我推荐的做法是用block内前一部分线程统一从global memory拷贝数据到shared memory再执行__syncthreads()保证所有数据就位后所有warp再从shared memory读取。虽然这里没有double buffer性能不会到极致但结构清晰好调试。关键代码如下__global__ void deepgemm_smem_kernel(const half* A, const half* B, float* C, int M, int N, int K) { __shared__ half As[64][64]; __shared__ half Bs[64][64]; int block_m blockIdx.x * 64; int block_n blockIdx.y * 64; // 协作搬运 A 块 for (int idx threadIdx.x; idx 64 * 64; idx blockDim.x) { int row idx / 64; int col idx % 64; if (block_m row M col K) { As[row][col] A[(block_m row) * K col]; } else { As[row][col] __float2half(0.0f); } } // 协作搬运 B 块按行读但B本身是列主序需要看清楚布局 for (int idx threadIdx.x; idx 64 * 64; idx blockDim.x) { int row idx / 64; int col idx % 64; if (col N row K) { Bs[row][col] B[row * N block_n col]; } else { Bs[row][col] __float2half(0.0f); } } __syncthreads(); // 每个warp计算一个 16x16 子块 int warp_id threadIdx.x / 32; int warp_m (warp_id / 4) * 16; // 4个warp一排 int warp_n (warp_id % 4) * 16; wmma::fragmentwmma::matrix_a, 16, 16, 16, half, wmma::row_major a_frag; wmma::fragmentwmma::matrix_b, 16, 16, 16, half, wmma::col_major b_frag; wmma::fragmentwmma::accumulator, 16, 16, 16, float c_frag; wmma::fill_fragment(c_frag, 0.0f); for (int k 0; k 64; k 16) { wmma::load_matrix_sync(a_frag, As[warp_m][k], 64); wmma::load_matrix_sync(b_frag, Bs[k][warp_n], 64); wmma::mfma_sync(c_frag, a_frag, b_frag, c_frag); } wmma::store_matrix_sync(C[(block_m warp_m) * N block_n warp_n], c_frag, N, wmma::mem_row_major); }这个版本已经能跑出正确结果。注意几个细节shared memory里做padding可以避免bank conflict但现在没有加搬运过程中对越界位置填0是很重要的因为深度学习模型里M不一定恰好是64的倍数。3.4 正确性验证与性能测量验证正确性不能只靠眼睛看。我的做法分两步。第一步在Python端用PyTorch生成随机输入和期望输出导出成二进制文件。第二步C端读取二进制跑内核再把结果写回文件和PyTorch的结果对比。对比时的容差要注意FP16输入下GEMM的输出误差在1e-2量级都算正常。用torch.allclose(atol1e-2, rtol1e-2)或者自己写一个相对误差统计不要一上来就用1e-5的绝对精度去卡否则你会被硬件累加顺序差异折磨死。性能测量用CUDA EventcudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); kernelgrid, block(...); cudaEventRecord(stop); cudaEventSynchronize(stop); float ms 0.f; cudaEventElapsedTime(ms, start, stop);麻烦的是要排除第一次的冷启动和显存带宽干扰。我一般连续跑十几次取稳态平均值。如果要做正经benchmark建议直接跑一个“造数据、搬数据、计算”的完整流程然后减去纯搬运时间这才是内核的真实算力。3.5 加入double buffering第一次性能跃升接下来一个关键优化是双缓冲这也是DeepGEMM“自动生成”的第一个典型受益点。双缓冲思路很简单把shared memory分成两块如A_buf[0]和A_buf[1]在计算第k块数据的同时用异步拷贝指令去搬第k1块数据。在Ampere之后可以用cp.async指令直接做异步全局到共享内存的拷贝不需要经过寄存器。这会涉及代码里流水线的复杂排布。我不建议新手直接上双缓冲而是先把单缓冲的版本调正确、调稳定再逐段加双缓冲。很多DeepGEMM项目在双缓冲阶段出的bug不是算法不对而是同步逻辑错了比如忘记在循环末尾重新__syncthreads()来确保下一轮数据拷贝完成。一旦出现这种bug最典型的现象是小矩阵跑得对大矩阵偶尔错而且每次错的位置还不一样。4. 常见问题与排查技巧实录4.1 输出全是垃圾值正确性问题这是我遇到的最高频问题。排查顺序很重要先确认输入数据符合预期。把A、B打印成二进制在Python端读出来比对别信内存里的“应该是对的”。检查fragment的布局假设。很多人会把wmma::row_major和wmma::col_major搞混结果是八竿子打不着。B矩阵在数学上是KxN如果内存存储是行主序GEMM里的B其实是“从行主序的KxN矩阵里取一列作为col_major”这个转换是初学者最容易晕的点。检查leading dimension。如果lda没有对齐到16的倍数Tensor Core的load会直接读到错误地址。如果共享内存版出错把tile size调成和warp计算完全一样的“1个warp一个块”逐段调试。我还有一个比较土但很有效的调试法把wmma计算退化成CPU端的一个小矩阵乘把同样数据喂给CPU对比单个tile的输出。只要定位到是哪一块的输出不对问题范围就缩小一半。4.2 性能上不去跑不满Tensor Core性能上不去先别急着怪编译器。用NSight Compute看一眼几个硬指标sm__pipe_tensor_cycles_activeTensor Core管线活跃占比。如果很低说明计算本身没有占满SM。l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld共享内存bank conflict次数。高的话swizzle没做对。l1tex__throughputL1吞吐。如果这个瓶颈卡住大概率是数据搬运和计算没有重叠好。我见过最典型的性能杀手是shared memory的padding配置错了。当你访问As[row][K_stride]时如果K_stride恰好等于32一个warp的宽度的某个倍数就会造成严重的bank冲突。解决办法是在shared memory每行末尾多分配一个half元素让stride变成奇数。这个优化带来的性能差异可能有30%~50%。另外不要忽视warp调度。如果你让8个warp同时从一个shared memory地址读取同样数据没问题这是广播但如果不同warp各自访问不同的“列”且stride相等bank冲突就会爆发。DeepGEMM的swizzle模块核心工作之一就是避免这种规律性访问。4.3 WMMA的坑为什么有时明明写了HMMA速度还是慢有个常见误解是“用了WMMA就自动用上了Tensor Core”。实际上wmma::load_matrix_sync和wmma::store_matrix_sync本身不是Tensor Core操作它们只是数据搬运。Tensor Core只体现在mfma_sync上。所以如果你的循环里搬运占比很高性能还是会被访存卡住。另外IMHAinteger MMA和HMMA在目标指令上不同。FP16的WMMA在编译后应该出现类似.mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32的PTX指令。如果你用cuobjdump --dump-sass检查发现大量FFMA而不是HMMA说明输入类型发生了隐式转换比如你把half当成了float用。调试时可以用以下命令快速查看SASScuobjdump -sass deepgemm_demo gemm.sass grep -i HMMA\|FFMA gemm.sass | head -20如果真没有HMMA检查你的代码里是不是不小心把half用在普通算术运算中导致编译器回退。正确做法是保持half只在load / mma时出现不要随便和float混用。4.4 精度问题TF32与FP16的取舍在训练场景FP16的精度范围有限经常出现Inf/NaN。DeepGEMM这类项目一般会支持TF32Tensor Float 32。TF32是A100引入的它的尾数只有10位但指数范围和FP32一样所以在数值稳定性上比FP16好很多。代价是计算速度可能只有FP16 Tensor Core的1/2左右具体取决于硬件。我在做推理场景时一般偏好BF16因为动态范围很大不容易溢出。但BF16的精度比FP16还低尾数7位在有些场景下Attention输出误差会放大。所以DeepGEMM生成器里的“精度选择”不应该只对应一个宏定义而是要跟上层算子融合的“敏感性分析”结合。比如GELU这种平滑函数对精度不太敏感但Softmax之后的GEMM对误差就敏感得多。一个小技巧在做精度敏感算子的GEMM时把A和B同时拆成高位和低位做两次Tensor Core乘加即“split-K”的精度增强版。误差能大幅降低代价是计算时间几乎翻倍。这个模式可以作为生成器的可选开关默认关需要高精度时再开。4.5 自动调优搜索空间过大的烦恼一旦让DeepGEMM的编译器自动搜索超参新手很容易陷入“搜索空间爆炸”的泥潭。如果搜索空间里有10个可调维度每个维度取5个候选就会有9万多组合跑一遍能把你等到怀疑人生。我的建议是分阶段。第一阶段只调“block size”和“k-block size”两个维度第二阶段固定这两个最优值再调“pipeline stage”和“线程数”第三阶段才调“swizzle模式”。三个阶段的搜索时间加起来比一次性全空间搜索还短而且因为每一步的评估都建立在较优的基础上不容易被噪声带偏。还要注意评估流程本身的稳定性。GPU频率浮动会影响结果我一般会锁定clocks并做多轮取median。如果发现连续两轮结果差异超过5%优先检查是不是机器上跑了别的任务。5. 朝真正可维护的算子库演进5.1 从“内核生成”到“算子库封装”当DeepGEMM的生成器稳定之后自然要考虑怎么给上层框架用。我不建议直接让模型代码调用CUDA kernel太痛苦。做一个C的算子接口把GEMM、融合激活、LayerNorm打包成一个DeepGEMMOp上层只需要传tensor元数据和计算图配置内部负责生成或选择最合适的内核。我见过一个比较成功的做法是在算子库里内置一个小型“缓存”——同一个shape、同一组融合配置生成出来内核和调优结果直接缓存下来。这样第一次调用可能花几十秒做autotune但后续调用直接命中缓存零额外开销。这个缓存很适合接在训练框架的autotune流程里。5.2 多架构适配与降级策略DeepGEMM的一大优势是多架构支持。我建议在生成器里维护一张“特性表”比如Amperesm_80支持WMMA、cp.async、TF32Hoppersm_90支持WGMMA、tensor memory acceleratorBlackwellsm_100可能还会有新的矩阵指令对于同一份DSL描述编译器根据目标架构选择不同的后端指令。如果目标架构不支持Tensor Core就直接生成FFMA内核虽然慢但至少能跑。这种降级策略在推广框架时特别好用因为用户的卡不可能全都一样。5.3 从GEMM到更广义的矩阵运算当GEMM的生成器成熟后你会发现可以很自然扩展到Strided Batch GEMM多头Attention里的批量矩阵乘、量化GEMMINT8权重和FP16激活混合、以及稀疏GEMM。它们的核心调度框架是同一套只是fragment类型和数据搬运动作不太一样。这方面我踩过的一个坑是在扩展Batch GEMM时直接把batch维当成M维来处理结果在batch数量较大时L2缓存局部性急剧下降。正确做法是让共享内存同时缓存多个batch的块让block调度也考虑batch维的连续性。这个问题在DSL层面体现为“tiling维度不全”我当时整整调了一周才想明白。最后想说的话说实话写DeepGEMM这类项目最大的收获倒不是那个最终的内核性能数字而是彻底搞清楚了“编译器优化”和“手写优化”之间的关系。以前我总觉得手写CUDA内核是“硬核”的象征但实际参与到自动生成的流程后才发现真正的能力是把GEMM的优化经验抽象成规则、把规则交给机器去搜索组合。这个过程需要你既懂Tensor Core的底层细节又懂调度器的搜索策略。对想深入GPU算子的同学来说这是一条非常值得走的路。如果让我再给一个建议的话那就是不要一上来就追着“达到cuBLAS的95%”这种目标跑先把你自己的内核在正确性和稳定性上打磨好再谈性能。一个在常见shape下稳定跑到cuBLAS七八成、但能任意融合算子的生成内核在实际项目里远比一个“某些shape很强、但换个shape就崩”的硬编码内核有价值。DeepGEMM这条路本质上就是让你有底气对上层说给我一个矩阵乘描述我就能给你一个能用、够快、能融合的内核。