operator_task_package/flashattn_task_package/kernel_ops/FlashAttentionKVCacheDecode

FlashAttention KV Cache Decode

本目录保存 FlashAttention paged KV cache GQA decode 算子的所有优化实现和实验产物。

命名规则

• TileLang：tilelang/opt_NNN_<summary>.py
• CUDA MACA：cudamaca/opt_NNN_<summary>.cu
• Triton：triton/opt_NNN_<summary>.py
• 当前 OJ 提交入口：对应语言目录中的 run_kernel.py 或 run_kernel.cu
• 测试：tests/test_<language>_<scope>.py
• 性能驱动：tests/benchmark_<language>_<scope>.py
• Profiling 驱动：tests/trace_<language>_<case>.py

每次优化创建新的编号版本，不覆盖历史版本。编号文件是实现事实来源，优化结论记录在对应的 <backend>_opt.md 中。

工作原则

1. 先按正式题面建立正确性参考和边界测试。
2. 每个候选依次完成正确性、编译资源检查、稳定 A/B 性能测试和必要的 profiler 分析。
3. 原始日志保存在 results/<backend>/，文档中记录结论和原始产物路径。
4. 版本结论使用 retained、rejected、failed 或 inconclusive。
5. 当前容器的显存配额为 16 GB；按用户确认，其计算资源与 64 GB 容器相同。本地性能仍只用于方向判断，正式结论以 OJ 为准。

当前状态

TileLang opt_006_hybrid_shared.py 已通过OJ 14/14，总时间4.001 ms，是当前最佳完整版本；它按静态规格组合opt_004/005的有效shared策略。当前 run_kernel.py 仍与opt_001一致。