Skip to content

fix: replace numeric_limits::infinity() with literal to fix CUDA 12.8 NVRTC compilation #327

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
kuishou68 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

fix: replace numeric_limits::infinity() with literal to fix CUDA 12.8 NVRTC compilation #327

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<logits_dtype_t>::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[];
Expand Down