Skip to content

[Metal] Skip empty NAX GEMM output groups #3941

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
XXXXRT666 wants to merge 1 commit into
base: main
Choose a base branch
from
+21 −8
Open

[Metal] Skip empty NAX GEMM output groups #3941

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
10 changes: 10 additions & 0 deletions mlx/backend/metal/kernels/steel/gemm/gemm_nax.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,11 +45,17 @@ auto gemm_loop(
NAXTile<AccumType, TM, TN> Dtile;
Dtile.clear();

const bool has_output = sgp_sm > 0 && sgp_sn > 0;

int gemm_k_iterations_ = gemm_k_iterations_aligned;

STEEL_PRAGMA_NO_UNROLL
for (int kk0 = 0; kk0 < gemm_k_iterations_; kk0++) {
threadgroup_barrier(mem_flags::mem_none);
if constexpr (!kAlignedM || !kAlignedN) {
if (!has_output)
continue;
}

STEEL_PRAGMA_NO_UNROLL
for (int kk1 = 0; kk1 < BK; kk1 += SK) {
Expand Down Expand Up @@ -94,6 +100,10 @@ auto gemm_loop(

if constexpr (!kAlignedK) {
simdgroup_barrier(mem_flags::mem_none);
if constexpr (!kAlignedM || !kAlignedN) {
if (!has_output)
return Dtile;
}

const short rem_bk = K - gemm_k_iterations_ * BK;

Expand Down
19 changes: 11 additions & 8 deletions mlx/backend/metal/kernels/steel/gemm/kernels/steel_gemm_fused_nax.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,14 +200,17 @@ template <
params->gemm_k_iterations_aligned,
sgp_sm,
sgp_sn);
if (use_out_source) {
gemm_epilogue<kAlignedM.value, kAlignedN.value>(
Dtile, C, params, addmm_params, sgp_sm, sgp_sn);
}
if constexpr (kAlignedM && kAlignedN) {
Dtile.store(D, int(params->ldd));
} else {
Dtile.store_safe(D, int(params->ldd), short2(sgp_sn, sgp_sm));
if ((kAlignedM.value || sgp_sm > 0) &&
(kAlignedN.value || sgp_sn > 0)) {

@zcbenz zcbenz Aug 7, 2026

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

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

This is only skipping the epilogue, should the whole gemm_loop be skipped too?

@XXXXRT666 XXXXRT666 Aug 7, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

The GEMM work is already skipped inside gemm_loop, but the call itself cannot be skipped per SIMD group because the loop contains a threadgroup_barrier for every aligned-K iteration. An edge threadgroup can contain both useful and empty SIMD groups, so returning before the call would leave the useful groups waiting at the barrier.
Empty groups therefore still enter gemm_loop and participate in each barrier, but immediately continue before tile loads and matmad. The epilogue and store have no threadgroup-wide synchronization, so they can be skipped entirely.

if (use_out_source) {
gemm_epilogue<kAlignedM.value, kAlignedN.value>(
Dtile, C, params, addmm_params, sgp_sm, sgp_sn);
}
if constexpr (kAlignedM && kAlignedN) {
Dtile.store(D, int(params->ldd));
} else {
Dtile.store_safe(D, int(params->ldd), short2(sgp_sn, sgp_sm));
}
}
});
});
Expand Down
Loading