From be52cf660b5b75065abb05bf8946685cec5ab81e Mon Sep 17 00:00:00 2001 From: Anastasiia Filippova Date: Tue, 10 Feb 2026 23:17:28 +0100 Subject: [PATCH] register pressure (#3116) --- mlx/backend/cuda/quantized/fp_quantize.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/mlx/backend/cuda/quantized/fp_quantize.cu b/mlx/backend/cuda/quantized/fp_quantize.cu index 3b4e96ef..7554bef6 100644 --- a/mlx/backend/cuda/quantized/fp_quantize.cu +++ b/mlx/backend/cuda/quantized/fp_quantize.cu @@ -174,7 +174,7 @@ __global__ void fp_quantize_columnwise( auto block_idx = cg::this_thread_block().group_index(); auto idx_in_block = cg::this_thread_block().thread_index(); - constexpr int BLOCK_X = 32; + constexpr int BLOCK_X = 16; constexpr int BLOCK_Y = 32; constexpr int elem_per_byte = (bits == 8) ? 1 : 2; constexpr int bytes_per_group = group_size / elem_per_byte; @@ -323,7 +323,7 @@ fp_dequantize(const uint8_t* w, const uint8_t* scales, T* out, size_t size) { inline std::tuple get_columnwise_quantize_launch_args(size_t size, int group_size, int M, int K) { - constexpr int BLOCK_X = 32; + constexpr int BLOCK_X = 16; constexpr int BLOCK_Y = 32; int rows_per_block = BLOCK_X; int cols_per_block = BLOCK_Y * group_size;