Add jax-cudnn-frontend agent skill - #2273
Open
jessicadesilva wants to merge 4 commits into
Open
Conversation
…els in pure JAX workflows
Steboss
reviewed
Aug 19, 2026
| the kernel once with fixed seeds, prints **value-based** evidence (NaN | ||
| counts, checksums, fixed-position samples), and exposes every configuration | ||
| choice as an env-var toggle. This script is simultaneously your repro for | ||
| bug reports and your bisection harness. Allocate outputs with `jnp.empty` so |
Contributor
There was a problem hiding this comment.
A little correction, from JAX 0.11 jnp.empty may not be reliable as a NaN check - they say it may retrn genuinely uninitialized memory. Thus, an unwritten tile can contain any finite value and evade the jnp.isnan.
We could say to use jnp.full(..., jnp.nan), wdyt? could it be useful for this case?
Steboss
reviewed
Aug 19, 2026
Steboss
left a comment
Steboss
left a comment
Contributor
There was a problem hiding this comment.
I love this job, a few comments but overall ready for approval
| ## Keep the direct path alive | ||
|
|
||
| Alongside the JAX-native wrapper, maintain a `cute.compile`-style direct | ||
| invocation of the same kernel instance (concrete arrays via the DSL's |
Contributor
There was a problem hiding this comment.
I am not an expert on dlpack, but if I see from JAX doc they say that external mutation of a DLPack-backed JAX array may lead to undefined behaviour. Could this be a problem?
Comment on lines
+119
to
+137
| ```python | ||
| @cute.jit | ||
| def _launcher(stream, x, meta, out, aux, *, scale: float): | ||
| _zero_f32(aux).launch( # prologue: aux is accumulated into | ||
| grid=(cute.ceil_div(cute.size(aux), 256), 1, 1), | ||
| block=(256, 1, 1), stream=stream) | ||
| _kernel(x, out, aux, meta, cutlass.Float32(scale), stream) | ||
|
|
||
| @functools.partial(jax.jit, static_argnums=()) | ||
| def op(x, meta): | ||
| return cjax.cutlass_call( | ||
| _launcher, | ||
| output_shape_dtype=[ | ||
| jax.ShapeDtypeStruct(x.shape, x.dtype), # out | ||
| jax.ShapeDtypeStruct(aux_shape, jnp.float32), # aux — shape from | ||
| ], # vendor allocation! | ||
| softmax_scale=..., | ||
| )(x, meta) | ||
| ``` |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Adds a discovery-first procedural skill for wrapping, integrating, and validating cuDNN Frontend / CuTe DSL kernels from pure JAX, without PyTorch.
Example usage:
Using the
jax-cudnn-frontendskill: wrap the cuDNN Frontend CuTe DSL block sparse attention kernels for inference and training from pure JAX, exposing the kernel configuration knobs so we can autotune per problem shape.Lives at:
docs/agent-skills/jax-cudnn-frontend/so it's discoverable by any AI coding agent or human.Covers: environment discovery, locating the vendor's own kernel call site as the source of truth for tensor contracts, minimal standalone repro scripts, JAX wrapper validation, and a symptom-indexed debugging guide distilled from real integration failures.