https://github.com/jmmartinez updated https://github.com/llvm/llvm-project/pull/200395
From c10c73df5679df339afd9e2e77f3ec2827df0489 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Fri, 29 May 2026 11:57:30 +0200 Subject: [PATCH] [Clang][HIP] Guard declarations of cmath comparisons when using Microsoft's STL 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__ delcarations we provide in the header wrappers. This patch guards these declarations/definitions to use the ones coming from the STL. --- .../__clang_cuda_math_forward_declares.h | 6 ++ .../lib/Headers/__clang_hip_runtime_wrapper.h | 19 +++++ clang/test/Headers/hip-constexpr-cmath.hip | 70 +++++++++++++++++++ 3 files changed, 95 insertions(+) create mode 100644 clang/test/Headers/hip-constexpr-cmath.hip diff --git a/clang/lib/Headers/__clang_cuda_math_forward_declares.h b/clang/lib/Headers/__clang_cuda_math_forward_declares.h index 45fe1e5b1772d..a03e26048e7e8 100644 --- a/clang/lib/Headers/__clang_cuda_math_forward_declares.h +++ b/clang/lib/Headers/__clang_cuda_math_forward_declares.h @@ -94,21 +94,25 @@ __DEVICE__ bool isfinite(long double); #endif __DEVICE__ bool isfinite(double); __DEVICE__ bool isfinite(float); +#ifndef __HAS_CONSTEXPR_MATH_CMP2_FUNCTIONS __DEVICE__ bool isgreater(double, double); __DEVICE__ bool isgreaterequal(double, double); __DEVICE__ bool isgreaterequal(float, float); __DEVICE__ bool isgreater(float, float); +#endif #ifdef _MSC_VER __DEVICE__ bool isinf(long double); #endif __DEVICE__ bool isinf(double); __DEVICE__ bool isinf(float); +#ifndef __HAS_CONSTEXPR_MATH_CMP2_FUNCTIONS __DEVICE__ bool isless(double, double); __DEVICE__ bool islessequal(double, double); __DEVICE__ bool islessequal(float, float); __DEVICE__ bool isless(float, float); __DEVICE__ bool islessgreater(double, double); __DEVICE__ bool islessgreater(float, float); +#endif #ifdef _MSC_VER __DEVICE__ bool isnan(long double); #endif @@ -116,8 +120,10 @@ __DEVICE__ bool isnan(double); __DEVICE__ bool isnan(float); __DEVICE__ bool isnormal(double); __DEVICE__ bool isnormal(float); +#ifndef __HAS_CONSTEXPR_MATH_CMP2_FUNCTIONS __DEVICE__ bool isunordered(double, double); __DEVICE__ bool isunordered(float, float); +#endif __DEVICE__ long labs(long); __DEVICE__ double ldexp(double, int); __DEVICE__ float ldexp(float, int); diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index 19ce7a5d2c86b..b67aa498fc207 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -144,8 +144,27 @@ typedef __SIZE_TYPE__ size_t; #if defined(__HIPCC_RTC__) #include <__clang_hip_cmath.h> #else +// MSVC STL Update 202604L introduces constexpr math comparison functions. On +// HIP, constexpr implies __host__ and __device__. Our __device__ overloads +// clash with these definitions. The functions are: isfinite, isinf, isnan, +// isnormal, isunordered, isgreater, isgreaterequal, isless, islessequal, +// islessgreater. +// +// This assumes that the constexpr version provided in the STL doesn't call +// __host__ helpers. +#if defined(__HIP__) && __has_include(<version>) +// The <version> header is fairly recent. Since the problem is +// only present in recent versions of the STL we can safely ignore the +// workaround for older versions. +#include <version> +#if (defined(_MSVC_STL_UPDATE) && _MSVC_STL_UPDATE >= 202604L) +#define __HAS_CONSTEXPR_MATH_CMP2_FUNCTIONS 1 +#endif +#endif + #include <__clang_cuda_math_forward_declares.h> #include <__clang_hip_cmath.h> +#undef __HAS_CONSTEXPR_MATH_CMP2_FUNCTIONS #include <__clang_cuda_complex_builtins.h> #include <algorithm> #include <complex> 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) _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
