Summary
pygraph.execute() costs ~45-50 µs of host time per call on the python frontend for an SDPA graph with 13 tensors (4 ragged offsets + 2 cu_seq_lens + Q/K/V/O/Stats + Stats ragged offset), of which only ~14 µs is cudnnBackendExecute (measured by interposing the symbol the frontend resolves with dlsym). The remaining ~30 µs is the binding: python dict → std::unordered_map<int64_t, void*>, the override_uids/override_shapes/override_strides lists when execute-time override is used, workspace handling, and the handle lookup.
For comparison FlashInfer's FA2 path launches its kernel through TVM-FFI in ~14 µs total, so on small serving steps (decode-context-parallel verify, chunk remainders) the frontend glue is the difference between the two backends.
Environment
cudnn-frontend 1.30 (release pybind), cuDNN 9.27, torch 2.13, SM100 box; medians over 2000 calls with the GPU idle.
Suggestion
A pre-bound execution object: build the variant pack once from {uid: tensor} (or accept raw pointers / a preallocated int64_t*+void** pair), and let execute take only the pointers that changed (plus override shapes as a flat buffer). Same spirit as the stream caching already done in cudnn.Handle (set_stream is now 0.09 µs when unchanged vs 1.5 µs raw).
note to self: claude::61d24ed2-7c90-4a97-9cbb-b91ae18136fd — "flashinfer frost prefill/GEMM enablement" · cwd /home/scratch.yanxu_libs/flashinfer
Summary
pygraph.execute()costs ~45-50 µs of host time per call on the python frontend for an SDPA graph with 13 tensors (4 ragged offsets + 2 cu_seq_lens + Q/K/V/O/Stats + Stats ragged offset), of which only ~14 µs iscudnnBackendExecute(measured by interposing the symbol the frontend resolves withdlsym). The remaining ~30 µs is the binding: python dict →std::unordered_map<int64_t, void*>, theoverride_uids/override_shapes/override_strideslists when execute-time override is used, workspace handling, and the handle lookup.For comparison FlashInfer's FA2 path launches its kernel through TVM-FFI in ~14 µs total, so on small serving steps (decode-context-parallel verify, chunk remainders) the frontend glue is the difference between the two backends.
Environment
cudnn-frontend 1.30 (release pybind), cuDNN 9.27, torch 2.13, SM100 box; medians over 2000 calls with the GPU idle.
Suggestion
A pre-bound execution object: build the variant pack once from
{uid: tensor}(or accept raw pointers / a preallocatedint64_t*+void**pair), and letexecutetake only the pointers that changed (plus override shapes as a flat buffer). Same spirit as the stream caching already done incudnn.Handle(set_streamis now 0.09 µs when unchanged vs 1.5 µs raw).note to self: claude::61d24ed2-7c90-4a97-9cbb-b91ae18136fd — "flashinfer frost prefill/GEMM enablement" · cwd /home/scratch.yanxu_libs/flashinfer