Skip to content

[CUDA] Fix gather_qmm for experts past 2^32 weight elements #4738

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

Merged
zcbenz merged 3 commits into from
Oct 23, 2026
Merged

[CUDA] Fix gather_qmm for experts past 2^32 weight elements #4738

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
dc6c88b
Fix CUDA gather_qmm for experts past 2^32 weight elements
IamCooller Oct 20, 2026
c6fcf86
remote test as it is too slow and the change is trivial
zcbenz Oct 23, 2026
4ca67f4
nit
zcbenz Oct 23, 2026
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
4 changes: 2 additions & 2 deletions mlx/backend/cuda/device/gemm_sm70.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -220,9 +220,9 @@ CUTE_DEVICE void gemm_sm70_mainloop(
template <bool KMajor>
inline constexpr auto make_matrix_stride(int m, int k) {
if constexpr (KMajor) {
return cute::make_stride(k, cute::Int<1>{}, m * k);
return cute::make_stride(k, cute::Int<1>{}, int64_t(m) * k);
} else {
return cute::make_stride(cute::Int<1>{}, m, m * k);
return cute::make_stride(cute::Int<1>{}, m, int64_t(m) * k);
}
}

Expand Down
12 changes: 6 additions & 6 deletions mlx/backend/cuda/device/qmm_sm80.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -303,9 +303,9 @@ void qmm_sm80_kernel(
int l_coord = int(blockIdx.z);

// Define layouts (mixed).
auto dA = make_stride(k, Int<1>{}, m * k); // (dM,dK,dL)
auto dB = make_stride(k, Int<1>{}, n * k); // (dN,dK,dL)
auto dC = make_stride(n, Int<1>{}, m * n); // (dM,dN,dL)
auto dA = make_stride(k, Int<1>{}, int64_t(m) * k); // (dM,dK,dL)
auto dB = make_stride(k, Int<1>{}, int64_t(n) * k); // (dN,dK,dL)
auto dC = make_stride(n, Int<1>{}, int64_t(m) * n); // (dM,dN,dL)
auto S_layout = make_scales_layout<GroupSize>(n, k, l);

// Handle broadcasting.
Expand Down Expand Up @@ -388,9 +388,9 @@ void qmm_sm80_rhs_kernel(

// Define layouts (mixed). All gathered rows form one matrix; each run of
// equal rhs_indices inside the row tile reuses one expert tile.
auto dA = make_stride(k, Int<1>{}, m * k); // (dM,dK,dL)
auto dB = make_stride(k, Int<1>{}, n * k); // (dN,dK,dL)
auto dC = make_stride(n, Int<1>{}, m * n); // (dM,dN,dL)
auto dA = make_stride(k, Int<1>{}, int64_t(m) * k); // (dM,dK,dL)
auto dB = make_stride(k, Int<1>{}, int64_t(n) * k); // (dN,dK,dL)
auto dC = make_stride(n, Int<1>{}, int64_t(m) * n); // (dM,dN,dL)
auto S_layout = make_scales_layout<GroupSize>(n, k, group_count);

Tensor mA_mkl = make_tensor(make_gmem_ptr(A), make_shape(m, k, 1), dA); // (M,K,L)
Expand Down
8 changes: 4 additions & 4 deletions mlx/backend/cuda/quantized/qmm/qmm_sm90.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,10 +124,10 @@ void qmm_sm90(
const char* tag = "[quantized_matmul]";
auto [m, n, k, l, broadcast_b] = make_problem_shape(x, w, out);

auto dA = make_stride(int64_t(k), Int<1>{}, int64_t(m * k));
auto dB = make_stride(int64_t(k), Int<1>{}, int64_t(n * k));
auto dS = make_stride(Int<1>{}, int64_t(n), int64_t(n * k / group_size));
auto dD = make_stride(Int<1>{}, int64_t(n), int64_t(m * n));
auto dA = make_stride(int64_t(k), Int<1>{}, int64_t(m) * k);
auto dB = make_stride(int64_t(k), Int<1>{}, int64_t(n) * k);
auto dS = make_stride(Int<1>{}, int64_t(n), int64_t(n) * k / group_size);
auto dD = make_stride(Int<1>{}, int64_t(n), int64_t(m) * n);
if (broadcast_b) {
get<2>(dB) = 0;
get<2>(dS) = 0;
Expand Down
Loading