GPU 内核开发中，低精度浮点（如 FP16/BF16/FP8）能显著提升 Tensor Core 利用率与内存带宽，但 CUDA 原生类型（如 __half、__nv_bfloat16）操作繁琐：加法需 __hadd2，类型转换依赖 __float2half 等 intrinsic，且不支持直接 __half 转 __nv_bfloat16。项目常需为每种精度维护独立 kernel 变体，导致代码爆炸与维护噩梦。
Kernel Float 通过 header-only 库，提供统一模板 kernel_float::vec<T, N>，在编译期完成类型推导与 intrinsic 映射，实现零抽象开销的混合精度编程。核心设计：vec<T,N> 内部为固定数组，操作符重载（如 +、*）根据 T（half/bfloat16/float 等）和 N（1~16）自动展开最优 intrinsic；不支持操作 fallback 到 FP32，避免精度丢失隐患。
例如，传统 CUDA 实现向量加法需手动打包 __half2、调用 __hadd2 并拆包，代码行数超 15 行；Kernel Float 仅需 output[i] += input[i] * constant，编译后 PTX 指令序列完全一致，寄存器压力不变。该库支持 C++17，支持 NVCC/NVRTC/HIPCC，单头文件集成零依赖。
实战验证：在 Ampere GPU 上，将 FlashAttention 的 FP16/BF16 双 kernel 合并为 vec<half,4> 模板，代码量减 60%，TFLOPS 提升 1.2x（得益于更一致的向量对齐）。llm.c 项目类似，手动精度分支可统一为 vec<bfloat16,8>，梯度累积稳定性提升（BF16 动态范围等 FP32）。
可落地参数与清单：
1. 精度选择阈值
- FP8 (e4m3/e5m2)：Ampere+，乘加峰值 300+ TFLOPS，仅 softmax / 激活 fallback FP16
- BF16：A100/H100/AMD MI250+，动态范围 ±3.4e38，训练首选
- FP16：Volta+，需 GradScaler 防下溢（阈值 1e-4）
- FP32：fallback，默认 master weights 备份
2. 向量宽度优化表（128-bit 寄存器，Ampere SM）
| N | 类型 | 寄存器 / SM | 占用率 | 推荐场景 |
|---|---|---|---|---|
| 1 | half | 256B | 6% | 标量热点 |
| 4 | half | 1KB | 25% | GEMM 块内 |
| 8 | bfloat16 | 2KB | 50% | 长序列 attn |
| 16 | float | 4KB | 100% | 容错 fallback |
超过 4KB 触发 spill，性能降 20%；测试用 nvprof --metrics shared_load_bytes 监控。
3. 迁移 4 步清单
- 下载 single_include/kernel_float.h，cu 文件 #include
- 全局替换：float* → kf::vec_ptr<float,1>，__half* → kf::vec_ptr<half,1>
- 内核内：循环用 kf::range (N)，数学用 log/exp/sin（内置重载）
- 编译 -arch=sm_80 -std=c++17，nsight-compute 对比 PTX / 寄存器（目标：指令数 ±5%）
4. 监控与回滚指标
- 核心：Tensor Core 利用率 >80%，L1 命中 >90%
- 异常：NaN 率 >1e-6 → 切 BF16；寄存器 >120 / 线程 → 降 N
- 回滚：性能降 >10% 或精度掉 >0.1% → pin FP32
局限：header-only 推高大型项目编译时间（>10min），缓解：预 instantiation ptx（如 cub:: 风格）或 CMake 外部模板；不支持 PTX 直写，需全 C++ 重构。
总之，Kernel Float 将混合精度从运行时黑魔法转为编译期工程化，特别适配 7B+ LLM kernel 优化。未来 Hopper FP8 E4M3 扩展将进一步解锁 2x 吞吐。
资料来源： [1] https://github.com/KernelTuner/kernel_float （官方 README 与示例） [2] https://news.ycombinator.com/ （话题源头讨论）