[GROOT N1.7][RTX 5090] Unable to reproduce the reported 16.61 ms E2E latency

[YellowRiver2001]

Body

Description

The documentation reports 16.61 ms warm P50 E2E latency for GROOT N1.7 FP8 on an RTX 5090 using a 2-view base checkpoint:

https://github.com/flashrt-project/FlashRT/blob/main/USAGE.md#groot-n17-rtx

The result appears to come from commit:

1dd5717

I have not been able to reproduce this number. More importantly, the repository does not appear to include the exact benchmark fixture, timing script, or complete environment used for the 16.61 ms
result.

Environment

• GPU: NVIDIA GeForce RTX 5090
• Compute capability: SM120
• Driver: 595.84
• CUDA toolkit: 12.8.61
• PyTorch: 2.11.0+cu128
• cuBLASLt: 12.9.2.10
• Build target: sm_120a
• Precision: FP8
• Batch size: 1
• Denoising steps: 4
• Action horizon: T=40

The GPU was operating at approximately 2.9 GHz with no observed power or thermal throttling.

Results

Using a real 2-view LIBERO input generated through the GROOT N1.7 preprocessing path:

• pixel_features: (512, 1024)
• grid_thw: (2, 3)
• llm_input_embeds: (1, 156, 2048)

Measured warm P50 latency:

Measurement P50
━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━ ━━━━━━━━━━
Backbone CUDA graph 9.77 ms
────────────────────────────── ──────────
Action CUDA graph 11.15 ms
────────────────────────────── ──────────
Two graph replays 20.92 ms
────────────────────────────── ──────────
Full infer(..., aux=aux) API 21.24 ms
────────────────────────────── ──────────
Documented result 16.61 ms

This leaves a gap of approximately 4.6 ms / 28% versus the documented result.

I also tested with cuBLASLt 12.8.3 instead of 12.9.2, but the result was effectively unchanged.

Missing reproduction details

Could you please provide the following?

1. Exact checkpoint repository and revision/hash.

2. The exact 2-view aux fixture, or a script for generating it.

3. Shapes of all tensors in the fixture, especially:

   • pixel_features
   • grid_thw
   • llm_input_embeds
   • visual_pos_masks
   • rope_cos / rope_sin

4. Exact prompt, normalized state, and initial noise used.

5. Complete benchmark script or command.

6. Number of warmup and timed iterations.

7. Whether latency was measured using CUDA events or wall-clock timing.

8. Whether the timed call was exactly:

   model.infer(
       state_normalized,
       initial_noise=initial_noise,
       aux=aux,
       use_dit_graph=True,
   )

9. CUDA, PyTorch, cuBLAS/cuBLASLt, driver, and compiler versions.

10. GPU power limit and clock configuration.

11. Expected backbone and action-graph latency breakdown.

Ideally, a self-contained script such as:

benchmarks/groot_n17_rtx_latency.py

together with a downloadable fixture or fixture-generation script would make the 16.61 ms result independently reproducible.

Thanks!

[LiangSu8899]

Hi! Thanks for the detailed environment and latency breakdown.

The CUDA 12.8+ requirement in the installation guide is the general minimum for the broader Blackwell family, including GeForce Blackwell and Thor-class targets. It does not guarantee that every model-specific FP8 cuBLASLt descriptor is supported by every CUDA 12.x release.

RTX 5090/SM120 is more specific. NVIDIA documents that tensor-scaled FP8 GEMMs on Blackwell GeForce, compute capability 12.x, require TN format—A transposed and B non-transposed:

• [cuBLAS 12.9 — Narrow Precision Data Types Usage](https://docs.nvidia.com/cuda/archive/12.9.1/cublas/index.html#narrow-precision-data-types-usage)
• [CUDA 12.8 — Blackwell cuBLAS support](https://docs.nvidia.com/cuda/archive/12.8.0/cuda-toolkit-release-notes/index.html#cublas-library)

and also documents this class of failure in the installation guide:

• Known runtime issue: cuBLASLt FP8 heuristic code=15

As noted there, the CUDA major version alone is not sufficient to determine FP8 support. The exact cuBLASLt runtime patch version, GPU architecture, GEMM shape/layout, scaling mode, and epilogue must be checked.

At commit 1dd5717, GROOT N1.7 calls fp8_descale_fp16, whose cuBLASLt descriptor uses CUBLAS_OP_N for both operands. The extension can be compiled successfully for sm_120a with CUDA 12.8.61, but the first FP8 ViT GEMM may fail at runtime:

M=512, N=1024, K=1024
CUBLAS_STATUS_NOT_SUPPORTED (code=15)

This behavior was observed with both cuBLASLt 12.8 and 12.9. It is a cuBLASLt heuristic/runtime limitation, not a failure to compile FP8 or sm_120a kernels. The FP16 path works in the same CUDA 12.8 environment.

I also compared 1dd5717 with the later GROOT control-loop changes. The relevant implementations in gemm_runner.cu and pipeline_rtx_fp8.py did not change, and no automatic FP16 or CUTLASS fallback was added.

This leaves three likely explanations:

1. The extension was actually built against CUDA 13

   PyTorch may report 2.11.0+cu128, while flash_rt_kernels.so links against libcublasLt.so.13. PyTorch’s CUDA tag does not determine the CUDA runtime used by an independently compiled extension.

2. The measured frontend was the full-FP16 frontend

   The CUDA 12.8 FP16 measurements for the reported shapes were:

   FP16 backbone:       9.21 ms
   FP16 action graph:  10.41 ms
   Combined:           19.66 ms
   

   These are close to the reported values:

   Backbone graph:      9.77 ms
   Action graph:       11.15 ms
   Combined:           20.92 ms
   

3. Driver 595.84 changes the algorithms exposed by cuBLASLt 12.9

   This cannot yet be excluded. However, the first two possibilities are more consistent with NVIDIA’s documented FP8 restrictions for Blackwell GeForce under cuBLAS 12.x.

Could you run the following in the exact environment used for the 21.24 ms result?

python - <<'PY'
import torch
import flash_rt.flash_rt_kernels as kernels

print("torch:", torch.__version__)
print("torch CUDA:", torch.version.cuda)
print("GPU:", torch.cuda.get_device_name(0))
print("extension:", kernels.__file__)
print("frontend:", type(model._pipe))
PY

Then inspect the extension path printed above:

EXTENSION=/exact/path/to/flash_rt_kernels.so

readelf -d "$EXTENSION" | grep NEEDED
ldd "$EXTENSION" | grep -E 'cublas|cudart|libcuda'
cuobjdump -lelf "$EXTENSION" | grep -E 'sm_120|sm_120a' | head
sha256sum "$EXTENSION"

The expected frontend for this FP8 path is:

GrootN17TorchFrontendRtxFP8

Please also load the exact libcublasLt path reported by ldd:

import ctypes

path = "/absolute/path/from/ldd/libcublasLt.so"
lib = ctypes.CDLL(path)
lib.cublasLtGetVersion.restype = ctypes.c_size_t
print("actual cuBLASLt version:", lib.cublasLtGetVersion())

For comparable steady-state timing, use one graph-capture call, five warmup calls, and 30 CUDA-event measurements:

model.infer(
    state_normalized,
    initial_noise=initial_noise,
    aux=aux,
    use_dit_graph=True,
)

for _ in range(5):
    model.infer(
        state_normalized,
        initial_noise=initial_noise,
        aux=aux,
        use_dit_graph=True,
    )

torch.cuda.synchronize()

times = []
for _ in range(30):
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    start.record()
    model.infer(
        state_normalized,
        initial_noise=initial_noise,
        aux=aux,
        use_dit_graph=True,
    )
    end.record()
    end.synchronize()
    times.append(start.elapsed_time(end))

times.sort()
print("P50:", times[len(times) // 2], "ms")
print("mean:", sum(times) / len(times), "ms")
print("range:", times[0], times[-1], "ms")

The currently validated RTX 5090 FP8 configuration is:

CUDA Toolkit: 13.0.88
cuBLASLt:     13.1.0.3
Build target: sm_120a
Container:    NVIDIA NGC PyTorch 25.10

This configuration successfully executes the FP8 descriptor and reproduces the documented two-view result at approximately 16.58 ms.