Staging by Anerudhan · Pull Request #2 · NVIDIA/cudnn-frontend

Anerudhan · 2021-03-20T03:30:09Z

Changes in pull request:

Fix compilation warnings reported with -Wall and -Wextra flags
Support for backward activations dx = f(dy, X).
Support for lower_clip, upper_clip, lower_clip_slope and alpha and beta paramters for relu, elu, softplus and swish.
Added additional checks during build phase. Such as for bDesc being nullptr etc.
Improved error checking for xDesc, yDesc depending on whether the operation is convolution or pointwise.
Add matmul descriptor
Add conv_scale_bias_add_relu and matmul_bias_gelu sample
Comparison between frontend and backend
Fix compilation issue in samples for gcc-5
New sample for HEUR_B

YangXu1990uiuc

Changes lgtm thank you!

Completes the remove/relocate edits from the code review that the inline suggestion blocks could not express on their own: - NVIDIA#2: remove the now-duplicate same_hdim t2r_dKV4 wait before dKV1 - NVIDIA#4: remove the now-duplicate same_hdim t2r_dKV01 wait before dKV3 - NVIDIA#6: switch same_hdim dKV2/dKV3 drain to split T2R + reduce_dKV_from_reg Also repairs line-drift from batch-applying the earlier suggestions (their anchors shifted after the first suggestion inserted lines): - NVIDIA#3: remove a leftover dangling `if cutlass.const_expr(not self.same_hdim_kv):` with no body (IndentationError) left because the original lines were not replaced - NVIDIA#5: remove the leftover old dKV0/dKV1 if/else block that sat before the new unified block, which would otherwise double the reduce/atomic_add and the t2r_dKV01 barrier arrival Final state is the intended split-T2R + signal drain refactor: every reduce site loads dKV into registers, fences, signals the MMA warp that the TMEM columns are free, then does the global atomic_add; same_hdim and not-same_hdim now follow the same protocol. dq is bitwise-identical to ef50ddc; head_dim 576 and 512 both verified race-free (compute-sanitizer racecheck 0).

Completes the remove/relocate edits from the code review that the inline suggestion blocks could not express on their own: - NVIDIA#2: remove the now-duplicate same_hdim t2r_dKV4 wait before dKV1 - NVIDIA#4: remove the now-duplicate same_hdim t2r_dKV01 wait before dKV3 - NVIDIA#6: switch same_hdim dKV2/dKV3 drain to split T2R + reduce_dKV_from_reg Also repairs line-drift from batch-applying the earlier suggestions (their anchors shifted after the first suggestion inserted lines): - NVIDIA#3: remove a leftover dangling `if cutlass.const_expr(not self.same_hdim_kv):` with no body (IndentationError) left because the original lines were not replaced - NVIDIA#5: remove the leftover old dKV0/dKV1 if/else block that sat before the new unified block, which would otherwise double the reduce/atomic_add and the t2r_dKV01 barrier arrival Final state is the intended split-T2R + signal drain refactor: every reduce site loads dKV into registers, fences, signals the MMA warp that the TMEM columns are free, then does the global atomic_add; same_hdim and not-same_hdim now follow the same protocol. dq is bitwise-identical to ef50ddc; head_dim 576 and 512 both verified race-free (compute-sanitizer racecheck 0). Signed-off-by: zky <51477259+zkyue@users.noreply.github.com>

…m 576) (#396) * Fix latent TMEM WAR race in DSA backward SM100 dKV drain (head_dim 576 path) In the not-same_hdim_kv path (head_dim 576 / head_dim_v 512), the reduce warps read dKV2/dKV3 from TMEM inside store_dKV, after their t2r_dKV4_done arrive. The MMA warp's next-iteration dKV0/dKV1 gemms overwrite the same TMEM columns (tmem_dKV2_offset == tmem_dKV0_offset, tmem_dKV3_offset == tmem_dKV1_offset) with no happens-before edge to those reads: mma_reduce_dKV_pipeline has 2 stages, so the part1 producer_acquire of generation 3i+3 only orders against the consumer_release of generation 3i+1 (the dKV4 generation), and no named barrier separates part2's TMEM reads (generation 3i+2) from the issue of the next part1's overwrites. Correctness currently depends on the tensor core still draining queued work when the reduce warps issue their loads. The window reproduces deterministically under adversarial delay: a 1 us spin inserted in the reduce warps between the two part2 store_dKV calls (no other change) corrupts dkv columns [384:512) on every run at S=2048 H=64 D=576 topk=512 bf16 (rel L2 vs fp32 autograd 1.0e0 vs 5.6e-3 baseline); a 10 us spin before both calls corrupts [256:384) as well. Delays at already-synchronized points (after part1's register-staged T2Rs, or before the MMA warp's dKV0 issue) never corrupt. Fix, mirroring part1's existing register-staging pattern: T2R dKV2/dKV3 into registers, fence, arrive on a new t2r_dKV23_done named barrier (id 9), and only then run the global-memory atomic reduction; the MMA warp waits on that barrier before issuing the next iteration's dKV0 (skipped on the first iteration, balanced by a final arrive after the loop, like t2r_dKV4_done). Register peak is unchanged (part1 already holds two fragments). The same_hdim_kv (head_dim 512) path is untouched and compiles identically. Measured on B200 (S=8192 H=64 topk=1024 bf16, ABAB paired): head_dim 576: 5.398 -> 5.840 ms median (+8.0%), the cost of no longer overlapping the next iteration's dKV issue with the dKV2/dKV3 readback; head_dim 512: +0.03% (noise). dq stays bitwise-identical to develop on both head dims; dkv/d_sink match the fp32 autograd reference at the baseline rel L2 on both, and the adversarial-delay matrix that trips develop is clean on the fixed kernel. Signed-off-by: zky <51477259+zkyue@users.noreply.github.com> * Apply suggestions from code review Co-authored-by: Jie Fang <jief@nvidia.com> Signed-off-by: zky <51477259+zkyue@users.noreply.github.com> * Complete review remove/relocate edits and repair suggestion line-drift Completes the remove/relocate edits from the code review that the inline suggestion blocks could not express on their own: - #2: remove the now-duplicate same_hdim t2r_dKV4 wait before dKV1 - #4: remove the now-duplicate same_hdim t2r_dKV01 wait before dKV3 - #6: switch same_hdim dKV2/dKV3 drain to split T2R + reduce_dKV_from_reg Also repairs line-drift from batch-applying the earlier suggestions (their anchors shifted after the first suggestion inserted lines): - #3: remove a leftover dangling `if cutlass.const_expr(not self.same_hdim_kv):` with no body (IndentationError) left because the original lines were not replaced - #5: remove the leftover old dKV0/dKV1 if/else block that sat before the new unified block, which would otherwise double the reduce/atomic_add and the t2r_dKV01 barrier arrival Final state is the intended split-T2R + signal drain refactor: every reduce site loads dKV into registers, fences, signals the MMA warp that the TMEM columns are free, then does the global atomic_add; same_hdim and not-same_hdim now follow the same protocol. dq is bitwise-identical to ef50ddc; head_dim 576 and 512 both verified race-free (compute-sanitizer racecheck 0). Signed-off-by: zky <51477259+zkyue@users.noreply.github.com> --------- Signed-off-by: zky <51477259+zkyue@users.noreply.github.com> Co-authored-by: Jie Fang <jief@nvidia.com>

Four env-bisectable levers on the v12 base (all-off reproduces v12); composition targets the two hardware-adjudicated ring legs from the vm5probe run2 evidence (MAT_ACQ:MAT_WAIT 70:30 back-pressure; leader per-pass software ~3us/tile is the largest single item; exchange seam ~1.5-1.7us; L2 read hit 95.4% killed the re-warm direction). - DSA_V15_L2X (default 1): P/dS exchange re-staged through an HBM ring (workspace_pds, [slot 2][cta 2][kind 2] x 8KB per token; hot set ~4.9MB at 74 resident clusters). Math publishes whole P/dS images in two waves through a new 8,192B pds_stage funded byte-exactly by deleting p_xchg + the write-dead ds_xchg; W18 becomes the pds producer (acquire moved from math -- the 2.2us MATH_PDS_ACQ park dies) and converts the ring into SMEM operands with five 1D bulk G2S fills, committing the leader's dQ gate only after ds_image bytes LAND (strictly stronger than v12). DSM sends, the count-128 pds_ready handshake and math's ring residence all retire. Leader code unchanged. workspace_pds is a trailing default-None param: pre-v15 harness call sites keep working; L2X=1 without the ring raises at build time (see V15_RUNNER_NOTES.md for the 2-line harness patch). - DSA_V15_REGSWAP (default 2): warpgroup-uniform variants only (setmaxnreg is warpgroup-granular): B = W16-19 uniform 64 / reduce 120; C = W16-19 uniform 56 / gather 40. Pool exactly 61,440 in all settings. - DSA_V15_DQ_MERGE (default 1): both kdq credits waited back-to-back (ROUTE_K acquires AND commits them upfront -- no producer cycle, unlike the condemned dVdK merge), releases staggered per round. - DSA_V15_ALLTMA (default 1): own-half DSM bulk retired via the pre-existing pure-TMA fallback (run2: DSM legs 1.7x slower mean, 2.25x worse p90; also retires the spec-UB own-CTA DSM bulk). Adversarially reviewed (protocol + DSL refuters + merge, zero fatals): byte identity of all three operand buffers re-derived from layouts; wait_group(0) write-completion ordering for the cross-proxy S2G->G2S edge confirmed; review fix NVIDIA#1 applied (gen_ready mbars slot ping-pong -- restores the structural <=1-ahead publish bound v12 had); review fix NVIDIA#2 applied (batch>1 ring aliasing guard). IKET at L2X=1: 28 names (retired MATH_PDS_ACQ/MATH_BAR1/ROUTE_P/ROUTE_dS, added W18_PDS); L2X=0 bisect builds are release/correctness-only (31 names > 29 cap). Expected: realistic 8.5-8.9ms, optimistic 7.7-8.1 (v12 = 11.654, baseline = 8.14). Runner protocol: V15_RUNNER_NOTES.md (v15_run1). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

dsa_bwd_sm100_2cta_v32.py: the fine-grained all-f32 transposed form (h_c=64 / D_c=128 / dual dKV slot) per T3细粒度全f32转置设计_20260803 + T32_BUILD_ADDENDUM (dS dual-image dq_b repair, third-audit errata). Five transposed GEMMs, kv128-bundle chase gather, 12-gen natural- layout supply ring, mb_dqb cluster gates (errata NVIDIA#2 arming), dual-slot drains x8/bundle, dr-major leader schedule, v17a-orientation dQ epilogue. Sections built incrementally with per-section py_compile; self-audit trailer + 11 V32-TODO audit pins in-file. IKET 25<=29. Untested; validation begins this commit. HANDOFF_20260803.md carries the full two-track session record. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

The same gap on the other side of the boundary. cuDNN's runtime-compiled engines launch through the driver, which reads the CALLING thread's context stack, and a thread that has done no CUDA work has nothing on it. Measured on a matmul+relu+relu graph (which routes to those engines), on a thread where cuDNN is the first CUDA call: before: cold #0/#1/NVIDIA#2 ctx 0x0 -> 0x0 cuCtxGetLimit returned error invalid device context (201) after: cold #0/#1/NVIDIA#2 ctx 0x0 -> 0x31275a20 OK Deterministic both ways, and one torch op on the thread beforehand hides it entirely -- the CUDA runtime binds the primary context as a side effect, and something normally does, which is why no framework has run into this. The precompiled engines launch with <<<>>> and are unaffected for the same reason. Placed at execute_plan_at_index, which the file already documents as the point all execute overloads funnel through, so backend and OSS paths are both covered once. The driver entry points are resolved through the runtime (cudaGetDriverEntryPointByVersion), so the front end still never links libcuda -- the approach cu_tensor_map_encode_tiled already uses and documents. NV_FE_CALL_TO_CU exists but is unused and would link libcuda directly in the non-dynamic-loading build. Same rung order as the Python side: a bound context is left alone (it is process-wide and the caller chose it), a real stream names its own context, and the default-stream handles (0, CU_STREAM_LEGACY, CU_STREAM_PER_THREAD) name none, so the runtime's device decides there.

The same gap on the other side of the boundary. cuDNN's runtime-compiled engines launch through the driver, which reads the CALLING thread's context stack, and a thread that has done no CUDA work has nothing on it. Measured on a matmul+relu+relu graph (which routes to those engines), on a thread where cuDNN is the first CUDA call: before: cold #0/#1/NVIDIA#2 ctx 0x0 -> 0x0 cuCtxGetLimit returned error invalid device context (201) after: cold #0/#1/NVIDIA#2 ctx 0x0 -> 0x3f67f670 OK Deterministic both ways, and one torch op on the thread beforehand hides it entirely -- the CUDA runtime binds the primary context as a side effect, and something normally does, which is why no framework has run into this. The precompiled engines launch with <<<>>> and are unaffected for the same reason. Placed at execute_plan_at_index, which the file already documents as the point all execute overloads funnel through, so backend and OSS paths are both covered once. Driver entry points are resolved through the runtime (cudaGetDriverEntryPointByVersion), so the front end still never links libcuda -- the approach cu_tensor_map_encode_tiled already uses and documents. Same rung order as the Python side: a bound context is left alone, a real stream names its own context, and the default-stream handles name none, so the runtime's device decides there. Review: the dynamic-loading lookup is non-throwing (get_cuda_symbol throws when the library or symbol is missing, and this runs in a static initializer), guarded the way the rest of the headers guard exceptions; the backend test skips instead of failing where no engine serves the fused graph. Comments trimmed throughout -- the rationale and the measurements live in the PR.

…n both sides of the boundary (#626) * fix(device): ensure the RIGHT context, not merely a context ensure_current_context returned as soon as ANY context was current, so a thread already bound to another GPU's context kept it. The legacy default stream (handle 0) carries no context of its own -- it resolves against whatever is current -- so under a foreign context the work runs on THAT context's GPU, where the pointers are invalid: an async fault at some later sync rather than an error at the launch. A real stream does carry its context and a cross-context launch is rejected outright, so only the stream-0 path is silent, and stream 0 is exactly what torch's default stream is. Resolve the target rather than accept the incumbent: the stream's context when the stream names one, else the caller's device. execute() passes the handle's ordinal, so the FE path no longer asks the runtime which GPU it is on -- Handle.device owns that since #612 -- and cudaGetDevice() stays only as the fallback for a caller that cannot name a device. Cost measured on parley: 59 ns per execute in situ, 0.1% of a 79 us GDN op. test_ensure_current_context.py covers the cold thread, the foreign-device replacement, stream-wins-over-device, and the steady-state no-op. * fix(execute): bind a context on the calling thread in C++ too The same gap on the other side of the boundary. cuDNN's runtime-compiled engines launch through the driver, which reads the CALLING thread's context stack, and a thread that has done no CUDA work has nothing on it. Measured on a matmul+relu+relu graph (which routes to those engines), on a thread where cuDNN is the first CUDA call: before: cold #0/#1/#2 ctx 0x0 -> 0x0 cuCtxGetLimit returned error invalid device context (201) after: cold #0/#1/#2 ctx 0x0 -> 0x3f67f670 OK Deterministic both ways, and one torch op on the thread beforehand hides it entirely -- the CUDA runtime binds the primary context as a side effect, and something normally does, which is why no framework has run into this. The precompiled engines launch with <<<>>> and are unaffected for the same reason. Placed at execute_plan_at_index, which the file already documents as the point all execute overloads funnel through, so backend and OSS paths are both covered once. Driver entry points are resolved through the runtime (cudaGetDriverEntryPointByVersion), so the front end still never links libcuda -- the approach cu_tensor_map_encode_tiled already uses and documents. Same rung order as the Python side: a bound context is left alone, a real stream names its own context, and the default-stream handles name none, so the runtime's device decides there. Review: the dynamic-loading lookup is non-throwing (get_cuda_symbol throws when the library or symbol is missing, and this runs in a static initializer), guarded the way the rest of the headers guard exceptions; the backend test skips instead of failing where no engine serves the fused graph. Comments trimmed throughout -- the rationale and the measurements live in the PR. * build: declare the cuTile runtime as a [cutile] extra The cuTile linear-attention engines import cuda.tile, which was not declared anywhere -- not an extra, not requirements.txt -- so whether they run at all depended on the environment happening to have it. That is also why test_execute_from_a_thread_with_no_cuda_context silently covers only the FROST half on most machines: the cuTile engines decline in check_support when the import fails, and the test skips. Base cuda-tile only. Its [tileiras] extra pins cuda-toolkit>=13.2,<13.4, and that upper bound would cap the whole environment's toolkit and shut out CUDA 12 entirely -- the same reason nvidia-cutlass-dsl is not pinned to the FROST floor here. Without it cuda.tile falls back to a system tileiras, consistent with this package already leaving GPU wheels to the user. Resolution checked: `.[cutedsl,cutile]` resolves in one pass and adds exactly one package (cuda-tile 1.5.0) with nothing downgraded -- base cuda-tile requires only typing-extensions. The python_version marker keeps the extra resolvable on the declared 3.9 floor, which cuda-tile itself does not support. * perf(execute): probe the context before querying the stream The guard ran cudnnGetStream on every execute, though only the cold path needs a stream. Probe with one cuCtxGetCurrent instead and fetch the stream only when a context actually has to be established. Backend graph.execute() host time on parley, rebuilding the module for each: develop, no guard 10.805 us this PR, probe first 10.498 us / 10.683 us (two builds) The PR measured faster than develop both times, so the difference between builds is noise -- run-to-run spread alone is ~0.35 us across the 15 samples, and the guard's one cuCtxGetCurrent is ~106 ns. The unconditional version measured 10.788 us, i.e. also within noise: cudnnGetStream from C++ is nowhere near the ~1.5 us the Python path costs through pybind. Probing first is still the right shape, but it was not buying back a visible regression.

Anerudhan added 2 commits March 19, 2021 20:01

Copy over files from git-lab to external git.

dbd4eb0

Update documentation.

9f26fa4

Anerudhan requested a review from YangXu1990uiuc March 20, 2021 03:30

Anerudhan self-assigned this Mar 20, 2021

Anerudhan added 2 commits March 26, 2021 14:42

Changes from code review and new samples

4e0d2e4

Updating the documentation

0e084b9

github-pages Bot temporarily deployed to github-pages March 26, 2021 21:50 Inactive

YangXu1990uiuc approved these changes Mar 29, 2021

View reviewed changes

Anerudhan merged commit b4e1ad9 into main Mar 29, 2021

Anerudhan deleted the staging branch March 29, 2021 18:29

YJMSTR added a commit to YJMSTR/cudnn-frontend that referenced this pull request Jul 3, 2026

add compat for nvtx.range and ExternalStream (NVIDIA#2)

9edf6ef

lesaff-r mentioned this pull request Jul 22, 2026

[Bug]: Segmentation fault in libcudnn_engines_precompiled.so.9.4.0 when executing fused INT8 convolution with vectCount=4 #422

Open

1 task

YangXu1990uiuc mentioned this pull request Aug 17, 2026

frost(sdpa): add more features to sm120 frost fp8 sdpa_fwd kernel #595

Merged

3 tasks

Navigation Menu

Uh oh!

Staging - #2

Staging#2
Anerudhan merged 4 commits into
mainfrom
staging

Anerudhan commented Mar 20, 2021

Uh oh!

YangXu1990uiuc left a comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

Anerudhan commented Mar 20, 2021

Uh oh!

YangXu1990uiuc left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants