From 0decac0d96c42b49572498019f0a0e3600f50398 Mon Sep 17 00:00:00 2001 From: khluu Date: Wed, 3 Jun 2026 17:11:47 -0700 Subject: [PATCH] fix: resolve CUTLASS fmin compatibility for DeepSeek-V4 init Signed-off-by: khluu --- .../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 0eba82126af..7b329c941ca 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 @@ -320,11 +320,11 @@ class SparseAttnCompressNormRopeStoreC4Kernel: 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), ) @@ -978,11 +978,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), )