deepseek-ai/DeepGEMM
1
0
Fork 0
mirror of https://github.com/deepseek-ai/DeepGEMM.git synced 2026-11-05 00:47:26 +00:00
Code Issues 162 Projects Packages Wiki Activity

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

New issue
Open
opened 2026-10-26 21:16:56 +00:00 by bsakaguchi · 0 comments
bsakaguchi commented 2026-10-26 21:16:56 +00:00
Owner
Copy link

📋 Pull Request Information

Original PR: https://github.com/deepseek-ai/DeepGEMM/pull/327
Author: @kuishou68
Created: 5/5/2026
Status: 🔄 Open

Base: main ← Head: fix-cuda128-nvrtc-constexpr

📝 Commits (1)

  • f523804 fix: replace numeric_limits::infinity() with literal to fix CUDA 12.8 NVRTC compilation

📊 Changes

1 file changed (+2 additions, -1 deletions)

View changed files

📝 deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh (+2 -1)

📄 Description

Problem

On CUDA 12.8, NVRTC fails to compile smxx_clean_logits.cuh with:

error: expression must have a constant value
constexpr float neg_inf = -cute::numeric_limits<float>::infinity();

cuda::std::numeric_limits<T>::infinity() is not constexpr in CUDA 12.8 NVRTC environment.

Solution

Replace with raw float literal -1e38f, which is a true compile-time constant and mathematically equivalent to negative infinity for softmax masking purposes.

Testing

  • Fixes JIT compilation on CUDA 12.8 (H100)
  • Preserves numerical correctness for attention score masking
  • Compatible with both float32 and bfloat16 logits types

Closes #295

🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.

## 📋 Pull Request Information **Original PR:** https://github.com/deepseek-ai/DeepGEMM/pull/327 **Author:** [@kuishou68](https://github.com/kuishou68) **Created:** 5/5/2026 **Status:** 🔄 Open **Base:** `main` ← **Head:** `fix-cuda128-nvrtc-constexpr` --- ### 📝 Commits (1) - [`f523804`](https://github.com/deepseek-ai/DeepGEMM/commit/f5238042ece0be7e14fd43d6a789775ea7b9a6ef) fix: replace numeric_limits::infinity() with literal to fix CUDA 12.8 NVRTC compilation ### 📊 Changes **1 file changed** (+2 additions, -1 deletions) <details> <summary>View changed files</summary> 📝 `deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh` (+2 -1) </details> ### 📄 Description ## Problem On CUDA 12.8, NVRTC fails to compile `smxx_clean_logits.cuh` with: ``` error: expression must have a constant value constexpr float neg_inf = -cute::numeric_limits<float>::infinity(); ``` `cuda::std::numeric_limits<T>::infinity()` is not `constexpr` in CUDA 12.8 NVRTC environment. ## Solution Replace with raw float literal `-1e38f`, which is a true compile-time constant and mathematically equivalent to negative infinity for softmax masking purposes. ## Testing - Fixes JIT compilation on CUDA 12.8 (H100) - Preserves numerical correctness for attention score masking - Compatible with both float32 and bfloat16 logits types Closes #295 --- <sub>🔄 This issue represents a GitHub Pull Request. It cannot be merged through Gitea due to API limitations.</sub>
Sign in to join this conversation.
No Branch/Tag specified
Branches Tags
main
public-sync-20261014
0930-release
nv_dev
fix-mega-moe-fence
public-release-2609
public-release-2607
public-release-2606
mega-update
public-release-260416
ko3n1g/ci/build-wheels
nv_dev_f8e8fb5
nv_dev_950d31a
nv_dev_67fc648
nv_dev_c491439
nv_dev_4ff3f54
v2.1.1.post3
v2.1.1.post2
v2.1.1.post1
v2.1.1
v2.1.0
Labels
Clear labels   
pull-request
Mirrored from GitHub Pull Request

No labels
pull-request
Milestone
Clear milestone
No items
No milestone
Projects
Clear projects
No items
No project
Assignees
Clear assignees
No assignees
1 participant
Notifications
Due date
The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference
deepseek-ai/DeepGEMM#368
No description provided.