[GH-ISSUE #295] [Bug]: NVRTC JIT compilation fails on CUDA 12.8 (smxx_clean_logits.cuh: expression must have a constant value) #150

New issue

Open

opened 2026-10-26 21:13:43 +00:00 by bsakaguchi · 0 comments

bsakaguchi commented

2026-10-26 21:13:43 +00:00

Owner

Originally created by @LilySu on GitHub (Mar 25, 2026).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/295

Describe the bug

When running DeepGEMM on CUDA 12.8, the fp8_mqa_logits kernel fails to JIT-compile via NVRTC. The compiler throws an error stating that cute::numeric_limits<float>::infinity() cannot be used as a constexpr.

Error Log

NVRTC log: "kernel.cu": creating precompiled header file "kernel.pch"
/usr/local/lib/python3.12/dist-packages/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh(17): error: expression must have a constant value
      constexpr float neg_inf = -cute::numeric_limits<float>::infinity();
                                ^
/usr/local/cuda/include/cuda/std/detail/libcxx/include/limits(686): note #2703-D: cannot call non-constexpr function "cuda::std::__4::__libcpp_numeric_limits<float, true>::infinity"

Root Cause & Investigation

In CUDA 12.8 device code, the underlying implementation of cuda::std::numeric_limits<float>::infinity() inside libcxx/include/limits is missing the constexpr qualifier under NVRTC's stripped-down compilation environment.

What we tried before finding the fix

Cleared JIT cache (/root/.deep_gemm, /workspace/.deep_gemm_cache/cache).
Reinstalled DeepGEMM non-editable (pip install --no-build-isolation .).
Removed stale editable install artifacts (.egg-info, .pth files).
Set CPLUS_INCLUDE_PATH to CUTLASS headers (NVRTC ignores it).
Replaced with standard built-ins (__builtin_inff()) — NVRTC strips host-side built-ins and throws an undefined identifier error.
Replaced with native CUDA intrinsics (__int_as_float(0xff800000)) — NVRTC refuses to evaluate it as a constant expression.
Implemented a PyTorch try/except fallback using torch.einsum to bypass the crashed kernel — this successfully unblocked the model run, confirming the JIT failure was the only roadblock.
Only a -1e38f raw literal worked as a true compile-time constant for NVRTC.

Proposed Fix / Workaround

To maintain compatibility with CUDA 12.8 without requiring users to upgrade to 12.9, replacing the standard library call with a raw float literal fixes the NVRTC compilation instantly.

In deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh (Line 17):
Change from:
constexpr float neg_inf = -cute::numeric_limits<float>::infinity();

Change to:
constexpr float neg_inf = -1e38f;

Since -1e38f mathematically acts as negative infinity to mask out attention scores during the Softmax step, it preserves accuracy while bypassing the NVRTC constexpr limitation.

Environment

DeepGEMM Version: Version 2.3.0, commit d30fc36
GPU: 1x NVIDIA H100 SXM (80GB HBM3, SM90)
System Resources: 26 vCPUs, 251 GB Memory
Environment: Docker container (Template: runpod-torch-v280)
CUDA Version: 12.8 (System)
Python Version: 3.12
PyTorch Version: 2.8.0+cu128

Originally created by @LilySu on GitHub (Mar 25, 2026). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/295 ### Describe the bug When running DeepGEMM on CUDA 12.8, the `fp8_mqa_logits` kernel fails to JIT-compile via NVRTC. The compiler throws an error stating that `cute::numeric_limits<float>::infinity()` cannot be used as a `constexpr`. ### Error Log ```text NVRTC log: "kernel.cu": creating precompiled header file "kernel.pch" /usr/local/lib/python3.12/dist-packages/deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh(17): error: expression must have a constant value constexpr float neg_inf = -cute::numeric_limits<float>::infinity(); ^ /usr/local/cuda/include/cuda/std/detail/libcxx/include/limits(686): note #2703-D: cannot call non-constexpr function "cuda::std::__4::__libcpp_numeric_limits<float, true>::infinity" ``` ### Root Cause & Investigation In CUDA 12.8 device code, the underlying implementation of `cuda::std::numeric_limits<float>::infinity()` inside `libcxx/include/limits` is missing the `constexpr` qualifier under NVRTC's stripped-down compilation environment. ### What we tried before finding the fix - Cleared JIT cache (`/root/.deep_gemm`, `/workspace/.deep_gemm_cache/cache`). - Reinstalled DeepGEMM non-editable (`pip install --no-build-isolation .`). - Removed stale editable install artifacts (`.egg-info`, `.pth` files). - Set `CPLUS_INCLUDE_PATH` to CUTLASS headers (NVRTC ignores it). - Replaced with standard built-ins (`__builtin_inff()`) — NVRTC strips host-side built-ins and throws an undefined identifier error. - Replaced with native CUDA intrinsics (`__int_as_float(0xff800000)`) — NVRTC refuses to evaluate it as a constant expression. - Implemented a PyTorch `try/except` fallback using `torch.einsum` to bypass the crashed kernel — this successfully unblocked the model run, confirming the JIT failure was the only roadblock. - **Only a `-1e38f` raw literal worked as a true compile-time constant for NVRTC.** ### Proposed Fix / Workaround To maintain compatibility with CUDA 12.8 without requiring users to upgrade to 12.9, replacing the standard library call with a raw float literal fixes the NVRTC compilation instantly. In `deep_gemm/include/deep_gemm/impls/smxx_clean_logits.cuh` (Line 17): **Change from:** `constexpr float neg_inf = -cute::numeric_limits<float>::infinity();` **Change to:** `constexpr float neg_inf = -1e38f;` Since `-1e38f` mathematically acts as negative infinity to mask out attention scores during the Softmax step, it preserves accuracy while bypassing the NVRTC `constexpr` limitation. ### Environment * **DeepGEMM Version:** Version 2.3.0, commit d30fc36 * **GPU:** 1x NVIDIA H100 SXM (80GB HBM3, SM90) * **System Resources:** 26 vCPUs, 251 GB Memory * **Environment:** Docker container (Template: `runpod-torch-v280`) * **CUDA Version:** 12.8 (System) * **Python Version:** 3.12 * **PyTorch Version:** 2.8.0+cu128

No labels

pull-request

No milestone

No project

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#150

No description provided.

Rows
Columns