From f5238042ece0be7e14fd43d6a789775ea7b9a6ef Mon Sep 17 00:00:00 2001 From: cocoon Date: Tue, 5 May 2026 02:07:48 +0000 Subject: [PATCH] fix: replace numeric_limits::infinity() with literal to fix CUDA 12.8 NVRTC compilation - cute::numeric_limits::infinity() is not constexpr in CUDA 12.8 NVRTC - Using -1e38f literal as negative infinity workaround (mathematically equivalent for softmax masking) - Fixes JIT compilation failure in smxx_clean_logits kernel on CUDA 12.8 Closes #295 --- deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh b/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh index 2f66b980c5..b3fd90972f 100644 --- a/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh +++ b/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh @@ -17,7 +17,8 @@ void smxx_clean_logits(const uint32_t seq_len, const uint32_t seq_len_kv, const const uint32_t warp_idx = __shfl_sync(0xffffffff, threadIdx.x / 32, 0); constexpr uint32_t kAlignment = 16 / sizeof(logits_dtype_t); - const logits_dtype_t neg_inf = -cute::numeric_limits::infinity(); + // Use raw float literal instead of numeric_limits::infinity() to avoid NVRTC constexpr issues on CUDA 12.8 + const logits_dtype_t neg_inf = logits_dtype_t(-1e38f); // Allocate filled `-inf` shared memory extern __shared__ __align__(1024) logits_dtype_t smem_buffer[];