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.

alihassanijr · 2026-07-23T14:33:17Z

@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.

Would love to help @sayakpaul , but as is I'm struggling to find resources to build the standard wheels for 0.21.7 🤣

But cutting down with +PTX makes sense to me.
And no it shouldn't create any issues, because as far as libnatten is concerned, cutlass-fna/fmha compiles if SM80 and later, hopper-fna/fmha if SM90a, and blackwell-fna/fmha if SM100a/SM103a are specified.

Figuring out shared memory limit, supported tile shapes / configs, and the rest are all determined at runtime based on the actual device.

Beyond that it's just handled by the cuda compiler and how it builds fatbins, with the exception of arch guards we place, but again those are pretty relaxed for cutlass-fna/fmha:

** Disregard ** the fallthru error message -- it should read "built for SM80 - SM120"

sayakpaul · 2026-07-23T14:51:20Z

Would love to help @sayakpaul , but as is I'm struggling to find resources to build the standard wheels for 0.21.7 🤣

Is it related to kernels? To me, it seems like related to how you distribute the wheels?

Noted regarding the rest of what you said. Shall I wait for you till you build?

alihassanijr · 2026-07-23T17:40:35Z

Would love to help @sayakpaul , but as is I'm struggling to find resources to build the standard wheels for 0.21.7 🤣

Is it related to kernels? To me, it seems like related to how you distribute the wheels?

Noted regarding the rest of what you said. Shall I wait for you till you build?

Nope, unrelated to kernels. I just don't have a standard way of building wheels. Every release I try to find resources to build, sometimes it's easier and quicker, sometimes it's more involved.

For this release, I'm just going to run the full pipeline to make sure it's stable, and tag so that kernels can pick up the latest fixes and perf improvements, and then figure out where/how to build wheels.

sayakpaul · 2026-07-24T00:49:30Z

Sounds good! Keep me posted.

…

On Thu, 23 Jul 2026 at 23:10, Ali Hassani ***@***.***> wrote: *alihassanijr* left a comment (huggingface/kernels-community#1031) <#1031 (comment)> Would love to help @sayakpaul <https://github.com/sayakpaul> , but as is I'm struggling to find resources to build the standard wheels for 0.21.7 🤣 Is it related to kernels? To me, it seems like related to how you distribute the wheels? Noted regarding the rest of what you said. Shall I wait for you till you build? Nope, unrelated to kernels. I just don't have a standard way of building wheels. Every release I try to find resources to build, sometimes it's easier and quicker, sometimes it's more involved. For this release, I'm just going to run the full pipeline to make sure it's stable, and tag so that kernels can pick up the latest fixes and perf improvements, and then figure out where/how to build wheels. — Reply to this email directly, view it on GitHub <#1031?email_source=notifications&email_token=AFPE2TDOUGL53C5N2KHBWXT5GJE2VA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKMBWGE2TCNRWG4Z2M4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-5061516673>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AFPE2TBVL2OMMO5UZQQW7GT5GJE2VAVCNFSNUABGKJSXA33TNF2G64TZHMYTANZRGQ3TKNJSHE5US43TOVSTWNBZGAYTCOJSGI3TPILWAI> . Triage notifications, keep track of coding agent tasks and review pull requests on the go with GitHub Mobile for iOS <https://github.com/notifications/mobile/ios/AFPE2TE2KPFJ5N27TL3TVYD5GJE2VA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKMBWGE2TCNRWG4Z2M4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJKTGN5XXIZLSL5UW64Y> and Android <https://github.com/notifications/mobile/android/AFPE2TBQDHW2AH6EZKLM3SD5GJE2VA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTKMBWGE2TCNRWG4Z2M4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLTGN5XXIZLSL5QW4ZDSN5UWI>. Download it today! You are receiving this because you were mentioned.Message ID: ***@***.***>

alihassanijr · 2026-07-26T18:51:45Z

@sayakpaul 0.21.7 is tagged and live, wheels are also up:

https://github.com/SHI-Labs/NATTEN/releases/tag/v0.21.7

Header-only C++ changes (fna iterators/kernels, tokperm padding + H*D alignment), runtime tile-config selection now limits large-smem tiles to cc 80/90/100/103, and the three upstream test updates. No autogen changes between v0.21.6 and v0.21.7, so no regen needed. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

Release builds (9 variants/platform x ~2.5-3h serial) exceeded the 20h CI job timeout. Drop 7.0/7.5/8.6/8.9 from the core section, matching flash-attn2's arch set: sm_86/89 run sm_80 SASS via within-major forward compat, and NATTEN's runtime config selection (v0.21.7) restricts large-smem tiles to cc 80/90/100/103. Also let natten use the full 32 vCPUs of the build runners (2 jobs x 16 cores, up from the 2x12 default). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

sayakpaul · 2026-07-28T05:49:13Z

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

github-actions · 2026-07-28T05:49:26Z

Build request processed.

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

Dispatched (1):

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

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

sayakpaul and others added 2 commits July 28, 2026 11:03

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!

alihassanijr commented Jul 23, 2026

Uh oh!

sayakpaul commented Jul 23, 2026

Uh oh!

alihassanijr commented Jul 23, 2026

Uh oh!

sayakpaul commented Jul 24, 2026 via email

Uh oh!

alihassanijr commented Jul 26, 2026

Uh oh!

sayakpaul commented Jul 28, 2026

Uh oh!

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

Uh oh!

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

github-actions Bot commented Jul 28, 2026 •

edited

Loading