Skip to content

megablocks: add ROCm backend (AITER Triton grouped GEMM) - #916

Merged
danieldk merged 1 commit into
huggingface:mainfrom
Abdennacer-Badaoui:megablocks-rocm-add
Jun 4, 2026
Merged

danieldk merged 1 commit into
huggingface:mainfrom
Abdennacer-Badaoui:megablocks-rocm-add

Conversation

@Abdennacer-Badaoui

Copy link
Copy Markdown
Member

Adds AMD ROCm support to the megablocks kernel alongside the existing cuda/xpu/cpu backends.

  • New [kernel.megablocks_rocm] block (gfx942 + gfx950) reusing the same CUDA op sources via hipify; two small CUDA-safe guards (_LDG for __ldg, MEGABLOCKS_CUB for cub::DeviceHistogram) plug the gaps in torch's hipify symbol table.
  • CUTLASS grouped GEMM is CUDA-only; on ROCm grouped GEMM is provided by AITER's Triton gmm/ptgmm, vendored under torch-ext/megablocks/_grouped_gemm_triton/ with an adapter mapping megablocks' gmm(a, b, c, batch_sizes, trans_a, trans_b) semantics. grouped_gemm/backend.py dispatches by torch.version.hip.
  • torch_binding.cpp registers the 5 ops under ROCM_KERNEL (HIP → torch::kCUDA); CUTLASS gmm stays CUDA-only.
  • Drive-by: fixes a latent shadowing bug in megablocks/__init__.py where from . import layers rebinds ops to the python submodule, breaking the top-level megablocks.sort wrapper on both backends.

Validated on gfx942 (MI300X): ops correctness, grouped-GEMM autograd, full MegaBlocksMoeMLP forward; both via local cpp-extension build (torch 2.10) and via the real nix kernel-builder artifact (torch212-cxx11-rocm72) loaded in the matching dev shell.

@Abdennacer-Badaoui
Abdennacer-Badaoui requested a review from drbh as a code owner June 2, 2026 17:38
@danieldk

danieldk commented Jun 3, 2026

Copy link
Copy Markdown
Member

/kernel-bot build megablocks

@github-actions

github-actions Bot commented Jun 3, 2026 •

Copy link
Copy Markdown

Build request processed.

Command: /kernel-bot build megablocks
Mode: build only
Target branch: pr-916
PR head SHA: b117abfdca98060d98a6a47178529a4b86f60de1
Workflows: build.yaml, build-mac.yaml, build-windows.yaml

Dispatched (1):

Comment on lines +7 to +8
from ._ops import ops as _compiled_ops
from . import ops

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

can you explain why we need both _compiled_ops and ops exposed?

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Inside the megablocks package there are two different "ops" things that happen to share the same name:

  1. The C++ kernels (compiled .so) : sort, histogram, etc. as raw Torch ops. Fast, no autograd, takes all the args (e.g. sort(x, end_bit, x_out, iota_out) : 4 args).
  2. A Python folder also called ops/ : wrappers around those C++ kernels that add autograd support. Same names, different signatures (e.g. sort(x, end_bit=None) : 2 args, returns a tuple).

@Abdennacer-Badaoui Abdennacer-Badaoui Jun 3, 2026 •

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Before my change, the single name ops was bound to the compiled namespace, then silently overwritten to the submodule when from . import layers ran ; so half the file got one and half got the other depending on call order. That's the bug that broke megablocks.sort. Giving them distinct names makes the intent explicit and removes the order-dependence.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

thanks for the explanation and the fix!

@drbh drbh left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

looks good to me!

@danieldk

danieldk commented Jun 4, 2026

Copy link
Copy Markdown
Member

/kernel-bot merge-and-upload megablocks

@github-actions

github-actions Bot commented Jun 4, 2026 •

Copy link
Copy Markdown

Build request failed.

Command: /kernel-bot merge-and-upload megablocks
Mode: merge, build and upload
Target branch: ``
PR head SHA: b117abfdca98060d98a6a47178529a4b86f60de1
Workflows: `build.yaml, build-mac.yaml, build-windows.yaml`

Failure: Failed to merge PR before build/upload. Check mergeability and required checks.

@danieldk
danieldk merged commit f8f8b50 into huggingface:main Jun 4, 2026
4 of 6 checks passed
@danieldk

danieldk commented Jun 4, 2026

Copy link
Copy Markdown
Member

/kernel-bot merge-and-upload megablocks

@github-actions

github-actions Bot commented Jun 4, 2026 •

Copy link
Copy Markdown

Build request processed.

Command: /kernel-bot merge-and-upload megablocks
Mode: merge, build and upload
Target branch: ``
PR head SHA: b117abfdca98060d98a6a47178529a4b86f60de1
Workflows: `build.yaml, build-mac.yaml, build-windows.yaml`

Merge result: PR is already merged. Continuing with build/upload.

Dispatched (1):

sayakpaul pushed a commit that referenced this pull request Jun 5, 2026
sayakpaul added a commit that referenced this pull request Jun 5, 2026
sayakpaul added a commit that referenced this pull request Jun 5, 2026
* feat: allow kernel-bot to run security analysis.

* include security in the build.

* megablocks for rocm (#916)

* feat: update merge-and-upload permissions (#930)

* [Liger Kernels] Adding TiledMLP kernels (#918)

* chore: adding liger tiled mlp kernels

* liger-kernels: test with Triton ROCm binary wheels

* chore: align layer to pure

---------

Co-authored-by: Daniël de Kok <me@danieldk.eu>

* finegrained-fp8: MXFP4 recipes (#923)

* add fp4 matmul kernels for deepseek v4 flash

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* add missing source code..

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* update

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* update to w4a8

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* update

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* perf optimization add add benchmark_perf.py script

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>

* structure

* standarize all kernels and testing around the same api

* better testing framework

* fix batched with sentinels

* tests with sentinels as well

* test static and support ue8m0

* specify v2 as the version

* cleanup

* naming convention

* update readme

* remove bench

* format

* update readme

* trigger

* skip tests on < sm89

* drop sm89 testing, too error prone

* Use triton-rocm-fixes branch to pass build

* Update build.toml

* Fix Torch ops

* Remove custom path insertion

This breaks tests in kernel-builder devshells

---------

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>
Co-authored-by: Wang, Yi <yi.a.wang@intel.com>
Co-authored-by: Daniël de Kok <me@danieldk.eu>

* address review feedback

* Revert "megablocks for rocm (#916)"

This reverts commit e06068f.

* Revert "feat: update merge-and-upload permissions (#930)"

This reverts commit a913728.

* Revert "[Liger Kernels] Adding TiledMLP kernels (#918)"

This reverts commit 3d9bb2f.

* Revert "finegrained-fp8: MXFP4 recipes (#923)"

This reverts commit 9f9d5e1.

* security-and-build

---------

Signed-off-by: Wang, Yi <yi.a.wang@intel.com>
Co-authored-by: BADAOUI Abdennacer <106801897+Abdennacer-Badaoui@users.noreply.github.com>
Co-authored-by: drbh <david.richard.holtz@gmail.com>
Co-authored-by: Aritra Roy Gosthipaty <aritra.born2fly@gmail.com>
Co-authored-by: Daniël de Kok <me@danieldk.eu>
Co-authored-by: Ilyas Moutawwakil <57442720+IlyasMoutawwakil@users.noreply.github.com>
Co-authored-by: Wang, Yi <yi.a.wang@intel.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants