从GEMM到Transformer：手写CUDA算子优化的完整路线与性能调优实践

5. 写在Week4末尾的体会

终于把这一周的东西整理完了。说句实话，这周是我接触CUDA以来最“烧脑”但也最“上瘾”的一周——从最简单的GEMM开始，一路优化到能自己手写Transformer里几个关键算子的kernel，最后还能用Nsight Compute给性能瓶颈“定罪”。整个过程很像在做一道需要反复打磨的算法题，只不过这台“电脑”变成了带几千个核心的GPU。

这篇总结适合几类人看：正在系统学CUDA、想搞懂深度学习算子优化的人；被各种CUDA环境问题折磨到怀疑人生的人；以及准备面试推理引擎、算子优化岗位的同学。我不会上来就贴一堆高深的优化技巧，而是尽量把“为什么要这么做”“踩过的坑长什么样”都讲清楚。文章里所有的数值和现象，都是我这一周在自己机器上实测出来的，配置是RTX 3090 + Ubuntu 22.04 + WSL2，PyTorch 2.8.0 + CUDA 12.1组合包。你机器不同没关系，方法论完全通用。

1. 为什么Week4先死磕GEMM

1.1 GEMM是深度学习的“基本盘”

很多人一上来就想写FlashAttention那种炫酷的kernel，我真心不建议。你在Transformer里看到的绝大多数计算，本质上都能归约到GEMM：

• 全连接层就是GEMM；
• 卷积通过im2col或者隐式GEMM实现；
• Attention里的QK^T和score@V是GEMM；
• FFN的两层线性变换是GEMM；
• 包括多卡张量并行里的矩阵分块，底层也是GEMM的拆分逻辑。

所以GEMM优化是所有算子优化的“基本盘”。把GEMM的共享内存、bank conflict、向量化、双缓冲这套方法论吃透，再去看FlashAttention的tiling策略，你会觉得特别熟悉，因为它的核心就是在处理QK^T和PV这两个大GEMM中间夹了一个softmax。

这周我给自己定的目标是：把一个朴素的GEMM kernel从不到200 GFLOPS的实测性能，优化到接近8 TFLOPS。数字本身不重要，重要的是我真正理解了每一步优化到底在解决什么问题。

1.2 我的GEMM优化路线图：Naive、Tiling、向量化、双缓冲

先说下朴素版本长什么样。假设我们要算 C[M,N] = A[M,K] @ B[K,N]，最直接的想法是每个线程算一个输出元素：

__global__ void gemm_naive(const float* A, const float* B, float* C, int M, int N, int K) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if (row < M && col < N) { float sum = 0.0f; for (int k = 0; k < K; ++k) { sum += A[row * K + k] * B[k * N + col]; } C[row * N + col] = sum; } }

这段代码逻辑完全正确，但性能惨不忍睹。为什么？每个线程要读A的一整行和B的一整列，block里的线程之间没有任何数据复用。比如16x16的block，16行A和16列B，其实总共只涉及16+16=32段数据，但每个线程各自从全局内存取，重复读了16倍以上。全局内存带宽就那么大，时间全部浪费在等数据上。

第一步肯定是共享内存tiling。思想很直白：一个block先把需要用到的A子块和B子块搬进共享内存，然后block内的线程反复从共享内存取数计算。共享内存比全局内存快一个量级，关键是它让数据被“缓存”在了离计算最近的地方。

实现方式也不复杂——定义一个block负责计算一块 TILE_M x TILE_N 的输出，比如32x32。每个线程仍然算一个输出元素，但在算之前，整个block协作把A的一个32x32 tile和B的一个32x32 tile加载到共享内存：

#define TILE 32 __shared__ float As[TILE][TILE]; __shared__ float Bs[TILE][TILE]; int tx = threadIdx.x, ty = threadIdx.y; int row = blockIdx.y * TILE + ty; int col = blockIdx.x * TILE + tx; float sum = 0.0f; for (int k0 = 0; k0 < K; k0 += TILE) { As[ty][tx] = A[row * K + k0 + tx]; Bs[ty][tx] = B[(k0 + ty) * N + col]; __syncthreads(); for (int k = 0; k < TILE; ++k) { sum += As[ty][k] * Bs[k][tx]; } __syncthreads(); } C[row * N + col] = sum;

注意这里的两次__syncthreads()：第一次是确保共享内存里的数据写完了再开始读取计算；第二次是确保所有线程算完了再覆盖下一轮数据。没有这一步会出现典型的数据竞争，算出来的结果在边界处时对时错，而且很难复现，是最恶心的bug之一。

优化到这一步，性能大概能从200 GFLOPS涨到1-2 TFLOPS，但离硬件上限还差得远。接下来的关键操作是：

• 向量化加载：把float换成float4，一次读16字节。这对全局内存和共享内存都有好处，但要求地址16字节对齐，所以通常在K维度上确保长度是4的倍数。做的时候有个小坑：shared memory声明时最好用float4数组而不是float数组，否则编译器很难自动向量化。

• 双缓冲：共享内存的加载是有延迟的，如果一整个block在那里傻等数据算完再加载下一块，计算单元就空转了。双缓冲的思路是在计算当前tile的时候，提前把下一个tile的数据加载到另一块共享内存。在Ampere架构上可以用cp.async直接发出异步拷贝，不占用寄存器也不阻塞线程，这一招对隐藏内存延迟效果极其明显。

• 每个线程算多个输出：把thread coarsening加上，例如每个线程算4x4=16个输出元素。这样做的好处是增加数据复用率，同一个A子块的数据可以被算4次，同时减少block数量，调度开销也更低。

这一套组合拳打下来，我的GEMM实测性能从不到200 GFLOPS提升到了8 TFLOPS左右。虽然离3090的FP32峰值35 TFLOPS还有差距，但从“明显错误”到“能看”，这个跨越是巨大的。

1.3 Bank Conflict与Padding：一个细节拖垮整个kernel

如果只看上面那段tiling代码，你可能觉得单核利用率已经不错了。但我在实测时发现，SM的利用率经常只有20%上下，后来用Nsight Compute一查，罪魁祸首是共享内存bank conflict。

共享内存的硬件结构是把连续的存储空间切分成32个bank，每个bank每个时钟周期只能处理一个访问请求。一个warp有32个线程，如果它们同时访问同一个bank里的不同地址，硬件就必须把这些访问拆成多个周期串行执行，这就叫bank conflict。等于说共享内存带宽被打折了。

最常见的冲突场景就在矩阵转置或者按列访问矩阵元素时。比如Bs[k][tx]这种访问方式，在B以行主序存储时，warp里的32个线程访问的是同一行的不同列——也就是第32列、第33列、第34列……它们落在同一个bank里，于是一个本来一个周期就能完成的load，变成了32个周期。

解决办法特别简单粗暴：给共享内存数组的每一行加一个float的padding。把Bs[TILE][TILE]改成Bs[TILE][TILE + 1]。这样第i行的第j个元素和真实内存地址之间错开了一个float，warp访问同一行的不同列时，索引自然分散到不同bank，冲突就消失了。

这个细节让我明白一个道理：算力虽然重要，但访存路径上任何一个“看似无伤大雅”的设计，都可能让kernel性能掉一个数量级。所以后续每次优化，我都拿着Nvidia的Nsight Compute看一下shared__st_bank_conflicts、shared__ld_bank_conflicts两个指标，比盲调快太多。

2. 从GEMM到Transformer：把注意力机制拆成算子清单

2.1 Transformer的全部关键算子拆解

有了GEMM的底子，再看Transformer就轻松很多。我做的第一件事不是写代码，而是把整个Transformer前向计算拆成一张算子清单，然后按“计算密集”和“访存密集”分类：

• Embedding查表：Gather操作，访存密集，GPU上其实不太划算，更多是内存带宽受限；
• 位置编码（PE）：逐元素sin/cos，访存密集，计算量相对小；
• Q/K/V投影：三个GEMM，实际工程上通常拼成一个大的GEMM做，计算密集；
• Attention Score：Q与K^T的GEMM，再加scale和mask，计算密集；
• Softmax：逐行归一化，访存密集，是Attention里最容易忽略的瓶颈；
• Attention Output：score与V的GEMM，计算密集；
• Output Projection：又一个GEMM；
• LayerNorm：先求均值方差再归一化，访存密集；
• FFN：两个GEMM夹一个GELU激活，第一个GEMM计算量尤其大，因为中间维度通常是4倍。

按时间占比看，在Decoder里GEMM类算子能占到70%到90%的耗时，Softmax、LayerNorm、Embedding加起来占比不高，但它们在端到端推理里会造成频繁的kernel launch和显存读写，不能被忽视。所以我在Week4后半段把重点放在两件事上：手写一个融合的Softmax Attention kernel，以及把位置编码PE在GPU上的正确实现方式搞明白。

2.2 手写一个融合的Softmax Attention Kernel

先说说朴素的Attention实现方式：PyTorch里直接写softmax(Q @ K^T / sqrt(d)) @ V，这是最直观的写法，但性能上有一个很大的问题——QK^T的结果会被写回全局内存，softmax再读一遍，然后乘V再写一遍。这一来一回，中间矩阵[seq_len, seq_len]在HBM和SM之间反复横跳，如果seq_len是4096，光这个中间矩阵就要占64MB，非常浪费。

优化思路有两个方向。第一，把softmax融合到矩阵乘法里，减少中间矩阵的读写；第二，把QK^T、softmax、@V三者合到一个kernel里，这其实就是FlashAttention的雏形。

我先说softmax本身怎么在kernel里写。朴素softmax需要两遍扫描：第一遍找最大值，第二遍算exp和归一化。但开两个kernel显然不划算。在线softmax（online softmax）可以只扫一遍：

假设当前已经处理到第t个元素，维护当前最大值m_t和累积指数和l_t。读入新元素x后：

float m_new = fmaxf(m_old, x); l_new = l_old * expf(m_old - m_new) + expf(x - m_new);

这样一遍扫描就能拿到真实的max和sum。之后再从头读一遍算最终的exp(x - m_max) / l，或者如果后续还要做矩阵乘V，就边更新边累加。这个“边扫描边累积”的思路看起来简单，但它是FlashAttention能够节省访存的核心，没有它每次都要多读一遍输入。

我在自己写的Attention kernel里做了一件更激进的事：把QK^T的scale和mask也融合进去。每个block负责query子块和key子块的一小段点积，算完直接乘上scale，如果某个位置需要mask（比如padding），就直接把score设成-inf，这样exp之后天然是0，不需要额外的分支。然后在线softmax维护running max和sum，同时累加V的输出贡献。整个过程中间不会生成一份完整的[seq_len, seq_len]矩阵，显存占用低很多。

写这个kernel的时候有两个我特别想强调的坑：

一是数值稳定性。如果不减去max直接算exp(x)，当score很大时expf会溢出变成inf，softmax结果直接出NaN。这也是为什么在线softmax里必须维护running max的原因。

二是__syncthreads()的摆放位置。因为softmax的归一化涉及整个row的统计量，而每个row可能跨多个block协作计算，block之间需要同步（或者干脆让一个block处理一整行），所以要么限制每个block处理的行数，要么用grid.sync之类的机制。我前期踩的坑是不同block各算各的，最后归一化因子对不上，结果错得离谱。

最终这个融合kernel在我机器上比“分开写”的版本快了大概2.5倍，中间矩阵的显存占用直接省掉了。这个优化带来的爽感，比GEMM调优还强，因为你能感觉到自己的“算法意识”在起作用，而不只是堆优化技巧。

2.3 Embedding与位置编码（PE）在GPU上的正确打开方式

热词里很多人搜“transformer架构嵌入表示层pe计算”，说明这块对新手确实容易乱。原版Transformer的位置编码公式长这样：

PE(pos, 2i) = sin(pos / 10000^(2i/d_model)) PE(pos, 2i + 1) = cos(pos / 10000^(2i/d_model))

其中pos是token在序列里的位置，i是维度索引，范围是0到d_model/2 - 1。这个编码的本质是给每个位置生成一个“频率不同”的正弦波，让模型能够感知相对位置关系。

在GPU上实现PE最稳妥的做法是预计算一个[seq_len, d_model]张量，一次性拷到显存里，每个batch重复用。别在每次前向的时候在GPU上现算一遍，纯属浪费。用PyTorch预计算的代码几乎任何教程都会给：

import torch import math def positional_encoding(seq_len, d_model): pe = torch.zeros(seq_len, d_model, dtype=torch.float32) position = torch.arange(seq_len, dtype=torch.float32).unsqueeze(1) div_term = torch.exp( torch.arange(0, d_model, 2, dtype=torch.float32) * (-math.log(10000.0) / d_model) ) pe[:, 0::2] = torch.sin(position * div_term) pe[:, 1::2] = torch.cos(position * div_term) return pe

这段代码里有几个细节值得展开。

第一，div_term用的是exp(arange * (-log(10000)/d_model))，而不是直接算10000^(2i/d_model)再取倒数。原因是前者可以预先算好一个长度为d_model/2的向量，避免在GPU上每个位置都算一次pow(10000.0, ...)，这既是数值稳定的写法也是性能友好的写法。

第二，一定要用float32来算，不要用float16。虽然推理时模型权重可能是FP16，但PE里有个1/10000^(2i/d_model)，当i比较大时分母会变得巨大，FP16很容易精度不够，位置编码就是一堆噪声。

如果真要手写一个CUDA kernel来生成PE，也不复杂。每个线程负责一个(pos, dim)坐标，通过dim % 2判断是算sin还是cos。但更高明的做法是把div_term作为常量数组传入kernel，每个线程只需做一次position * div_term[dim / 2]和一次sin/cos调用，把重复计算全部去掉。这个“常量前移”的思想在算子优化里非常通用，几乎任何kernel都能用一遍。

3. CUDA环境与工具链：跑通算子的前提

3.1 多版本CUDA共存：PATH、LD_LIBRARY_PATH和软链接的坑

这一周我的机器上至少同时出现过CUDA 11.8、12.1、12.4三个版本，因为不同项目依赖的框架版本不一样。一开始我被环境问题折磨到想砸电脑，后来才总结出一套相对稳妥的管理方式。

先理解几个命令的区别。nvidia-smi右上角显示的是“当前驱动支持的最高CUDA版本”，这并不等于你系统里实际安装的CUDA Toolkit版本。nvcc --version显示的才是Toolkit的版本。很多人一看nvidia-smi说支持12.1，但nvcc --version还是11.8，就以为环境坏了，其实这完全正常，驱动和Toolkit本来就是两码事。

多版本共存时最忌讳的事情是手改/usr/local/cuda这个软链接。每次切换版本都去改软链接，然后全局改PATH和LD_LIBRARY_PATH，短期内能用，但很容易把正在跑的服务搞挂，而且切来切去迟早会忘记当前到底是哪个版本，排查问题时异常痛苦。

我现在的做法是每个项目一个conda环境，在环境里装对应版本的PyTorch和CUDA相关依赖。完全不需要在系统层面做全局切换。如果一个项目必须用系统Toolkit编译CUDA扩展，再单独指定CUDA_HOME和PATH：

export CUDA_HOME=/usr/local/cuda-12.1 export PATH=$CUDA_HOME/bin:$PATH export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH

如果你管理的是多用户服务器，想让所有用户默认都生效，可以写到/etc/profile.d/cuda.sh里。但注意路径不要写死，最好通过/usr/local/cuda这个软链接来指路，这样后续版本升级时只需要更新软链接。

额外提醒一个常见误区：LD_LIBRARY_PATH里的lib顺序会影响运行时加载的cudart、cublas版本，如果设置错了，即使nvcc --version是对的，运行程序时也可能加载到另一个版本的库。排查这类问题可以用ldd ./your_app看实际加载路径。

3.2 WSL2安装CUDA的正确姿势

现在我大部分调优工作都在WSL2里做，说下正确姿势。

Windows宿主机装好NVIDIA驱动后，WSL2里直接运行nvidia-smi就可以看到GPU信息，显示的还是和Windows相同的驱动版本。很多人误以为这样就算装好CUDA了，结果一编译代码发现nvcc不存在——没错，WSL2里的驱动由Windows共享，但CUDA Toolkit需要自己在Linux里装。

在WSL2里安装CUDA Toolkit，去官网选择Linux、WSL-Ubuntu的安装包就行，安装方式跟原生Ubuntu几乎一样。装的时候建议别用默认的/usr/local/cuda改软链接方式，直接在WSL的~/.bashrc里配置：

export PATH=/usr/local/cuda/bin:$PATH export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH

WSL2下的PyTorch组合，我实测比较稳的组合是Python 3.10 + PyTorch 2.8.0 + CUDA 12.1。去PyTorch官网get-started页面选好配置，复制它会给你生成的那条pip install命令，别自己随便拼。pip install torch默认装的是CPU版本，很多人装完发现torch.cuda.is_available()返回False，十有八九是这里出了问题。

安装完一定要跑这条命令验证：

python -c "import torch; print(torch.__version__, torch.version.cuda, torch.cuda.is_available())"

如果输出True，再跑torch.ones(1).cuda()确认实际能执行GPU计算。我在WSL2踩过最坑的一次是torch.cuda.is_available()返回True，但一跑真实模型就报no kernel image，问题出在PyTorch版本和GPU算力不匹配，这个下面第四节详细说。

3.3 用Nsight Compute确认优化方向

优化kernel最忌讳“盲人摸象”。我见过太多人一上来就疯狂调block大小、换循环顺序，结果性能纹丝不动，因为没有定位到真正的瓶颈。Nsight Compute（命令行工具ncu）就是解决这个问题的。

最基本用法：

ncu --set full ./your_app

它会跑一遍程序，然后给出一大堆性能指标。你不需要全看懂，重点关注几个：

• sm__throughput.avg.pct_of_peak_sustained_elapsed：SM计算单元的利用率；
• dram__throughput.avg.pct_of_peak_sustained_elapsed：显存带宽利用率；
• sm__warps_active.avg.pct_of_peak_sustained_active：活跃warp数，也就是occupancy；
• l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld：共享内存bank conflict计数。

判断方法很简单：如果dram__throughput接近80%以上，而sm__throughput只有20%，说明这个kernel是访存瓶颈（memory-bound），优化重点应该放在减少全局内存读写、增加数据复用、用向量化加载降低请求次数。反过来如果SM利用率高、DRAM利用率低，说明是计算瓶颈（compute-bound），考虑减少冗余计算、提高指令级并行。

我之前的GEMM kernel优化到1 TFLOPS左右就卡住了，怎么调block大小都没用。上ncu一看，dram__throughput已经98%，SM只有11%，根本原因就是朴素实现疯狂读全局内存。后续加上共享内存tiling后，DRAM利用率一下子降到30%，SM利用率升到60%以上，性能自然就上来了。

有一点要注意：ncu在某些环境会报ERR_NVGPUCTRPERM权限错误，一般用sudo ncu运行就能解决。如果你在容器里使用，需要额外开性能计数器权限。这属于环境配置问题，别在kernel代码上白费劲。

4. 新手必看：CUDA算子优化常见问题排查实录

4.1 “no kernel image is available for execution”到底在说什么

这个报错满屏都是，尤其是在新显卡上跑老PyTorch的时候。完整报错通常是：

torch.acceleratorerror: cuda error: no kernel image is available for execution on the device

它的真实含义是：CUDA在加载某个kernel时，找不到能在当前GPU上执行的“kernel image”。GPU的SASS指令是针对具体架构（SM版本）编译的，比如RTX 3060是sm_86，RTX 4060 Ti是sm_89。如果一个库编译时只包含了sm_80或sm_86的SASS，那它在sm_89的卡上就找不到可执行代码。

这就是为什么40系显卡用户经常遇到这个问题：新卡架构太新，老版本的PyTorch wheel里没有包含对应架构的内核。3060、4060 Ti这些卡，如果你装的PyTorch版本过老，也可能报同样错误，因为老版本根本没有sm_86或sm_89的预编译kernel。

解决办法按优先级排列：

1. 升级PyTorch到与显卡架构匹配的新版本，同时选择官方wheel对应的CUDA版本；
2. 如果必须用旧PyTorch，确保CUDA版本足够新（例如CUDA 12.x的wheel通常覆盖更多新架构），但这只是兜底；
3. 自己编译PyTorch扩展时，设置环境变量TORCH_CUDA_ARCH_LIST="8.9;8.6"或对应算力，确保编译产物包含目标架构。

排查之前，先跑一下python -c "import torch; print(torch.cuda.get_device_capability())"确认当前GPU的算力。只有知道目标算力，才能选对PyTorch版本。

4.2 CUDA Samples找不到、nvidia-smi与nvcc版本不一致

“CUDA Samples找不到”也是个高频问题。出现这个情况一般是安装Toolkit时没勾选Samples组件，或者用的是精简安装。正常Samples应该出现在/usr/local/cuda/samples，如果你没找到，可以根据系统直接拉取对应版本的cuda-samples，比如Ubuntu/Debian系：

apt-get install nvidia-cuda-samples

装完的路径通常在/usr/share/nvidia-cuda-samples，不在/usr/local/cuda下面，别傻找。RedHat系用dnf install nvidia-cuda-samples类似。

至于“nvidia-smi显示的CUDA版本和nvcc --version不一致”，我在3.1节说过这是正常的。再强调一次：nvidia-smi右上角是驱动支持的上限，nvcc --version是你当前Toolkit版本。两者不一样不代表环境坏了，只要满足“驱动版本 >= Toolkit所需版本”就行。驱动是向后兼容的，新驱动可以运行旧Toolkit编译出来的程序。

4.3 驱动、Toolkit、PyTorch、cuDNN的版本对应关系

我把这四者的关系整理成一张速查表，方便你排查问题时对照：

这里有个重要的认知：PyTorch的wheel包自己带了一套CUDA运行时依赖，所以即使你系统里没装CUDA Toolkit，也能跑import torch; torch.cuda.is_available()为True的模型推理。只有当你要从头编译CUDA扩展，比如自己写kernel或者编译flash-attn这类库时，才必须装完整的Toolkit。

对于4060 Ti这类新卡，很多教程会问你“4060ti支持的cuda版本”，其实这是问错了方向。40系卡在CUDA 11.8和12.x下都能跑，关键是PyTorch的wheel是否包含对应架构的SASS。版本选择应该以PyTorch官方支持矩阵为主，而不是单独看显卡。

4.4 我的排查顺序与工具清单

这一周我几乎每天都要排查各种环境问题，后来总结了一套固定顺序，能解决90%的“莫名其妙”的错：

1. 先确认GPU裸设备有没有被识别：nvidia-smi。这一步挂了，后面全白搭；
2. 确认PyTorch能不能看到GPU：python -c "import torch; print(torch.cuda.is_available())"；
3. 确认当前设备算力：torch.cuda.get_device_capability()；
4. 执行一个最小GPU算子：torch.ones(1).cuda()；
5. 如果需要编译扩展，再确认nvcc --version和torch.version.cuda是否能对齐，不对齐会有一堆链接报错；
6. 怀疑环境变量有问题时，用env | grep -i cuda和ldd看实际加载库路径。

整个过程我基本不动系统级的配置，所有依赖都锁在conda环境里。如果你发现自己经常要来回切换CUDA版本，我给的建议也很朴素：不要硬改全局环境，用conda或者venv隔离，每个项目一套环境。可能刚开始会觉得多占点磁盘，但换来的是半年不用为环境问题掉头发，这笔账怎么算都划算。

5. 写在Week4末尾的体会

这周最大的收获，反而不是某个具体kernel提速了多少，而是我终于建立了“性能问题要先量化再动手”的直觉。没上ncu之前，我总凭感觉觉得瓶颈在某个循环，花了一晚上重写，结果性能纹丝不动；上了profiler之后，两分钟就发现DRAM带宽利用率满了，算力根本没吃满。

另一个体会是：算子优化的核心永远在“访存”与“计算”的权衡。GEMM的tiling是在用共享内存换全局内存访问次数，FlashAttention是在用重计算换HBM带宽，PE预计算是在用显存换实时计算开销。想通这一层后，再看任何一篇优化论文或者源码，你都能猜到它大概在优化什么、为什么这样优化。

最后分享一个我这周才养成的小习惯：每次改kernel前，先保存上一个版本的性能数据，建议用一个简单的表格记下版本号、改动内容、SM利用率、DRAM利用率、耗时。改完一版就更新一行，不要凭感觉。这个习惯帮我少走了很多回头路——很多时候你觉得新版本更快，其实是测量误差，有个基线数据对比就不会自欺欺人了。

下周我打算沿着这条线继续深入FlashAttention的完整实现，以及Grouped GEMM在多头注意力里的应用。如果你也在学CUDA算子优化，或者刚被某个no kernel image报错逼疯，欢迎直接在评论区聊聊你踩过的坑，我也想知道大家还会遇到什么奇葩问题。