From 597bc1593635b2f5ee841b3b9162af64a04d2d28 Mon Sep 17 00:00:00 2001 From: Willow Lopez <100782273+Oxygen56@users.noreply.github.com> Date: Wed, 3 Jun 2026 13:07:10 +0800 Subject: [PATCH] fix: resolve CUTLASS fmin compatibility for DeepSeek-V4 init (#44236) Signed-off-by: Willow Lopez <100782273+Oxygen56@users.noreply.github.com> --- .../nvidia/ops/sparse_attn_compress_cutedsl.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/vllm/models/deepseek_v4/nvidia/ops/sparse_attn_compress_cutedsl.py b/vllm/models/deepseek_v4/nvidia/ops/sparse_attn_compress_cutedsl.py index 1e6c2ed829..ed16ca6d3b 100644 --- a/vllm/models/deepseek_v4/nvidia/ops/sparse_attn_compress_cutedsl.py +++ b/vllm/models/deepseek_v4/nvidia/ops/sparse_attn_compress_cutedsl.py @@ -370,11 +370,11 @@ class SparseAttnCompressNormRopeStoreC4Kernel: inv_scale = _recast_val((Uint32(254) - ue8m0) << Uint32(23), Float32) for pair in cutlass.range_constexpr(self.elems_per_lane // 2): elem = const_expr(pair * 2) - y0 = cute.arch.fmin( + y0 = cutlass.min( cute.arch.fmax(q[elem] * inv_scale, Float32(-self.fp8_max)), Float32(self.fp8_max), ) - y1 = cute.arch.fmin( + y1 = cutlass.min( cute.arch.fmax(q[elem + 1] * inv_scale, Float32(-self.fp8_max)), Float32(self.fp8_max), ) @@ -1026,11 +1026,11 @@ class SparseAttnNormRopeStoreKernel: bits = _recast_val(scale_raw, Uint32) ue8m0 = ((bits + Uint32(0x7FFFFF)) >> Uint32(23)) & Uint32(0xFF) inv_scale = _recast_val((Uint32(254) - ue8m0) << Uint32(23), Float32) - y0 = cute.arch.fmin( + y0 = cutlass.min( cute.arch.fmax(q0 * inv_scale, Float32(-self.fp8_max)), Float32(self.fp8_max), ) - y1 = cute.arch.fmin( + y1 = cutlass.min( cute.arch.fmax(q1 * inv_scale, Float32(-self.fp8_max)), Float32(self.fp8_max), )