[HIP] calling math function with user defined type causes ambiguity

[yxsamliu]

The following code causes ambiguity for HIP but not for CUDA.

#include "hip/hip_runtime.h"
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);
  }
}

https://godbolt.org/z/TdjvxP7q9

The error message is

In file included from /opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_runtime_wrapper.h:143:
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:398:20: error: call to '__test' is ambiguous
  398 |   typedef decltype(__test(declval<_Tp>())) type;
      |                    ^~~~~~
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:405:18: note: in instantiation of template class '__hip::__numeric_type<user_bfloat16>' requested here
  405 |           bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
      |                  ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:441:26: note: in instantiation of default argument for '__promote_imp<user_bfloat16, user_bfloat16, user_bfloat16>' required here
  441 | class __promote : public __promote_imp<_A1, _A2, _A3> {};
      |                          ^~~~~~~~~~~~~~~~~~~~~~~~~~~~
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:554:21: note: in instantiation of template class '__hip::__promote<user_bfloat16, user_bfloat16, user_bfloat16>' requested here
  554 |     typename __hip::__promote<__T1, __T2, __T3>::type>::type
      |                     ^
<source>:14:5: note: while substituting deduced template arguments into function template 'fma' [with __T1 = user_bfloat16, __T2 = user_bfloat16, __T3 = user_bfloat16]
   14 |     fma(a, b, b);
      |     ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:396:17: note: candidate function
  396 |   static double __test(long double);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:394:17: note: candidate function
  394 |   static double __test(double);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:393:17: note: candidate function
  393 |   static double __test(unsigned long long);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:392:17: note: candidate function
  392 |   static double __test(long long);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:391:17: note: candidate function
  391 |   static double __test(unsigned long);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:390:17: note: candidate function
  390 |   static double __test(long);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:389:17: note: candidate function
  389 |   static double __test(unsigned);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:388:17: note: candidate function
  388 |   static double __test(int);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:387:17: note: candidate function
  387 |   static double __test(char);
      |                 ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:386:16: note: candidate function
  386 |   static float __test(float);
      |                ^
/opt/compiler-explorer/clang-assertions-trunk-20240730/lib/clang/20/include/__clang_hip_cmath.h:385:19: note: candidate function
  385 |   static _Float16 __test(_Float16);
      |               

The reason is that __test is not a template member function of __hip::__numeric_type. When ambiguity happens during overloading resolution of fma, the ambiguity error won't be bypassed through SFINAE so that clang can abandon the unsuccessful overloading resolution of the fma template defined in clang wrapper header and take the user defined fma.

This is a design flaw of __hip::__numeric_type and can be fixed by introducing a template __test_impl member function to enable SFINAE.