如果某个操作直接报错或者输出跑到了 CPU那就值得深挖一下。脚本在做什么#核心逻辑非常朴素x torch.randn(128, 512, devicecuda) out some_operation(x) print(out.device) # 还在 cuda:0 上吗再配合 warmup 计时前后都torch.cuda.synchronize()以及 Python warning 捕获和错误分类就是一个完整的冒烟测试了。但这里有个容易误解的地方输出 tensor 仍在 GPU 上不等于底层一定调用了高性能 GPU kernel。PyTorch 的设备语义本身就倾向于让 CUDA tensor 的结果留在同一设备上缺少对应 accelerator 实现时常见行为是直接报错而不是悄悄生成 CPU tensor。所以输出还在 cuda 上这个结论的含金量其实比证明该算子有高性能 GPU kernel低得多。覆盖了哪些操作#脚本测了 9 类常见操作大概 55 个检查项类别典型操作可能涉及的 backend线性代数matmul, linear, bmm, einsumrocBLAS / ATen卷积conv2d, depthwise, transpose, poolingMIOpen / ATen归一化batch_norm, layer_norm, group_norm, rms_normMIOpen / ATen激活函数relu, gelu, silu, sigmoid, softmax, mishATen注意力手动 attention, SDPACK / ATen / math fallbackReductionsum, mean, max, argmax, topkATen逐元素add, mul, exp, log, clamp, whereATen形状/索引reshape, cat, gather, embeddingATen损失函数cross_entropy, mse_loss, BCEATen选这些操作的原因很简单它们基本覆盖了主流深度学习模型包括 Transformer里最常见的计算模式。够用但不代表完整 workload 的性能质量。计时怎么做#GPU 操作是异步提交的所以计时前后必须同步for _ in range(3): # warmup out fn() torch.cuda.synchronize() t0 time.perf_counter() for _ in range(20): out fn() torch.cuda.synchronize() elapsed (time.perf_counter() - t0) / 20 * 1000这些时间只能粗略参考。部分检查项会在 lambda 里创建随机输入计时会混入 tensor 分配和随机数生成的成本。要做严肃性能分析应该单独固定输入、用更长的 repeat、配合 profiler。实际跑出来什么样#在 RX 6650 XT (gfx1032) 自编译 ROCm/PyTorch 环境下结果类似Summary: 54/55 returned CUDA/HIP tensors | 1 errors | 1 warnings怎么理解这个结果54 个检查项成功返回了 CUDA/HIP tensor。这是好消息——说明这些操作在当前环境下不会直接崩掉。batch_norm报错了。MIOpen 跑 HIPRTC 编译的时候找不到type_traits这是 C 头文件缺失的问题装上 VS Build Tools 的 C 工具链大概率能修。sdpa能跑但 PyTorch 没启用 memory efficient attention走的是数学 fallback。功能没问题性能不是最优。但这个结果不能写成算子覆盖率 98%——这 54 个只是输出还在 GPU 上不是证明有高性能 GPU kernel。如果你看到有人拿类似的数字说ROCm 在 Windows 上已经完美支持了那个结论是过头的。错误诊断#脚本会对报错做自动分类直接在终端里给出建议。目前能识别的几类BatchNorm MIOpen/HIPRTC error— 最常见的问题。MIOpen 用 HIPRTC 运行时编译 kernel但 HIPRTC 在 Windows 上经常找不到 C 标准库头文件。stderr 里会看到fatal error: type_traits file not found。修复方法是装 VS Build Tools 并确认 HIPRTC 的 include path 配置正确。注意这不是 CPU fallback是编译直接失败。SDPA: no flash attn— PyTorch 编译时没开启 memory efficient attention / FlashAttention。SDPA 会走数学 fallback仍在 GPU tensor 上运行但对长序列、大 batch 的 LLM workload 来说性能差距明显。GPU OOM— 显存不够减小输入尺寸或关掉其他 GPU 进程。MIOpen runtime error— 其他 MIOpen 错误。可以试试清 kernel 缓存rmdir /s /q %LOCALAPPDATA%\miopen或者开MIOpen_ENABLE_LOGGING2看详细日志。这个脚本适合干什么#新搭好 ROCm/PyTorch 环境后快速确认常见操作不会直接报错。ROCm 或 PyTorch 升级后做回归检查。作为 profiler 定位的第一层筛查——先知道哪些操作有问题再深入分析。不适合拿来做算子覆盖率统计或者性能瓶颈最终结论。如果真的想定位性能问题#光看输出 device 是不够的。要搞清楚 LLM 推理为什么慢建议继续做这些用torch.profiler记录 CPU 和 CUDA/HIP activity 的时间分布。把 prefill 和 decode 分开计时——自回归 decode 的瓶颈和 prefill 完全不同。用真实 LLM 的 tensor shape 单独测linear、RMSNorm、RoPE、SDPA、KV cache 路径。开 MIOpen/rocBLAS 日志确认关键操作到底走了哪个库。