<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/101239>101239</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[HIP] calling math function with user defined type causes ambiguity
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
yxsamliu
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
yxsamliu
</td>
</tr>
</table>
<pre>
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.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMWEmP4zoO_jXKhajAluPYOeSQTnXw-jJoYGowR0O26VjvyZKfllTXHOa3DyQ7e6q6gwamEgRetFD8PlIUTWYM30rEJUm_EErffhjWCe4IpSR9njBnW6WX-9ZJqeq35UuL0Cgh1CuXW6hUjVAxZ9AA60q-ddy-QaM0_PHtO5TOglQ2vK__9byakuiZRKvxOo_G__BKEy4r4WoEQmnLe0I3Le8L7aTlHU5br1QYaax2lQVnUBdlIxSz8RxI9mXoBSiKVhlbFFAUNe54hUVxPpjQPDwRuvDTsuf9TNWjZlZpGLtzPyL5ctVdK1cKvOj3cg7P4SpZh6ZnFQ7LH1_PlL2tIzQdIzSvlDSXUBmha7jVUb7XURG6OF0UQKN1WgI7QXfg4ajaVqiSiaKAneI1WDS2GNS6knehIZDkGeJp1HiNyvBG_VtyMmWQFMCU42VxU53Dw7nDtNb2hiQrQjeEbraqLpWwU6W3hG7-Q-jmpf5z9-N79vfiFJX3XtRaaejQGLZF4Oam9G8SGi4QRp-sodGqA0I3qreEbirV9VygfsIfvVAatW8TTG6fmDGoLVfSPFnt5F9PNKKzKEsiQjeCl_txhG6obxrlE7opitBRnDh98apZ36OetiRZxbPEox03y_9Hj6pjtg2rJ4ucJCsakWQ1EOgfKiYEWAWEZkXh3YPQDPg-Eihn9sZMFjmQbA0A9q3HGhuosRL-mdB8PzX3bTsmSLIuXnqSfB03WPiHiWf-A0EiSb_-1_8-i5dZlHrbeHJ8qEN_5xK4NJZJy5lfCVQDFrteMItQCWbMwJgPcsmKJKuikK5DzatiQLk-D1eeiQw0_u3QWKyhRY17HmZRGngolRJhn13JKlaxlxAW2jHhfHydEzq_MZCeD7zB9afRPIu9-81_RnONDXPCAtNb16Ecjp5Adq9VpywWvOuvCKbrywPisuFoAq4vLTCLAzuDYQ8LgVexd6XgFVyu7k1C1-AJD7eEJF9hPIred_Hbv8-ySJrOvEXi33T8kZjfMsn1rkjT2Rhvhojjj154Z92ieBmsUbzQ8Z4c9sGwNU5f3t0Tydoop6txeOzZSU_JeW39eWJcaSy3zvrsqcbaVVgfOdq7rQEurYLGySrweBhAaBbOzQxI-uWV29arG4eNf0WXB_RuT3KjJz0kQhAf-fvJOf0AwSFZ-LgQZ6dsV0zWvPaU7Uk8nkXzEZuxzPJqzObgcAwJ5U0zpngnUD8d5ewulLOPUT4iwOQugMnHAJ0M3zY1BHv6ywM6Lr0LMf0Fx70C-ukY47swxndY9QENGt0FNvq5QR8PY764B2O--DWDPpTP5vldEPOPIXJpH9CK2V0Qs48hVi3TD4gx5AXzX8V4kReEtOgI8VAyejCM4ftz8asY03OMxWYs2Rxg7luu9uP572bR5KVF0MiMksAN2JZZIPNoFD2PfKNUFtgxn-2wK1EfE13VDDPe-zSeR1P4d4vypNjYsr5HaaB22ufUaodaKFb7Z41GCXci2Oey88hnsbbFExlDSehVSUIzCyVC-dZ7Q9VgW63ctoV_br79Y_UVjBpgBVN4ooGVTNY-RW8RnDSuqtCYxokPFPFDD8ocuaix4f5k43IUP1Z_oEVWowYma7DsLxyWMqgPMw7Cpue1Lm5CNQZq9DEWGsFe_fo_ozis5LGVCA3_gTWUb_6DRKvaVR7OiQEP9vWftsJPvjSpVYCS-WgxcDid1MukXiQLNsFlnFGaZDSN4km7nOcxpZhjs0jjiOZVFcVxnVZ5nMdsEZfphC_3m4VGs5zGUyyjdN40SCmLkrJqyCzCjnExFWLXTZXeTrgxDpdxFNNkMRGsRGHGYrfEVwi9Y7VbL_2kp9JtDZlFghtrjmIstyJUyf_49p2kz6Hy5Znw-_AINXyQnVnGM3pVH584LZYX5UtuW1dOK9X5MCB2-9tTr9WfWPmQEnQ1hG5GMLsl_V8AAAD__96Hauw">