diff --git a/mlx/backend/cuda/device/gemm_sm70.cuh b/mlx/backend/cuda/device/gemm_sm70.cuh index f66b890750..63b38aee56 100644 --- a/mlx/backend/cuda/device/gemm_sm70.cuh +++ b/mlx/backend/cuda/device/gemm_sm70.cuh @@ -220,9 +220,9 @@ CUTE_DEVICE void gemm_sm70_mainloop( template 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); } } diff --git a/mlx/backend/cuda/device/qmm_sm80.cuh b/mlx/backend/cuda/device/qmm_sm80.cuh index 476c2339cf..5fea5b51b2 100644 --- a/mlx/backend/cuda/device/qmm_sm80.cuh +++ b/mlx/backend/cuda/device/qmm_sm80.cuh @@ -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(n, k, l); // Handle broadcasting. @@ -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(n, k, group_count); Tensor mA_mkl = make_tensor(make_gmem_ptr(A), make_shape(m, k, 1), dA); // (M,K,L) diff --git a/mlx/backend/cuda/quantized/qmm/qmm_sm90.cu b/mlx/backend/cuda/quantized/qmm/qmm_sm90.cu index 7bc5d83437..546e1755db 100644 --- a/mlx/backend/cuda/quantized/qmm/qmm_sm90.cu +++ b/mlx/backend/cuda/quantized/qmm/qmm_sm90.cu @@ -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;