A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 2, 03:30 UTC
fix: resolve CUTLASS fmin compatibility for DeepSeek-V4 init (#44236)
Signed-off-by: Willow Lopez <100782273+Oxygen56@users.noreply.github.com>
Willow Lopez committed
Jun 3, 2026 at 13:07 UTC
597bc1593635b2f5ee841b3b9162af64a04d2d28
1 file changed
+4
-4
vllm/models/deepseek_v4/nvidia/ops/sparse_attn_compress_cutedsl.py
+4
-4
@@ -370,11 +370,11 @@ class SparseAttnCompressNormRopeStoreC4Kernel:
370
inv_scale = _recast_val((Uint32(254) - ue8m0) << Uint32(23), Float32)
371
for pair in cutlass.range_constexpr(self.elems_per_lane // 2):
372
elem = const_expr(pair * 2)
373
- y0 = cute.arch.fmin(
373
+ y0 = cutlass.min(
374
cute.arch.fmax(q[elem] * inv_scale, Float32(-self.fp8_max)),
375
Float32(self.fp8_max),
376
)
377
- y1 = cute.arch.fmin(
377
+ y1 = cutlass.min(
378
cute.arch.fmax(q[elem + 1] * inv_scale, Float32(-self.fp8_max)),
379
Float32(self.fp8_max),
380
)
@@ -1026,11 +1026,11 @@ class SparseAttnNormRopeStoreKernel:
1026
bits = _recast_val(scale_raw, Uint32)
1027
ue8m0 = ((bits + Uint32(0x7FFFFF)) >> Uint32(23)) & Uint32(0xFF)
1028
inv_scale = _recast_val((Uint32(254) - ue8m0) << Uint32(23), Float32)
1029
- y0 = cute.arch.fmin(
1029
+ y0 = cutlass.min(
1030
cute.arch.fmax(q0 * inv_scale, Float32(-self.fp8_max)),
1031
Float32(self.fp8_max),
1032
)
1033
- y1 = cute.arch.fmin(
1033
+ y1 = cutlass.min(
1034
cute.arch.fmax(q1 * inv_scale, Float32(-self.fp8_max)),
1035
Float32(self.fp8_max),
1036
)