[clang] [HIP] Fix __clang_hip_cmath.hip for ambiguity (PR #101341)

[Yaxun Liu via cfe-commits]

https://github.com/yxsamliu created https://github.com/llvm/llvm-project/pull/101341

If there is a type T which can be converted to both float and double etc but itself is not specialized for __numeric_type, and it is called for math functions eg. fma, it will cause ambiguity with test function of __numeric_type.

Since test is not template, this error is not bypassed by SFINAE. This is a design flaw of __numeric_type. This patch fixes clang wrapper header to use SFINAE to avoid such ambiguity.

Fixes: SWDEV-461604

Fixes: https://github.com/llvm/llvm-project/issues/101239

>From 07d6d9392d3ec9f6054f988252da75203b346fd4 Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Tue, 30 Jul 2024 16:51:23 -0400
Subject: [PATCH] [HIP] Fix __clang_hip_cmath.hip for ambiguity

If there is a type T which can be converted to both float and double etc but itself
is not specialized for __numeric_type, and it is called for math functions eg. fma,
it will cause ambiguity with test function of __numeric_type.

Since test is not template, this error is not bypassed by SFINAE. This is a design
flaw of __numeric_type. This patch fixes clang wrapper header to use SFINAE to
avoid such ambiguity.

Fixes: SWDEV-461604

Fixes: https://github.com/llvm/llvm-project/issues/101239
---
 clang/lib/Headers/__clang_hip_cmath.h    |  7 ++++++-
 clang/test/Headers/__clang_hip_cmath.hip | 19 +++++++++++++++++++
 2 files changed, 25 insertions(+), 1 deletion(-)

diff --git a/clang/lib/Headers/__clang_hip_cmath.h b/clang/lib/Headers/__clang_hip_cmath.h
index b52d6b7816611..7d982ad9af7ee 100644
--- a/clang/lib/Headers/__clang_hip_cmath.h
+++ b/clang/lib/Headers/__clang_hip_cmath.h
@@ -395,7 +395,12 @@ template <class _Tp> struct __numeric_type {
   // No support for long double, use double instead.
   static double __test(long double);
 
-  typedef decltype(__test(declval<_Tp>())) type;
+  template <typename _U>
+  static auto __test_impl(int) -> decltype(__test(declval<_U>()));
+
+  template <typename _U> static void __test_impl(...);
+
+  typedef decltype(__test_impl<_Tp>(0)) type;
   static const bool value = !is_same<type, void>::value;
 };
 
diff --git a/clang/test/Headers/__clang_hip_cmath.hip b/clang/test/Headers/__clang_hip_cmath.hip
index ed1030b820627..0c9ff4cdd7808 100644
--- a/clang/test/Headers/__clang_hip_cmath.hip
+++ b/clang/test/Headers/__clang_hip_cmath.hip
@@ -87,3 +87,22 @@ extern "C" __device__ float test_sin_f32(float x) {
 extern "C" __device__ float test_cos_f32(float x) {
   return cos(x);
 }
+
+// Check user defined type which can be converted to float and double but not
+// specializes __numeric_type will not cause ambiguity diagnostics.
+struct user_bfloat16 {
+  __host__ __device__ user_bfloat16(float);
+  operator float();
+  operator double();
+};
+
+namespace user_namespace {
+  __device__ user_bfloat16 fma(const user_bfloat16 a, const user_bfloat16 b, const user_bfloat16 c) {
+    return a;
+  }
+
+  __global__ void test_fma() {
+    user_bfloat16 a = 1.0f, b = 2.0f;
+    fma(a, b, b);
+  }
+}