Agent 帮我写 CUDA：一个 Softmax Attention 算子的开发全记录
副标题： 从 Naive 到 Tiled 的四步优化——GTX 1660 Ti 上的人工智能辅助算子开发实战
一、引子：写了 9 篇博客，终于引出了重头戏环节：“算子开发”
这个agent系列写到了第 10 篇，覆盖了部署、推理引擎、KV Cache、RAG、Agent 工作机制……作为一个软件工程师，最重要的是让他参数到我们的生产工作流缓解，今天探讨的主题就是人+agent提升工作效率实践演示：算子开发。
这件事其实挺矛盾的：我用大模型做 Agent、跑推理、搭 RAG，但大模型最底层的计算（Attention）是怎么在 GPU 上实现的？如果我自己写一个会怎样？
更关键的是：现在 Agent 能不能帮我写 CUDA？
这篇文章就做两件事：
- 从零开发一个 Softmax Attention 的 CUDA 算子——Naive → Tiling → 向量化加载 → 调优，完整走一遍
- 用 Agent 辅助整个过程——让它写代码、分析性能、提优化方案，我来 review、决策、集成
整篇文章的所有代码和实验数据，来自我们真实跑在 GTX 1660 Ti（6GB） 上的结果。没有模拟、没有估算，每一毫秒都是实测的。
二、实验环境
硬件
| 项目 | 值 |
|---|---|
| GPU | NVIDIA GeForce GTX 1660 Ti |
| 架构 | Turing，Compute Capability 7.5 |
| 显存 | 6 GB GDDR6 |
| 显存带宽 | 288 GB/s（理论峰值） |
| FP32 算力 | 5.44 TFLOPS（理论峰值） |
软件
| 项目 | 值 |
|---|---|
| CUDA | 12.2 |
| PyTorch | 2.x（CUDA 版） |
| 精度 | float32 全精度 |
Agent 配置
整个开发过程中，Agent（Claude Code）负责：
- 生成 CUDA kernel 代码
- 编写编译脚本和测试框架
- 分析性能瓶颈并提出优化方向
- 运行 occupancy 参数扫描
我（人类）负责：
- 给出每个阶段的需求
- review 生成的代码
- 决策采用哪个优化方案
- 整理和分析性能数据
参考基线
我们用 PyTorch 原生的 F.scaled_dot_product_attention 作为正确性参照物：
import torch.nn.functional as F
def baseline_attention(q, k, v):
# B, H, N, D -> softmax(QK^T / sqrt(D)) @ V
return F.scaled_dot_product_attention(q, k, v)
所有自己写的 kernel 都必须通过 torch.allclose(atol=1e-2, rtol=1e-2) 验证。
测试配置
所有 benchmark 使用统一的张量形状配置：
| 参数 | 含义 | 测试值 |
|---|---|---|
| B | batch size（批量大小） | 2 |
| H | num heads（注意力头数） | 8 |
| N | sequence length（序列长度） | 64 ~ 4096 |
| D | head dimension（head_dim，每头的向量维度） | 32 / 64 / 128 |
关于 D： 全文提到的 D 均指 head dimension（head_dim），即 Attention 计算
softmax(QK^T / √D)V中每个头的向量维度。它决定了一次 score 计算的点积长度，也直接影响 shared memory 用量和 kernel 的最优 tile 参数。主流模型中 D 通常为 64 或 128。
四个内核版本总览
这篇文章依次开发了四个版本的 Attention kernel，它们的核心区别在于数据流：
| 版本 | 数据流 | 核心思想 |
|---|---|---|
naive |
Global Memory → Register | 每行 query 一个 warp，K/V 全部从显存直读 |
tiled |
Global Memory → Shared Memory → Register | Block 内多 warp 共享 K/V tile，减少显存读取 |
vec_load |
Global Memory(float4) → Shared Memory → Register | 用 128 位向量指令加载数据，提高显存带宽利用率 |
vec_padded |
Global Memory(float4) → Shared Memory(padded) → Register | Shared memory 加 padding，消除 bank conflict |
三、Step 1：Naive 版本——先让模型跑起来
需求
我给 Agent 的第一个指令很简单：
“写一个 Softmax Attention 的 CUDA kernel。输入 Q/K/V 是 [B, H, N, D] float32。用 online softmax 保证数值稳定。每个 warp 处理一行 query。不要用 shared memory。”
Agent 生成的代码
Agent 直接给出了 naive.cu，核心 kernel 大约 80 行：
__global__ void naive_attention_kernel(
const float* __restrict__ Q,
const float* __restrict__ K,
const float* __restrict__ V,
float* __restrict__ O,
int B, int H, int N, int D
) {
int bid = blockIdx.x;
int row = bid % N;
bid /= N;
int head = bid % H;
int batch = bid / H;
int lane = threadIdx.x;
float* q_ptr = Q + ((int64_t)batch * H + head) * N * D + row * D;
// ...加载 Q 到寄存器, 遍历所有 K/V, online softmax, 写回 O
}
设计要点：
- Grid:
B * H * N个 block，每个 block = 一个 warp（32 线程） - 每个 warp 处理一行 query：遍历所有 K/V 位置，计算 score → softmax → 加权累加
- Online softmax：在单次遍历中同时完成 max 计算、exp 归一化和输出累积，不需要存储中间 S 矩阵
- 纯 global memory：所有数据直接从显存读取，不做 tiling
编译和测试
Agent 同时生成了 Makefile 和 PyTorch wrapper：
# Makefile 关键部分
nvcc --shared -O3 -use_fast_math \
-gencode arch=compute_75,code=sm_75 \
-o libnaive_attention.so kernels/naive.cu
然后用 Python 调用编译好的 .so，与 PyTorch baseline 对比：
测试规模: B=2, H=4, N=64, D=32
最大绝对误差: 8.34e-07
平均绝对误差: 4.52e-08
不匹配元素: 0/16384 (0.0000%)
allclose(atol=1e-2, rtol=1e-2): ✅ PASS
第一版就正确了。 这就是 Agent 辅助开发的第一层价值：写一次性正确的代码。不需要我手写 CUDA 的 grid/block 索引计算、不需要我 debug online softmax 的数值稳定性——Agent 一次性搞定。
四、Step 2：Tiled 版本——让 Agent 分析瓶颈
Agent 的分析
Naive 版本跑通后，我问 Agent：“这个 kernel 的瓶颈在哪？”
Agent 的回答很清晰：
"Naive 版本的每个 warp 都会从 global memory 重复读取整个 K 和 V 矩阵。 有 N 行 query，K 和 V 就被读了 N 次。当 N 很大时（比如 4096），K 和 V 的 global memory 读取量是实际数据量的 4096 倍。**
解决方案：用 shared memory 做 tiling。 一个 block 处理多行 query（而不是一行），block 内的所有 warp 共享一个 K/V tile。这样 K 和 V 只需要从 global memory 读一次。"
Tiled Attention 的核心思想
Tiled Attention 的计算流程如下：
每个 block 加载 Q_tile [TILE_ROWS × D] 到 shared memory
↓
循环遍历 K/V 的 tile:
加载 K_tile [TILE_KV × D] → shared memory
加载 V_tile [TILE_KV × D] → shared memory
↓
block 内 TILE_ROWS 个 warp 各处理一行 query
每个 warp 计算 Q_row · K_tile 的 score
→ online softmax 跨 tile 合并 → 累积到输出
↓
下一个 tile...
↓
所有 tile 处理完毕 → 归一化 → 写回 O
这里最核心的难点是 online softmax 的跨 tile 合并——传统 softmax 需要先算完所有 score 才能算分母，但 online softmax 允许在遍历过程中逐步更新：
for 每个 K/V 位置:
score = dot(Q_row, K[pos]) / sqrt(D)
old_max = max_score
max_score = max(max_score, score)
alpha = exp(old_max - max_score) // 缩放旧输出
beta = exp(score - max_score) // 当前 score 权重
denom = denom * alpha + beta // 更新分母
output = output * alpha + beta * V // 更新输出
不需要等所有 score 算完才开始，后面的结果可以 rescale 到前面的统计量上。
Agent 在这步的作用
Tiled Attention 的实现在 Agent 辅助下一次通过。但让我 review 时，我关注了几个容易出错的地方：
- 跨 tile 的 alpha/beta rescale——逻辑对不对？
- 最后一个不完整 tile 的边界处理——N 不是 TILE_KV 的倍数时？
- __syncthreads() 的位置——有没有 race condition？
Agent 在这三个地方都处理正确了。我的 review 花了大约 5 分钟确认。
正确性测试
测试 3 组配置全 PASS：
测试 1/3 [ B=2 H=4 N= 64 D=32 ] naive: ✅ tiled: ✅
测试 2/3 [ B=2 H=8 N=128 D=64 ] naive: ✅ tiled: ✅
测试 3/3 [ B=1 H=12 N=256 D=64 ] naive: ✅ tiled: ✅
naive 和 tiled 的输出完全一致（max_diff=0），因为两者实现了相同的 online softmax 算法。
五、Step 3：向量化加载与 Bank Conflict 消除
float4 向量化加载
Agent 分析 tiled 版本的 global memory 访问模式后，建议了下一步优化：
“K/V 从 global memory 加载到 shared memory 时，每次只读 4 字节（一个 float）。如果改用 float4（16 字节），一次指令可以加载 4 个 float，减少指令数并利用 Turing 架构的 128-bit 对齐访问。”
改动很集中——只修改 global → shared 的加载部分：
// 改之前：标量加载
for (int i = tid; i < TILE_KV * D; i += blockDim.x) {
K_tile[i] = K[head_base + global_row * D + (i % D)];
}
// 改之后：float4 向量化加载
float4* K_src = (float4*)(K + head_base + kv_start * D);
float4* K_dst = (float4*)K_tile;
for (int i = tid; i < (TILE_KV * D) / 4; i += blockDim.x) {
K_dst[i] = K_src[i];
}
Bank Conflict Padding
Shared memory 有 32 个 bank，跨 stride 访问时如果多个线程访问同一 bank 的不同地址就会发生 bank conflict。当 D 是 32 的倍数时，这个问题特别明显。
解决方案：把 shared memory 的 stride 从 D 改成 D + 4（加 padding），让相邻行的同列元素落在不同 bank：
// 改之前：stride = D
float* K_tile = smem + TILE_ROWS * D;
// 改之后：stride = D + 4（加 padding）
int stride = D + 4;
float* K_tile = smem + TILE_ROWS * stride;
实际效果
两个优化加完后，我们有了四个版本的 kernel：
| 版本 | 描述 | 正确性 |
|---|---|---|
naive |
Global memory only，无优化 | ✅ |
tiled |
Shared memory tiling | ✅ |
vec_load |
Tiled + float4 向量化加载 | ✅ |
vec_padded |
Tiled + float4 + bank conflict padding | ✅ |
12 组测试（4 kernel × 3 组规模）全部 PASS。
数据流对比：四个版本的本质区别
光看代码不理解为什么优化收益这么小，画出数据流就很清楚了：
Naive — 全部从显存直读
Q[行,:] ──→ 寄存器 ─→ dot(Q,K) ─→ softmax ─→ acc(O)
↑ ↑
K[位置,:] ──────────────┘ │
V[位置,:] ──────────────────────────────┘
↑ └→ O[行,:] → 写回显存
└── 循环遍历 N 个 K/V 位置
所有数据 Global Memory → Register 直通，不经过 shared memory。每行 query 独立遍历一遍 K/V，N 行就把 K/V 读 N 遍。但 L2 cache 会缓解大部分重复读取。
Tiled — 引入 shared memory 做中转站
┌─ Q_tile[TR×D] (shared mem) ──→ 各 warp 读自己的 Q 行
Global Memory ──→ ├─ K_tile[KV×D] (shared mem) ──→ 所有 warp 共享读取
├─ V_tile[KV×D] (shared mem) ──→ 所有 warp 共享读取
└─ 循环加载下一个 K/V tile
K/V 先从 global memory 批量加载到 shared memory，block 内所有 warp 共享。K/V 只需从 global memory 读一次，而不是 N 次。代价是多了 __syncthreads() 同步屏障。
vec_load / vec_padded 的数据流和 Tiled 完全一样，区别仅在于 global → shared 这一步用 float4（一次 16 字节 vs 一次 4 字节），以及 shared memory 布局加了 padding 消除 bank conflict。路径没变，只是路更宽、内部不卡了。
六、Step 4：Occupancy 调优——找到最优参数
Tiled 版本有两个关键参数：
- TILE_ROWS：每个 block 处理几行 query（也是 block 内的 warp 数）
- TILE_KV：每轮加载多少个 K/V 位置到 shared memory
它们决定了 shared memory 用量和 block 并发度——直接影响 occupancy。
Agent 跑了一个参数扫描（occupancy_scan.py），测试 TILE_ROWS × TILE_KV 的 9 种组合：
D=64, N=512:
TILE_ROWS TILE_KV Block SMEM Lat(ms)
2 16 64 9KB 9.84
4 16 128 10KB 5.91
8 16 256 11KB 4.92 ← 最优
8 32 256 19KB 5.27
8 64 256 36KB 10.98
TILE_ROWS=8、TILE_KV=16 胜出——256 线程的 block 提供了足够的并行度，而 16 的 TILE_KV 保持了合理的 shared memory 占用。
一个意外的发现
但问题是：最优参数不是对所有场景都适用。
我把同样的扫描扩展到 D=32 和 D=128：
D=32: 最优 TILE_ROWS=8, TILE_KV=32 (SMEM=10KB)
D=64: 最优 TILE_ROWS=8, TILE_KV=16 (SMEM=11KB)
D=128: 最优 TILE_ROWS=8, TILE_KV=8 (SMEM=12KB)
虽然 TILE_ROWS 始终是 8，但 TILE_KV 随 D 增大而减小：D 越大，每个元素占的 shared memory 越多，需要更小的 tile 来维持 occupancy。
这个发现的意义在于：如果一个算子库只设置了固定 TILE_KV，它在不同 head_dim 上的性能会有显著差异。 这也是为什么我们最终采用了 D 自适应的编译策略——编译时根据 head_dim 选择不同的 TILE_KV。
七、完整性能对比
这是整篇文章最关键的部分。我们用 D 自适应最优参数，在 GTX 1660 Ti 上对比了 4 个 kernel 在 3 个 head_dim × 7 个序列长度的完整性能。
关于带宽效率： 下面的表不展示 BW 效率——因为我们的数据来自 Python 计时器而非硬件性能计数器，无法获知真实的 global memory transaction 次数。kernel 的 latency 大部分来自计算而非显存访问，BW 效率在这里没有参考意义。
D=32（Tiling 的舒适区）
| N | Naive(ms) | Tiled(ms) | VecLoad(ms) | VecPad(ms) | 最优加速比 |
|---|---|---|---|---|---|
| 64 | 0.11 | 0.10 | 0.09 | 0.09 | 1.17x |
| 128 | 0.39 | 0.30 | 0.28 | 0.28 | 1.40x |
| 256 | 1.20 | 0.94 | 0.88 | 0.87 | 1.38x |
| 512 | 4.57 | 3.75 | 3.42 | 3.41 | 1.34x |
| 1024 | 18.28 | 14.90 | 13.65 | 13.54 | 1.35x |
| 2048 | 71.57 | 59.55 | 54.68 | 54.25 | 1.32x |
| 4096 | 289.52 | 242.24 | 223.26 | 220.16 | 1.32x |
Tiled → VecLoad → VecPad 的增量收益微乎其微（~0.02ms），说明在 D=32 时 shared memory bank conflict 和标量加载并不是主要瓶颈。主要收益来自 Tiling 本身。
D=64（模糊地带）
| N | Naive(ms) | Tiled(ms) | VecLoad(ms) | VecPad(ms) | 最优加速比 |
|---|---|---|---|---|---|
| 64 | 0.12 | 0.14 | 0.12 | 0.12 | 0.94x |
| 128 | 0.45 | 0.37 | 0.34 | 0.34 | 1.35x |
| 256 | 1.32 | 1.48 | 1.29 | 1.33 | 1.03x |
| 512 | 5.02 | 5.85 | 5.04 | 5.05 | 1.00x |
| 1024 | 19.55 | 23.03 | 19.95 | 19.97 | 0.98x |
| 2048 | 77.83 | 91.31 | 79.16 | 80.27 | 0.98x |
| 4096 | 356.36 | 370.74 | 320.96 | 323.29 | 1.10x |
大部分场景加速比 ≈ 1.0x。注意 Tiled 比 Naive 还慢（0.85-0.96x），但加上 VecLoad 后基本拉平——说明 D=64 时标量加载是额外瓶颈，float4 向量化弥补了这部分。
D=128（退化区）
| N | Naive(ms) | Tiled(ms) | VecLoad(ms) | VecPad(ms) | 最优加速比 |
|---|---|---|---|---|---|
| 64 | 0.13 | 0.21 | 0.18 | 0.18 | 0.76x |
| 128 | 0.49 | 0.69 | 0.50 | 0.52 | 0.98x |
| 256 | 1.44 | 2.37 | 1.94 | 1.98 | 0.74x |
| 512 | 5.37 | 9.27 | 7.45 | 7.61 | 0.72x |
| 1024 | 21.33 | 36.75 | 29.55 | 30.20 | 0.72x |
| 2048 | 90.13 | 146.00 | 120.02 | 122.30 | 0.75x |
| 4096 | 362.77 | 581.24 | 476.84 | 486.92 | 0.76x |
Tiled 全面退化 25-45%。 原因：D=128 时每个 block shared memory 占用 ~20KB，每个 SM 最多并发 2 个 block——occupancy 太低，延迟隐藏能力不足，tile 加载和 __syncthreads() 的开销超过了节省 global memory 带宽的收益。
可视化一览
D=32: Tiled ██████████████████████ 1.20x VecPad █████████████████████████████████ 1.35x ✅
D=64: Tiled ████████████████ 0.86x VecPad █████████████████████ 1.00x ≈
D=128: Tiled ██████████ 0.62x VecPad ████████████ 0.76x ❌
为什么优化收益这么小——Roofline 分析
这是整篇文章最关键的教训，也是我在写这篇博客时和 Agent 反复讨论最多的地方。
从 Roofline 模型看，小 N 时确实是访存受限：
对于 Tiled 版本（理论最小访存），attention 的算术强度（AI）= FLOP / 字节数：
AI_tiled = (4 × B × H × N² × D) ÷ (4 × B × H × N × D × 4)
= N / 4
1660 Ti 的 ridge point = 5.44 TFLOPS / 288 GB/s ≈ 18.9。
| N | AI（Tiled） | Ridge | 理论瓶颈 |
|---|---|---|---|
| 64 | 16 | 18.9 | 🟡 Memory-bound |
| 128 | 32 | 18.9 | 🟢 Compute-bound |
| 256 | 64 | 18.9 | 🟢 Compute-bound |
| 4096 | 1024 | 18.9 | 🟢 Compute-bound |
理论预判：N=64 时 Tiling 应该显著有效（memory-bound），N≥128 后收益递减（compute-bound）。
但实测让我们发现了"黑盒"——L2 cache：
Naive 版本的实际 global memory 流量远小于理论值。因为 Naive 顺序扫描 K/V 的访问模式让 GPU 的 L2 cache 发挥了巨大作用：
- N=64 时：单个 (B=2, H=8) 的 K+V 总大小仅 64×32×4×2×16 = 256KB，完全在 1660 Ti 的 1.5MB L2 cache 范围内。Naive "读了 N 遍"的数据，第 2 遍起全是 L2 hit。
- 数据量越大，Naive 要读的遍数越多，但 L2 命中率随数据总量增加而下降，所以 N 越大 Tiling 的收益本应越大——但 N 越大 AI 也越大、kernel 逐渐变成 compute-bound，Tiling 省带宽的意义同时也在减弱。
这两个效应互相抵消，结果是 在任何 N 和 D 组合下，Tiling 的收益都不超过 1.4x。
这个结果说明了什么？
| 问题 | 答案 |
|---|---|
| Tiling 有错吗？ | 没错——在更大规模或更旧架构上收益更大 |
| 我们的实现有问题吗？ | 大概率没有——4 个版本趋势一致，各 D 的 Occupancy 扫描正常 |
| 为什么 L2 效果这么好？ | 1660 Ti 的 1.5MB L2 + attention 的顺序访存模式是天然匹配 |
| 结论是什么？ | 优化的代价要小于硬件已经帮你做的，才值得上。 |
硬件比你想象的更聪明。 这是贯穿整篇文章的核心教训——不要假设硬件不做优化，先测量，再决定要不要手工干涉。
一个失败的旁支实验： 我们还试了 ping-pong 双缓冲（把 K/V 缓冲区翻倍来 overlap 计算和加载），结果反而慢了 7-9%。原因很简单——Turing 没有
cp.async，省一次__syncthreads()的收益远小于缓冲区翻倍导致的 occupancy 下降。这件事本身也有教育意义：同一个优化技巧在不同架构上的效果可能完全相反，不测量就上优化就是赌博。
八、Agent 在算子开发中的角色
回到开头的那个问题：Agent 能不能帮我写 CUDA？
能，但需要正确的使用方式。
Agent 擅长的
| 能力 | 我们的体验 |
|---|---|
| 生成一次性正确的代码 | Naive、Tiled、VecLoad 三个版本都是一次编译通过 |
| 提出优化方向 | Agent 主动指出 tiling 可以解决重复读取问题、float4 可以提升带宽 |
| 分析瓶颈 | 能解释 naive 的瓶颈在哪、为什么 tiling 能缓解 |
| 跑参数扫描 | 自动化编写并运行 occupancy scan，比手工快得多 |
| 避免 low-level bug | 索引计算、同步 barrier 位置、边界条件——这些人类容易手误的地方，Agent 很稳 |
不过需要说明：Agent 之所以能这么"稳"，不是因为它的推理能力有多强，而是因为 CUDA 的生态太好了。 NVIDIA 有最完整的编程指南、数十个官方 Sample、GTC 演讲、StackOverflow 上百万个问答——这些全部在大模型的训练数据里。tiling、warp reduce、online softmax、bank conflict 规避，每一种模式都在几十种不同写法中被反复讨论过。Agent 不需要"理解"硬件原理，它只需要在给定约束条件（CC 版本、SMEM 大小）下匹配到训练数据中见过的最优模式。换成一套完全闭源的芯片架构，没有公开文档、没有论坛讨论、没有开源参考实现，Agent 的表现会大幅下降。所以"一次写对"这个能力是有边界的——它只对训练数据覆盖充分的生态成立。
Agent 不擅长的
| 不足 | 例子 |
|---|---|
| 判断优化是否值得 | Agent 不会告诉你"D=128 的 tiling 会退化"——它只会有理有据地推荐 tiling，但最终需要实测验证 |
| 理解硬件的实际行为 | Agent 知道 L2 cache 的理论，但不会预测到 Naive 的 L2 命中率和 Tiled 在 D=128 的 occupancy 退化——它倾向于假设最坏情况 |
| 选择在"复杂度和性能"之间取舍 | Agent 倾向于不断叠加优化，不会自主判断"这些优化已经够用、不值得继续了" |
一个人机协作的建议流程
需求 → Agent 生成代码 → 人类 Review → 编译运行
↓
人类判断"有瓶颈吗？" → 如果否 → 完成
↓ 如果是
问 Agent "瓶颈在哪？" → 给出方案
↓
人类选择方案 → 继续循环
关键决策点永远是"人类判断"——Agent 是很好的执行者，但还不是好的"停止决策者"。它不知道什么时候该停。
现实中的算子开发：这个流程已经是成熟方案
这套"Agent 写代码 + 人类决策方向"的模式，已经在芯片公司大规模实践。他们在自研芯片上做 attention、conv、layerNorm 的算子开发，流程和我们完全一样——输入芯片的架构参数，让 Agent 生成基础代码，人类工程师 review 和调优。区别只在于他们的手册是保密的，而我们的 Turing CC 7.5 参数是公开的。
为什么能成立？因为 CUDA 算子有大量可复用的模式——tiling、online softmax、float4 向量化——这些都在大模型的训练数据里。所以下一家芯片公司面试算子开发工程师时，问题可能不再是"手写一个高性能 Attention"，而是**“怎么让 Agent 在你的芯片上生成高性能 Attention，然后你怎么验证和调优它”**。
九、总结
实验结论
- Naive CUDA Attention 在 GTX 1660 Ti 上加速 tiling 仅 1.3x（D=32）——不要过早优化，先测量
- Tiling 只在 D=32 时有效（1.3-1.4x）——D=64 收益归零，D=128 全面退化
- 最优参数是 D 相关的——TILE_KV 需要随 head_dim 自适应调整
- Attention 在 1660 Ti 上是 compute-bound 而不是 memory-bound——Roofline 模型算完 AI ≈ 42 > ridge 18.9，省带宽的优化方向本身就不对
- Agent 能写出正确、高效的 CUDA 代码——但"该不该优化"的判断还得人来
给开发者的建议
| 场景 | 建议 |
|---|---|
| D ≤ 32 | 用 Tiled 版本，TILE_ROWS=8, TILE_KV=32 |
| D = 64 | 用 Naive 就够了，tiling 没多少收益 |
| D ≥ 128 | 必须用 Naive，tiling 会拖慢 30%+ |
| 不确定时 | 先跑 benchmark + Roofline 分析，再决定优化方向 |
最后的思考
这篇博客和系列前面的文章不同——它没有告诉你"怎么做得更好"，而是告诉你"有时候做得更少反而更好"。
Tiling 是 Attention 算子的经典优化，理论上能省 N 倍的显存访问。但实际在 GTX 1660 Ti 上，硬件已经替你挡了一部分（L2 cache），Naive 版本本身已经相当高效。最终的优化收益远小于理论预期。
这不代表优化没用——而是说明优化的第一步应该是测量，而不是猜测。 没有 benchmark 支撑的优化，本质上只是信仰行为。
这也是 Agent 辅助开发的一个映射：Agent 能给你无限多的优化建议，但它不会替你做 benchmark。测量、判断、决策——这些还是人的事。
但如果你反过来思考，这恰恰是它强大的地方：Agent 把算子开发的门槛从"手写每一行"降到了"review 和决策"。 对于芯片公司、算法团队、甚至个人开发者来说，这意味着同一个团队能覆盖的算子种类和硬件平台可以翻几倍。正如我们在这篇文章中做到的——一个人 + 一个 Agent，一天时间，跑完了 GTX 1660 Ti 上从 Naive 到 Tiled 再到参数调优的完整算子开发流程。这在五年前是不可想象的。
Agent 垂直技术社区，欢迎活跃、内容共建。
更多推荐
5
0- 0
已为社区贡献3条内容
所有评论(0)