https://github.com/jmmartinez updated https://github.com/llvm/llvm-project/pull/201563
From bfd55859d5f1473f93dc2411f216d79e1cc4e4b4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Thu, 4 Jun 2026 14:01:45 +0200 Subject: [PATCH 1/2] [Pre-commit test] --- clang/test/Headers/hip-constexpr-cmath.hip | 70 ++++++++++++++++++++++ 1 file changed, 70 insertions(+) create mode 100644 clang/test/Headers/hip-constexpr-cmath.hip diff --git a/clang/test/Headers/hip-constexpr-cmath.hip b/clang/test/Headers/hip-constexpr-cmath.hip new file mode 100644 index 0000000000000..053f2d487b32f --- /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: not %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) From f41c97219ee234a19cfe08317ac96e736de35422 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Thu, 4 Jun 2026 14:06:41 +0200 Subject: [PATCH 2/2] [Clang][HIP] Include `__clang_cuda_math_forward_declares.h` before `<cmath>` This patch should fix the following error on windows: https://github.com/ggml-org/llama.cpp/issues/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. --- clang/lib/Headers/__clang_hip_runtime_wrapper.h | 7 ++++++- clang/test/Headers/hip-constexpr-cmath.hip | 2 +- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index 19ce7a5d2c86b..4b8cffd86f044 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -111,6 +111,12 @@ __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 +150,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 index 053f2d487b32f..8b2c3187bf82c 100644 --- a/clang/test/Headers/hip-constexpr-cmath.hip +++ b/clang/test/Headers/hip-constexpr-cmath.hip @@ -12,7 +12,7 @@ // RUN: -target-cpu gfx906 -verify %t/main.hip -fcuda-is-device -o /dev/null // // Test with the 202604 stl -// RUN: not %clang_cc1 -include __clang_hip_runtime_wrapper.h \ +// 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 \ _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
