Author: Yaxun (Sam) Liu Date: 2024-08-02T18:37:22-04:00 New Revision: db1d3b23a37c7a57fa8b9e5bc94e1b22e278d361
URL: https://github.com/llvm/llvm-project/commit/db1d3b23a37c7a57fa8b9e5bc94e1b22e278d361 DIFF: https://github.com/llvm/llvm-project/commit/db1d3b23a37c7a57fa8b9e5bc94e1b22e278d361.diff LOG: [HIP] Fix __clang_hip_cmath.hip for ambiguity (#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 Added: Modified: clang/lib/Headers/__clang_hip_cmath.h clang/test/Headers/__clang_hip_cmath.hip Removed: ################################################################################ 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); + } +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
