[Clang][HIP] Include `__clang_cuda_math_forward_declares.h` before `<cmath>` by jmmartinez · Pull Request #201563 · llvm/llvm-project

jmmartinez · 2026-06-04T12:23:01Z

In HIP, constexpr functions are treated as both __host__ and __device__.

A new version of the MS STL shipped with the build tools version
14.51.36231 has constexpr definitions for some cmath functions when the
compiler in use is Clang (this gets worse when C++23 is in use).

These definitions conflict with the __device__ declarations we provide
in the header wrappers.

There is a workaround for this: We do not mark constexpr
functions that are defined in a system header as
__host__ and __device__ if there is a previous __device__
declaration.

By moving __clang_cuda_math_forward_declares.h before <cmath> is
included we're able to benefit from this behavior.

This fixes error like this one ggml-org/llama.cpp#22570
even for recent versions of C++.

This patch should address C++23 concerns where more functions are marked as constexpr.
However, we should still figure out how to provide __device__ constexpr versions for those functions, or not provide ours and use the system's if they exist.

This patch replaces #200395

llvmorg-github-actions · 2026-06-04T12:23:39Z

@llvm/pr-subscribers-clang

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

In HIP, constexpr functions are treated as both __host__ and __device__.

A new version of the MS STL shipped with the build tools version
14.51.36231 has constexpr definitions for some cmath functions when the
compiler in use is Clang (this gets worse when C++23 is in use).

These definitions conflict with the __device__ declarations we provide
in the header wrappers.

There is a workaround for this: We do not mark constexpr
functions that are defined in a system header as
__host__ and __device__ if there is a previous __device__
declaration.

By moving __clang_cuda_math_forward_declares.h before <cmath> is
included we're able to benefit from this behavior.

This fixes error like this one ggml-org/llama.cpp#22570
even for recent versions of C++.

This patch replaces #200395

Full diff: https://github.com/llvm/llvm-project/pull/201563.diff

2 Files Affected:

(modified) clang/lib/Headers/__clang_hip_runtime_wrapper.h (+4-1)
(added) clang/test/Headers/hip-constexpr-cmath.hip (+70)

diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
index 19ce7a5d2c86b..72494fdda4ec9 100644
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -111,6 +111,10 @@ __attribute__((weak)) inline __device__ void free(void *__ptr) {
 #endif //__cplusplus
 
 #if !defined(__HIPCC_RTC__)
+// We must include the forward declarations before cmath is included.
+// Otherwise `constexpr` functions in <cmath> would be implicitely __host__ __device__.
+// Declaring the __device__ verison before allows us to overload them with __device__ versions (this behavour is only valid for system headers).
+#include <__clang_cuda_math_forward_declares.h>
 #include <cmath>
 #include <cstdlib>
 #include <stdlib.h>
@@ -144,7 +148,6 @@ typedef __SIZE_TYPE__ size_t;
 #if defined(__HIPCC_RTC__)
 #include <__clang_hip_cmath.h>
 #else
-#include <__clang_cuda_math_forward_declares.h>
 #include <__clang_hip_cmath.h>
 #include <__clang_cuda_complex_builtins.h>
 #include <algorithm>
diff --git a/clang/test/Headers/hip-constexpr-cmath.hip b/clang/test/Headers/hip-constexpr-cmath.hip
new file mode 100644
index 0000000000000..8b2c3187bf82c
--- /dev/null
+++ b/clang/test/Headers/hip-constexpr-cmath.hip
@@ -0,0 +1,70 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: rm -rf %t
+// RUN: split-file %s %t
+//
+// Test with the pre-202604 stl
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %t/stl/include \
+// RUN:   -internal-isystem %t/stl/2025/include \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -verify %t/main.hip -fcuda-is-device -o /dev/null
+//
+// Test with the 202604 stl
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %t/stl/include \
+// RUN:   -internal-isystem %t/stl/2026/include \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -verify %t/main.hip -fcuda-is-device -o /dev/null
+
+#--- stl/include/stl_common
+_CLANG_BUILTIN2(isunordered)
+_CLANG_BUILTIN2(isgreater)
+_CLANG_BUILTIN2(isgreaterequal)
+_CLANG_BUILTIN2(isless)
+_CLANG_BUILTIN2(islessequal)
+_CLANG_BUILTIN2(islessgreater)
+
+#define FP_NAN 0
+#define FP_INFINITE 1
+#define FP_ZERO 2
+#define FP_SUBNORMAL 3
+#define FP_NORMAL 4
+#--- stl/2025/include/version
+#define _MSVC_STL_UPDATE 202508L
+#--- stl/2025/include/cmath
+#define _CLANG_BUILTIN2(NAME) \
+    bool NAME(float x, float y) noexcept { \
+        return __builtin_##NAME(x, y); \
+    }
+
+#include <stl_common>
+#--- stl/2026/include/version
+#define _MSVC_STL_UPDATE 202604L
+#--- stl/2026/include/cmath
+#define _CLANG_BUILTIN2(NAME) \
+    bool constexpr NAME(float x, float y) noexcept { \
+        return __builtin_##NAME(x, y); \
+    }
+
+#include <stl_common>
+#--- main.hip
+// expected-no-diagnostics
+
+#include <cmath>
+
+#define TEST_CMP(NAME) \
+    bool __attribute__((device)) test_device_ ## NAME (float x) { \
+        bool b = NAME(x, x); \
+        return b; \
+    }
+
+TEST_CMP(isunordered)
+TEST_CMP(isgreater)
+TEST_CMP(isgreaterequal)
+TEST_CMP(isless)
+TEST_CMP(islessequal)
+TEST_CMP(islessgreater)

llvmorg-github-actions · 2026-06-04T12:23:40Z

@llvm/pr-subscribers-backend-x86

Author: Juan Manuel Martinez Caamaño (jmmartinez)

Changes

In HIP, constexpr functions are treated as both __host__ and __device__.

A new version of the MS STL shipped with the build tools version
14.51.36231 has constexpr definitions for some cmath functions when the
compiler in use is Clang (this gets worse when C++23 is in use).

These definitions conflict with the __device__ declarations we provide
in the header wrappers.

There is a workaround for this: We do not mark constexpr
functions that are defined in a system header as
__host__ and __device__ if there is a previous __device__
declaration.

By moving __clang_cuda_math_forward_declares.h before <cmath> is
included we're able to benefit from this behavior.

This fixes error like this one ggml-org/llama.cpp#22570
even for recent versions of C++.

This patch replaces #200395

Full diff: https://github.com/llvm/llvm-project/pull/201563.diff

2 Files Affected:

(modified) clang/lib/Headers/__clang_hip_runtime_wrapper.h (+4-1)
(added) clang/test/Headers/hip-constexpr-cmath.hip (+70)

diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
index 19ce7a5d2c86b..72494fdda4ec9 100644
--- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h
@@ -111,6 +111,10 @@ __attribute__((weak)) inline __device__ void free(void *__ptr) {
 #endif //__cplusplus
 
 #if !defined(__HIPCC_RTC__)
+// We must include the forward declarations before cmath is included.
+// Otherwise `constexpr` functions in <cmath> would be implicitely __host__ __device__.
+// Declaring the __device__ verison before allows us to overload them with __device__ versions (this behavour is only valid for system headers).
+#include <__clang_cuda_math_forward_declares.h>
 #include <cmath>
 #include <cstdlib>
 #include <stdlib.h>
@@ -144,7 +148,6 @@ typedef __SIZE_TYPE__ size_t;
 #if defined(__HIPCC_RTC__)
 #include <__clang_hip_cmath.h>
 #else
-#include <__clang_cuda_math_forward_declares.h>
 #include <__clang_hip_cmath.h>
 #include <__clang_cuda_complex_builtins.h>
 #include <algorithm>
diff --git a/clang/test/Headers/hip-constexpr-cmath.hip b/clang/test/Headers/hip-constexpr-cmath.hip
new file mode 100644
index 0000000000000..8b2c3187bf82c
--- /dev/null
+++ b/clang/test/Headers/hip-constexpr-cmath.hip
@@ -0,0 +1,70 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: rm -rf %t
+// RUN: split-file %s %t
+//
+// Test with the pre-202604 stl
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %t/stl/include \
+// RUN:   -internal-isystem %t/stl/2025/include \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -verify %t/main.hip -fcuda-is-device -o /dev/null
+//
+// Test with the 202604 stl
+// RUN: %clang_cc1 -include __clang_hip_runtime_wrapper.h \
+// RUN:   -internal-isystem %S/../../lib/Headers/cuda_wrappers \
+// RUN:   -internal-isystem %t/stl/include \
+// RUN:   -internal-isystem %t/stl/2026/include \
+// RUN:   -internal-isystem %S/Inputs/include \
+// RUN:   -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-unknown \
+// RUN:   -target-cpu gfx906 -verify %t/main.hip -fcuda-is-device -o /dev/null
+
+#--- stl/include/stl_common
+_CLANG_BUILTIN2(isunordered)
+_CLANG_BUILTIN2(isgreater)
+_CLANG_BUILTIN2(isgreaterequal)
+_CLANG_BUILTIN2(isless)
+_CLANG_BUILTIN2(islessequal)
+_CLANG_BUILTIN2(islessgreater)
+
+#define FP_NAN 0
+#define FP_INFINITE 1
+#define FP_ZERO 2
+#define FP_SUBNORMAL 3
+#define FP_NORMAL 4
+#--- stl/2025/include/version
+#define _MSVC_STL_UPDATE 202508L
+#--- stl/2025/include/cmath
+#define _CLANG_BUILTIN2(NAME) \
+    bool NAME(float x, float y) noexcept { \
+        return __builtin_##NAME(x, y); \
+    }
+
+#include <stl_common>
+#--- stl/2026/include/version
+#define _MSVC_STL_UPDATE 202604L
+#--- stl/2026/include/cmath
+#define _CLANG_BUILTIN2(NAME) \
+    bool constexpr NAME(float x, float y) noexcept { \
+        return __builtin_##NAME(x, y); \
+    }
+
+#include <stl_common>
+#--- main.hip
+// expected-no-diagnostics
+
+#include <cmath>
+
+#define TEST_CMP(NAME) \
+    bool __attribute__((device)) test_device_ ## NAME (float x) { \
+        bool b = NAME(x, x); \
+        return b; \
+    }
+
+TEST_CMP(isunordered)
+TEST_CMP(isgreater)
+TEST_CMP(isgreaterequal)
+TEST_CMP(isless)
+TEST_CMP(islessequal)
+TEST_CMP(islessgreater)

github-actions · 2026-06-04T12:24:29Z

✅ With the latest revision this PR passed the C/C++ code formatter.

…cmath>` This patch should fix the following error on windows: ggml-org/llama.cpp#22570 In HIP, constexpr functions are treated as both __host__ and __device__. A new version of the MS STL shipped with the build tools version 14.51.36231 has constexpr definitions for some cmath functions when the compiler in use is Clang. These definitions conflict with the __device__ declarations we provide in the header wrappers. There is a workaround for this: It is possible to overload constexpr functions **that are defined in a system header** by declaring a __device__ version before. By moving `__clang_cuda_math_forward_declares.h` before `<cmath>` is included we're able to benefit from this behavour.

jmmartinez · 2026-06-08T15:43:25Z

Ping !

yxsamliu

LGTM. This matches the CUDA wrapper ordering and uses the existing forward declaration mechanism as intended. The exact-signature check keeps the impact limited, so I do not see a blocking issue here.

jmmartinez · 2026-06-09T10:20:09Z

Thanks !

…cmath>` (llvm#201563) In HIP, `constexpr` functions are treated as both `__host__` and `__device__`. A new version of the MS STL shipped with the build tools version 14.51.36231 has `constexpr` definitions for some `cmath` functions when the compiler in use is Clang (this gets worse when C++23 is in use). These definitions conflict with the `__device__` declarations we provide in the header wrappers. There is a workaround for this: We do not mark `constexpr` functions [_that are defined in a system header_](https://github.com/llvm/llvm-project/blob/03127a03860b9d8cb440fe8f51c00647f45eb8be/clang/lib/Sema/SemaCUDA.cpp#L877) as `__host__` and `__device__` if there is a previous `__device__` declaration. By moving `__clang_cuda_math_forward_declares.h` before `<cmath>` is included we're able to benefit from this behavior. This fixes error like this one ggml-org/llama.cpp#22570 even for recent versions of C++. This patch should address C++23 concerns where more functions are marked as `constexpr`. However, we should still figure out how to provide `__device__ constexpr` versions for those functions, or not provide ours and use the system's if they exist. This patch replaces llvm#200395

C++23 (P0533R9) makes the standard math functions constexpr. HIP treats constexpr as implicitly __host__ __device__, so clang's __clang_cuda_math_forward_declares.h can no longer overload them and every HIP translation unit fails with "__device__ function 'isfinite' cannot overload __host__ __device__ function". Fixed upstream by llvm/llvm-project#201563, which includes the forward-declares before <cmath>. The fix is NOT in ROCm 7.14, and verified by reading the shipped header rather than by version arithmetic, it is NOT in ROCm 10.0 either -- that release branch pins an llvm from seven weeks after the fix merged and still has the old order. It IS in ROCm 10.1 nightlies. So projects that must stay on an older ROCm need a workaround, and one that expires cleanly. This module generates a copy of the INSTALLED wrapper with the two includes reordered and puts it ahead of the resource dir on the include path. The language standard is untouched, so host and device translation units keep identical language and standard-library configuration. The measured alternatives do not: -std=c++20 and -D_HAS_CXX23=0 both leave HIP and host code with different STL configurations, an ODR hazard that MSVC cannot even detect_mismatch on, and -D_CONSTEXPR23= does not work at all because the STL redefines it. It is built to retire itself, since a shim that outlived the bug would shadow a good header with a stale copy: generated at configure time and never checked in; empty if a plain probe already compiles; empty if the installed header is already ordered correctly; and re-probed with the shim in place, discarded with a warning if it did not help. It also acts only on this one diagnostic, so unrelated build failures are left alone. Nothing is guarded on Windows. MSVC's STL is merely the first to implement P0533R9; ROCm/llvm-project#285 notes libstdc++ doing so will break Linux identically, so a WIN32 guard or a ROCm version check would miss it. Probe for the capability instead. Promoted straight to the trunk at the maintainer's request rather than waiting for the CubbyFlow and LichtFeld-Studio branches to merge, so the pair of Windows build assets lands as one usable feature instead of half on each branch. This was written with the assistance of an AI coding agent. Test Plan: Verified in four directions on a gfx1151 Windows host, plus a negative control so the pass is not vacuous. Test 2 injects the REAL ROCm 10.1 header, extracted from the shipped rocm-sdk-core wheel, via -resource-dir. ``` # 1. ROCm 7.14 at C++23 -> shim generated # 2. simulated ROCm 10.1 -> "no shim needed", nothing written # 3. ROCm 7.14 at C++20 -> "no shim needed" # 4. real HIP target, C++23 -> builds with the shim applied # 5. control: test 4 with the shim disabled -> fails on the same header cmake -S . -B b -G Ninja \ -DCMAKE_HIP_COMPILER=<rocm>/lib/llvm/bin/amdclang++.exe \ -DCMAKE_HIP_COMPILER_FORCED=TRUE -DCMAKE_HIP_ARCHITECTURES=gfx1151 cmake --build b ``` Pass -G Ninja: with no generator CMake picks Visual Studio, which forces cl.exe and silently ignores the compilers passed above. ``` python3 utils/check.py ```

TheRock ROCm 10.1.0 (published 2026-10-05) becomes the pin, with 10.0.0 kept below it as the fallback. It ships AMD clang 24, whose stock __clang_hip_runtime_wrapper.h includes __clang_cuda_math_forward_declares.h before <cmath> (llvm/llvm-project#201563): the MSVC 14.51 overload clash the patched wrapper worked around is fixed upstream, and the stock header compiles the validation TU and all of ggml-hip clean. 02-build.ps1 now force-includes nothing on clang >= 24; patches\hip\23 stays for the 10.0.0 fallback. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

Clang 24 (TheRock ROCm 10.1.0) fixed the MSVC <cmath> clash upstream (llvm/llvm-project#201563) and the framework installs the newest stable dist, so the clang 23 copy under patches\hip\ had no dist left to serve. 02-build.ps1 and pcie-bench's build.ps1 now fail fast below clang 24 instead of force-including a wrapper; the copy stays in git history. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

[Pre-commit test]

bfd5585

jmmartinez requested a review from yxsamliu June 4, 2026 12:23

jmmartinez self-assigned this Jun 4, 2026

llvmorg-github-actions Bot added clang Clang issues not falling into any other category backend:X86 clang:headers Headers provided by Clang, e.g. for intrinsics labels Jun 4, 2026

jmmartinez mentioned this pull request Jun 4, 2026

[Draft][Clang][HIP] Guard declarations of cmath comparisons when using Microsoft's STL #200395

Closed

jmmartinez force-pushed the users/jmmartinez/fix_hip_build_in_windows_electric_boogaloo branch from 22b18bd to f41c972 Compare June 4, 2026 12:25

This was referenced Jun 4, 2026

Compile bug: Seemingly random failures when building HIP/ROCm Windows ggml-org/llama.cpp#22570

Closed

[DO NOT REVIEW, PSDB ONLY][Clang][HIP/CUDA] Guard declarations of cmath comparisons when using Microsoft's STL ROCm/llvm-project#2726

Closed

yxsamliu approved these changes Jun 8, 2026

View reviewed changes

jmmartinez merged commit 01d3932 into main Jun 9, 2026
10 checks passed

jmmartinez deleted the users/jmmartinez/fix_hip_build_in_windows_electric_boogaloo branch June 9, 2026 10:20

This was referenced Jun 16, 2026

[Issue]: Windows: clang headers conflict with MSVC v145 ROCm/llvm-project#2669

Closed

[Clang][HIP] Fix math header conflict with MSVC v145 #203094

Closed

IIIIIllllIIIIIlllll mentioned this pull request Jul 31, 2026

Support TheRock LLVM layout in Windows HIP build script 0xShug0/audio.cpp#160

Merged

chinmaydd mentioned this pull request Aug 10, 2026

[Issue]: #include <cmath> fail to compile with c++23 when using hip ROCm/llvm-project#285

Closed

This was referenced Sep 23, 2026

Port current SOURCE_CANONICAL snapshot to RX 6750 XT / gfx1031 after publication Sakushey/gfx1030-dlss-nr-research#26

Open

Build reproducibility: preflight the Windows HIP compiler invocation Sakushey/gfx1030-dlss-nr-research#41

Open

BlueKingMuch mentioned this pull request Oct 1, 2026

HIP: build with Visual Studio 2026's C++ library (MSVC 14.51+) Niko1221/Strata#373

Closed

pfultz2 mentioned this pull request Oct 19, 2026

fix: remove redundant VS 2022 install in windows-gpu CI job ROCm/AMDMIGraphX#5411

Closed

2 tasks

Navigation Menu

Repository navigation

[Clang][HIP] Include `__clang_cuda_math_forward_declares.h` before `<cmath>` - #201563

[Clang][HIP] Include `__clang_cuda_math_forward_declares.h` before `<cmath>`#201563
jmmartinez merged 2 commits into
mainfrom
users/jmmartinez/fix_hip_build_in_windows_electric_boogaloo

jmmartinez commented Jun 4, 2026 •

edited

Loading

Uh oh!

llvmorg-github-actions Bot commented Jun 4, 2026

Uh oh!

llvmorg-github-actions Bot commented Jun 4, 2026

Uh oh!

github-actions Bot commented Jun 4, 2026 •

edited

Loading

Uh oh!

jmmartinez commented Jun 8, 2026

Uh oh!

yxsamliu left a comment

Uh oh!

Uh oh!

jmmartinez commented Jun 9, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

jmmartinez commented Jun 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmorg-github-actions Bot commented Jun 4, 2026

Uh oh!

llvmorg-github-actions Bot commented Jun 4, 2026

Uh oh!

github-actions Bot commented Jun 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

jmmartinez commented Jun 8, 2026

Uh oh!

yxsamliu left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

jmmartinez commented Jun 9, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

jmmartinez commented Jun 4, 2026 •

edited

Loading

github-actions Bot commented Jun 4, 2026 •

edited

Loading