这次适配的模型是 Hy-Embodied-0.5-VLA-RoboTwin， 设备是一张消费级 RTX 5060 Ti 16 GB。
在同一张卡、同一个模型、同一组输入 shape 下： 官方原实现的完整决策延迟是 602 ms，FlashRT 优化后是 56.2 ms； 动作余弦相似度 0.996841（门槛 0.99），连续 5 次重放输出位相等。
602 → 56.2 ms官方原实现 → FlashRT 优化后RTX 5060 Ti 16 GB 单卡 ｜ 一次完整决策 ｜ action cosine 0.996841
602 → 56.2
官方原实现 → 优化后（ms）
154.8 → 56.2
FlashRT 首版 → 交付版（ms）
0.996841
action cosine（门槛 0.99）
5/5
连续重放位相等
01 · WORKLOAD跑的是什么模型，跑在什么设备上一次 observation → 一个 40-token 的相对动作块（10 步 flow denoise），完整链路是：
preprocess(resize / normalize / pad) # 3 路相机 × 6 帧历史 + ViT(27 blocks, K=6, N=196, heads=16, head_dim=72) + merger + [vision | text] prefix permutation + prefill(32 blocks, S_p=181, n_vis=147) + 10-step flow denoise(32 blocks, M=41)
输入是 3 路相机 × 6 帧、480×640 的图像、20 维 state，以及固定随机种子的 noise (1, 40, 32)。 固定 shape，无持久 KV cache；稳态判据是「冷启动后重复同一决策，尾部延迟漂移 < 10%」。 两种输入来源都测过（合成输入 / 真实 RoboTwin 观测），延迟一致—— 单次决策的耗时与输入数值无关。
项
配置
GPU
单卡 NVIDIA RTX 5060 Ti 16 GB，sm_120，36 SMs
显存带宽（实测）
HBM ≈ 360 GB/s，L2 ≈ 1.7 TB/s
主机链路
PCIe x8；实测持续 H2D ≈ 13.7 GB/s
CUDA / 容器
CUDA 13.0 toolkit（nvcc 13.0），Linux 容器
Python / PyTorch
Python 3.12；PyTorch 2.14.0+cu130
算子库
CUTLASS（sm120 tcgen05 / W4A4 NVFP4）、FlashAttention-2（SM80 模板）、pybind11
构建选项
GPU_ARCH=120、FLASHRT_ENABLE_HYVLA=ON、FA2_ARCH_NATIVE_ONLY=ON
编译
必须 -j1：31 GB 内存下 NVFP4 CUTLASS 的 .cu 并行编译会 OOM
02 · RESULTS性能：三张表说清楚2.1 延迟阶梯测量方式：预热 5 次、计时 20 次，每次调用前后都做 torch.cuda.synchronize()，取中位延迟。
实现
一次完整决策延迟（中位）
官方原实现（未做任何 FlashRT 优化）
≈ 602 ms
FlashRT 首版（V0）
154.8 ms
FlashRT 交付版（V4）
56.2 ms
说明：交付版的 56.2 ms 为 p10 56.0 / p90 56.3；首版首次测得 154.8 ms，绑定 commit 复测 155.3 ms，差异属测量噪声。
2.2 RTF 换算（实时率，越小越好）RTF = 单次推理时间(s) × 控制频率 / action horizon(40)。官方未公布控制频率，因此给出 10 / 20 / 30 Hz 三档参考：
阶段
@10 Hz
@20 Hz
@30 Hz
官方原实现 602 ms
0.151
0.301
0.452
FlashRT 首版 154.8 ms
0.0387
0.0774
0.1161
FlashRT 交付版 56.2 ms
0.0140
0.0281
0.0422
2.3 正确性与确定性以官方原实现在同一批观测上输出的 raw action chunk 为 ground truth，FlashRT 交付版 action_cosine = 0.996841，高于本次门槛 0.99。连续 5 次重放比对输出，5/5 位相等，确定性可复现。官方是纯 BF16，FlashRT 交付版用 FP8 / NVFP4 混合精度，因此 action 输出不是逐位相同，而是 0.996841 —— 这是刻意花掉的精度预算，不是误差累积。正确性必须在实际交付的那条执行路径上测：V1 曾因正确性未在交付路径上单独校验而重开修复。03 · STAGES五个阶段：154.8 → 111.6 → 79.0 → 71.6 → 56.2每一步都有对应的 commit 与当次的 action cosine，可以逐段复现、逐段对比：
阶段
commit
延迟
cosine
核心手段
V0
b6926433
154.8 ms
基准
sm120 可运行封装；热路径仍走框架 fallback
V1
b0b73a64
111.6 ms
0.999901
FP8 block-128 GEMM、融合激活、原生 attention、去框架算子
V2
ee4d6008
79.0 ms
0.999701
NVFP4 W4A4 权重量化、GEMM epilogue 融合、热路径框架算子 222 → 0
V3
e7143f63
71.6 ms
0.999755
ViT fc2 融合 NVFP4 producer
V4
19c2da00
56.2 ms
0.996841
expert NVFP4、大 M GEMM tile 路由、FA2 small-N、fc1 epilogue 融合、ViT hd72 去 padding、producer/norm 重写、host 流水
从 V1 到 V4，延迟从 111.6 ms 一路降到 56.2 ms，几乎每一档都是在「精度还留在门槛之上」的前提下换来的。
04 · V4 SEQUENCEV4 的 14 步：每一毫秒都有出处V4 净改进 −15.5 ms（−21.6%），由 14 次提交累积而成。 下表右侧两列是每一步之后的延迟与 cosine —— 可以清楚看到精度是怎么被花掉、又在哪一步补回来的：
#
优化内容
commit
延迟
cosine
1
expert qkv/gu 走 NVFP4 W4A4 + 融合 NVFP4 producer
26a8d98d
71.6 → 69.6
0.998109
2
全 NVFP4 expert（o/dn 走手写 cksplit small-M kernel）+ pinned host H2D staging
fc200fc3
69.6 → 69.0
0.995839
3
ViT fc2 GELU → NVFP4 producer v2（register-resident）
00cab2f4
69.0 → 66.8
0.995839
4
FA2 small-N query-tile 特化（hd128 denoise <64,64,4> + hd96 ViT <64,32,4>）
bc03172f
66.8 → 64.3
0.995123
5
residual RMSNorm → NVFP4 producer v3（register-cached）
1bdde864
64.3 → 63.5
0.995123
6
大 M 无 bias NVFP4 GEMM 走 <128,256,128> pingpong
7135efa8
63.5 → 62.6
0.995123
7
pingpong + bias 融合变体（ViT qkv/proj/fc2 大 M bias 形状）
4616f346
62.6 → 61.7
0.995123
8
prefix large-N GEMM tuning
dd4c849f
61.7 → 61.0
0.995123
9
prefix down-proj SwiGLU + NVFP4 producer 融合
1bea8068
61.0 → 60.8
0.995123
10
输入 cast/H2D 流水 + FA2 prepare_q 融合进 RoPE（去 320 次 launch）
321f0a8e
60.8 → 59.8
0.995123
11
ViT fc1 producer 融进 GEMM epilogue（bias + tanh-GELU + NVFP4 quant）
c4194a63
59.8 → 57.7
0.997593
12
ViT patch embed：cuDNN stride-16 conv → im2col bf16 matmul
e8709985
57.7 → 57.4
0.996841
13
ViT temporal-mix kernel bf16x2 向量化（位相等）
d6344956
57.4（噪声内）
0.996841
14
ViT head-dim 72 去 padding（见第 5 节）
19c2da00
57.4 → 56.1
0.996841
cosine 的变化轨迹本身就是一条线索：第 4 步降到 0.995123（FP4 路由的精度边界）， 第 11 步因为 epilogue 用 fp32 回升到 0.997593，第 12 步因 GEMM 累加顺序不同降到 0.996841， 第 14 步位相等故不变 —— 每一步都始终高于本次门槛 0.99。
05 · CASE一个值得单独说的优化：ViT head-dim 72 去 paddingViT 的头维本来是 72，但历史上被零填充到 96，只是为了启用 FA2 的 tile —— 多出 24/96 = 25% 的纯浪费，体现在 qkv GEMM 的 N（4608 vs 3456）、 proj-gather 的维度（96 vs 72）和 temporal-mix 的流量上。
转机来自一次更正：FA2 的 static_assert(kHeadDim % 32 == 0) 只约束编译期的 kHeadDim； 运行时把 head_dim=72 喂给现成的 <96,64,32,4> tile 时， is_even_K 会判为 false，kernel 只加载 72 列、其余补零。隔离实测：D=72 为 157.2 µs / cos 0.999998， 对比 D=96 的 155.4 µs，每次调用只差约 2 µs。
于是这是一个纯 Python / config 的改动，不需要重编任何 kernel： 去掉 qkv 权重与 bias 的 72→96 padding，attention 直接走现成 tile。 成对交替实测整机延迟 57.44 → 56.21 ms（−1.23 ms）， 且 action cosine 0.996841 与 padded 路径位相等（因为被 pad 的行列是精确的 0），5 次重放位相等。
顺带记一笔被否掉的路：曾想为 D=72 手写一个 wmma attention， 结果被 sm_120 消费级卡的动态 shared memory 上限卡住 （cudaDevAttrMaxSharedMemoryPerBlockOptin=101376，约 99 KB，而设计需要约 177 KB）。 最后还是发现现成 tile 通过 is_even_K 本来就支持 D=72 —— 不需要自己写。
06 · CENSUS56.2 毫秒花在哪里：内核热点交付版本上的内核调用统计（调用次数 / 累计耗时）：
内核
调用次数
累计耗时
cutlass_gemm（ViT / prefill / 大 M）
1054
23.25 ms
expert o/dn cksplit<32,4,2>
640
6.38 ms
Memcpy HtoD
5
3.99 ms
ViT FA2 flash_fwd<96,64,32,4>
27
3.67 ms
ViT add_layer_norm_to_nvfp4 producer
46
3.33 ms
ViT/prefill FA2 flash_fwd<128,64,64,4>
320
2.83 ms
residual_add_rms_norm_to_nvfp4（v3）
693
1.94 ms
silu_mul_merged_grouped32
384
1.31 ms
ViT proj_gather<72>
27
1.19 ms
rope_qknorm_kvwrite_qb
320
1.18 ms
prefill_attn
32
1.16 ms
gather_o_nvfp4
320
0.98 ms
temporal_mix_v2
6
0.62 ms
可以看到，交付路径上已经没有任何框架算子在热路径里（V2 阶段把 222 个降到 0）， 剩下的耗时几乎全是 GEMM、attention 与融合 producer —— 也就是说，进一步的收益要来自 更细的 tile 路由和更低比特的精度路径，而不是「去掉某个算子」。
07 · HONESTY已探明的地板，和试过但没走通的路性能报告里最容易省掉的部分，恰恰是最有价值的部分。这次我们把「试了但不用」的路线也一并记录：
项
结论
expert M=41 的小 M GEMM
CUTLASS persistent 方案已近最优；CUTLASS W4A4 在 M=41 反而比 FP8 慢约 13%，改用自研 cksplit 小 M kernel 才拿到收益
ViT small-N attention
对 <96,64,32,4> 做了完整 trait sweep（kBlockM / kBlockN / kNWarps），production tile 就是最优，没有更好的配置
自写 D=72 wmma attention
被 sm_120 消费级卡约 99 KB 的动态 smem 上限阻塞（设计需约 177 KB），最终证实现成 tile 已支持 D=72，无需自写
host-resize-before-H2D
因 parity 崩溃（cos 0.970968）否决
FP8 / NVFP4 路由、prefix 大 N 调优、NVFP4 fc1/fc2（V1 阶段）
测为 non-winning，或精度不达标（cos 0.998093 < 0.999）
08 · LIMITS使用时的几点说明单卡、单 shape、稳态。所有数字来自一张 RTX 5060 Ti、固定 shape 的单次决策，未覆盖动态 batching、多请求并发、其它 GPU 架构。精度不是逐位相同：官方纯 BF16，FlashRT 交付版为 FP8 / NVFP4 混合精度，action_cosine = 0.996841。且这是观测上的余弦相似度，不等于真实闭环任务成功率。计时噪声约 1–5%，GPU 可能被其它租户共享（曾观测到 5.9–8.4 GB 外部占用）。跨机器的绝对值会有波动，比例更稳健。复现前提：需要 FlashRT 仓库（含 HyVLA RoboTwin 后端）与官方 Hy-Embodied-0.5-VLA-RoboTwin 权重 / norm_stats；换 commit 或换环境，数字可能变化。报告中不包含任何密钥、账号或私有数据路径，脚本只读模型与本地产物。
09 · TRY IT怎么用起来：FlashHub 上可以直接获取这套适配已经整理进 FlashHub。你不需要自己从零复现上面那 14 步 —— 在 FlashHub 上可以获取使用方式、环境要求与测试入口，按页面说明跑通即可， 欢迎实测并反馈你机器上的数字。
在 FlashHub 获取使用方式查看 HyVLA 适配的模型信息、运行环境与测试方式，跑完把你的延迟与 cosine 贴回来。
访问 flashhub.top →
flashhub.top · 具身全场景解决方案平台
想自己复测？最小步骤准备 CUDA 13.0 + Python 3.12 + torch 2.14.0+cu130 + CUTLASS + FlashAttention-2 环境。跑官方基线脚本，确认得到约 600 ms 量级的完整决策延迟。按首版 / 交付版的 commit 构建 FlashRT（注意 -j1）。分别测首版与交付版（交付版加 --fp8-block128，并设置 PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True）。跑正确性脚本，确认 cosine ≥ 0.99；再做 5 次重放确认确定性。记录 GPU 型号、驱动、CUDA、torch 版本、commit、显存占用与 p10 / p90，便于横向比较。10 · YOUR MODEL希望 FlashHub 支持别的模型？留言告诉我们HyVLA 只是这次的适配对象。真正决定「下一个适配谁」的，是大家手上正在用、也真正卡在性能上的那些模型。 如果你有特定模型希望 FlashHub 支持，欢迎留言说明：
模型名称与来源（开源仓库 / 权重链接，或型号说明）；你要跑在什么设备上（例如这张 5060 Ti、Orin / Thor、国产加速卡等）；当前的延迟和你的目标（现在多少毫秒、希望压到多少）；使用场景（VLA 决策、视觉骨干、世界模型，还是别的推理链路）。留言方式：在 FlashHub 站内留言，或在本篇评论区直接说明—— 信息越具体，越容易判断这条路值不值得排进来。
站内入口： flashhub.top
同一个模型，同一张消费级显卡，602 ms → 56.2 ms
使用方式与测试入口都在 FlashHub 上，欢迎实测、欢迎拍砖、欢迎点名下一个模型。