natten: port NATTEN v0.21.6 to kernel-builder by sayakpaul · Pull Request #1031 · huggingface/kernels-community

sayakpaul · 2026-07-16T10:42:19Z

Port https://natten.org/ for an upcoming cool release

Vendors SHI-Labs/NATTEN v0.21.6 (MIT) as a kernel-builder edition-5 CUDA kernel targeting kernels-community/natten: - csrc/ vendored with committed autogen output (144 TUs, upstream's 'default' policy); scripts/regen.sh + generate_build_toml.py make upstream bumps a one-command regen. - pybind11 bindings replaced with TORCH_LIBRARY registrations (torch-ext/torch_binding.cpp, 37 out-variant ops, limited-API safe); torch/extension.h stripped from static and autogen sources. - Full natten Python package (minus profiler) vendored with relative imports; _libnatten rewritten on top of the generated _ops module with fake (meta) registrations for torch.compile. - Kernel sections split by arch: core (reference + CUTLASS 2.X FNA/FMHA, SM70-SM120), hopper (9.0a), blackwell (10.0a, cuda-minver 12.8 so cu126 variants still build; natten.h derives the backend gates from CUDART_VERSION so static dispatchers stub out consistently). - Upstream test suite vendored + fast kernels_ci parity subset. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

CI failed on all variants with a fixed-output hash mismatch fetching the CUTLASS 4.5.2 source: nix-builder's recorded hash predates whatever the v4.5.2 tag resolves to now, and natten is the first cutlass_4_5 consumer. Pin the kernel-builder input to huggingface/kernels@ce92bd77 (branch fix-cutlass-4-5-hash), which corrects the hash; re-lock to main once the upstream PR merges. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

Upstream NATTEN's helpers.h and the autogen templates use torch::kFloat16 & co., which come from the torch C++ frontend headers that upstream pulled in transitively via torch/extension.h. With that include stripped for the limited API, the autogen TUs only see ATen and fail to compile (https://github.com/huggingface/kernels-community/actions/runs/29492671685). Switch to the identical at:: constants (kFloat32->kFloat, kFloat16->kHalf, kBFloat16->kBFloat16, zeros->zeros) in helpers.h and the autogen scripts, and regenerate. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

kernel-builder installs the package contents flat into the build variant directory, and test runners put that directory on PYTHONPATH. natten's top-level types module therefore shadowed the standard library's types module, and Python failed during site initialization before any test ran (https://github.com/huggingface/kernels-community/actions/runs/29493798395). Rename the module to _types (the only stdlib-colliding name in the package), rewrite the 28 internal import sites, and alias it back as natten.types in __init__ so the upstream import surface (from natten.types import ...) is unchanged, including for the vendored tests. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

These files are generated per-build by kernel-builder create-pyproject and were committed by mistake in the previous commit. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

kernel-builder's compat shim executes the package under a path-derived module name and only copies globals into `natten`, so real submodule imports (from natten.types / natten.utils.testing / ... import X) raised ModuleNotFoundError in the CI test shell. Install a meta-path finder (only when `natten` in sys.modules is our own compat shim) that resolves natten.* to the already-loaded hash-named modules, preserving module identity and avoiding double execution via parent __path__ lookups. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

sayakpaul · 2026-07-17T15:42:34Z

/kernel-bot build-and-stage natten --branch v0

github-actions · 2026-07-17T15:42:54Z

Build request processed.

Command: /kernel-bot build-and-stage natten --branch v0
Mode: build and stage
Target branch: v0
PR head SHA: f56f26d3018b4d000f40b04551e0ab9b7586988f
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (2):

natten (build-mac.yaml): https://github.com/huggingface/kernels-community/actions/runs/29593249107
natten (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/29593250728

alihassanijr · 2026-07-17T15:52:05Z

+    "9.0",
+    "10.0",


I noticed SM90 and SM100 have separate builds that compile with their arch-specific features (sm90a and sm100a). What's the reason for this separation?

The sections are how kernel-builder scopes -gencode flags: cuda-capabilities applies per [kernel.*]
section, so the hopper/blackwell sources get compiled only as sm_90a / sm_100a while the
legacy+FNA/FMHA core builds for the full 7.0–12.0 range — same effect as the per-source arch flags in
upstream's setup.py, just expressed in kernel-builder's schema.

The blackwell section additionally carries cuda-minver = "12.8", so it's skipped when building the CUDA 12.6 variants (where the SM100 CUTLASS kernels can't compile) without dropping the rest of the kernel. Each section also passes its NATTEN_WITH_*_FNA define to match the #ifdef gates in the autogen TUs; the core dispatchers derive the same defines from CUDART_VERSION so dispatch stays consistent with what was built.

alihassanijr · 2026-07-17T15:53:44Z

+cuda-capabilities = [
+    "10.0a",


Suggested change

cuda-capabilities = [

"10.0a",

cuda-capabilities = [

"10.0a",

"10.3a",

I'd make this change to support Blackwell Ultra ([G]B300) -- but only if the wheel is being built with CTK 12.9 or later. Sm103 was added in 12.9 release.

NATTEN reference for arches supported based on CTK version:
https://github.com/SHI-Labs/NATTEN/blob/main/scripts/packaging/pkg_helpers.bash#L43

Here's the blocker for that: huggingface/kernels#720

CUTLASS 4.5 turned cute::SM100_MMA_F8F6F4_SS from a tag type (paired with a variadic MMA_Traits specialization) into a class template, same shape as SM100_MMA_F16BF16_SS. The to_tiled_mma_sm100_ts overloads in the fna/fmha Blackwell collectives still pattern-matched the 4.4 traits form, failing on every blackwell TU (first exercised by the cu128+ release builds; PR builds only cover cu126, where the blackwell section is skipped via cuda-minver). Match the templated atom directly, exactly like the adjacent F16BF16 overload, which mirrors what the CUTLASS 4.5 collective builder emits via make_tiled_mma. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

sayakpaul · 2026-07-18T02:48:04Z

/kernel-bot build-and-stage natten --branch v0

github-actions · 2026-07-18T02:48:20Z

Build request processed.

Command: /kernel-bot build-and-stage natten --branch v0
Mode: build and stage
Target branch: v0
PR head SHA: 73e4fe133875789587406b9ee981842280fcd4cc
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (2):

natten (build-mac.yaml): https://github.com/huggingface/kernels-community/actions/runs/29627685072
natten (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/29627685759

sayakpaul · 2026-07-19T15:13:42Z

/kernel-bot build-and-stage natten --branch v0

github-actions · 2026-07-19T15:13:54Z

Build request processed.

Command: /kernel-bot build-and-stage natten --branch v0
Mode: build and stage
Target branch: v0
PR head SHA: 73e4fe133875789587406b9ee981842280fcd4cc
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (2):

natten (build-mac.yaml): https://github.com/huggingface/kernels-community/actions/runs/29692416453
natten (build.yaml): https://github.com/huggingface/kernels-community/actions/runs/29692416939

sayakpaul · 2026-07-20T02:37:57Z

@alihassanijr would it be possible for you to help us perform the "building and publishing" step? Our CI is running out: https://github.com/huggingface/kernels-community/actions/runs/29692416939 :D

Otherwise, I think we could do something like so (to keep moving things forward):

Reduce the core natten section's arch list (natten/build.toml currently has 8 entries: 7.0, 7.5, 8.0, 8.6, 8.9, 9.0, 10.0, 12.0 — every one is a -gencode pass on each of the 144 autogen TUs). flash-attn2 in this repo ships just [8.0, 9.0, 10.0, 12.0], relying on within-major SASS forward compatibility (sm_80 code runs on 8.6/8.9). I think dropping 8.6+8.9 alone cuts ~25% of the core section's nvcc work; also dropping 7.0/7.5 (no V100/T4) further reduces it.

Caveat: sm_86/89 have less shared memory than sm_80, so this is safe only if NATTEN's runtime config selection respects the device's actual smem limit.

Not sure if this applies to natten but we recently did #1032 which helped us cut down the build time of FA3 significantly.

github-actions Bot added the chore Version bumps, releases, misc maintenance label Jul 16, 2026

natten: add to KERNEL_SOURCE_MAPPING

755ffee

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

sayakpaul marked this pull request as draft July 16, 2026 10:45

sayakpaul and others added 5 commits July 16, 2026 16:27

natten: remove accidentally committed create-pyproject artifacts

38c1ffe

These files are generated per-build by kernel-builder create-pyproject and were committed by mistake in the previous commit. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

sayakpaul mentioned this pull request Jul 17, 2026

Improve API-change check: gate on version bump and allow additive-only changes #1030

Open

alihassanijr reviewed Jul 17, 2026

View reviewed changes

sayakpaul mentioned this pull request Jul 17, 2026

kernel-builder: support SM103 (Blackwell Ultra / B300) capabilities, e.g. 10.3a huggingface/kernels#720

Open

sayakpaul mentioned this pull request Jul 18, 2026

feat: prefer reading the build toml from current ref #982

Merged

moon-bot-app Bot mentioned this pull request Jul 18, 2026

Improve how we vendor kernels that have upstreams (reduce divergence / maintenance burden) huggingface/kernels#721

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

natten: port NATTEN v0.21.6 to kernel-builder#1031

natten: port NATTEN v0.21.6 to kernel-builder#1031
sayakpaul wants to merge 8 commits into
mainfrom
natten

sayakpaul commented Jul 16, 2026

Uh oh!

sayakpaul commented Jul 17, 2026

Uh oh!

github-actions Bot commented Jul 17, 2026 •

edited

Loading

Uh oh!

alihassanijr Jul 17, 2026

Uh oh!

sayakpaul Jul 17, 2026

Uh oh!

alihassanijr Jul 17, 2026

Uh oh!

sayakpaul Jul 17, 2026

Uh oh!

sayakpaul commented Jul 18, 2026

Uh oh!

github-actions Bot commented Jul 18, 2026 •

edited

Loading

Uh oh!

sayakpaul commented Jul 19, 2026

Uh oh!

github-actions Bot commented Jul 19, 2026 •

edited

Loading

Uh oh!

sayakpaul commented Jul 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Uh oh!

Conversation

sayakpaul commented Jul 16, 2026

Uh oh!

sayakpaul commented Jul 17, 2026

Uh oh!

github-actions Bot commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

alihassanijr Jul 17, 2026

Choose a reason for hiding this comment

Uh oh!

sayakpaul Jul 17, 2026

Choose a reason for hiding this comment

Uh oh!

alihassanijr Jul 17, 2026

Choose a reason for hiding this comment

Uh oh!

sayakpaul Jul 17, 2026

Choose a reason for hiding this comment

Uh oh!

sayakpaul commented Jul 18, 2026

Uh oh!

github-actions Bot commented Jul 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sayakpaul commented Jul 19, 2026

Uh oh!

github-actions Bot commented Jul 19, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

sayakpaul commented Jul 20, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

github-actions Bot commented Jul 17, 2026 •

edited

Loading

github-actions Bot commented Jul 18, 2026 •

edited

Loading

github-actions Bot commented Jul 19, 2026 •

edited

Loading