[PATCH] D77954: [CUDA][HIP] Fix overload resolution issue for device host functions

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rjmccall.
yaxunl edited the summary of this revision.

Currently clang fails to compile the following CUDA program in device compilation:

  __host__ int foo(int x) {
       return 1;
  }
  
  template<class T>
  __device__ __host__ int foo(T x) {
      return 2;
  }
  
  __device__ __host__ int bar() {
      return foo(1);
  }
  
  __global__ void test(int *a) {
      *a = bar();
  }

This is due to foo is resolved to the `__host__ foo` instead of `__device__ __host__ foo`.
This seems to be a bug since `__device__ __host__ foo` is a viable callee for foo whereas
clang is unable to choose it.

nvcc has similar issue

https://cuda.godbolt.org/z/bGijLc

Although it only emits a warning and does not fail to compile. It emits a trap in the code
so that it will fail at run time.

This patch fixes that.


https://reviews.llvm.org/D77954

Files:
  clang/lib/Sema/SemaOverload.cpp
  clang/test/SemaCUDA/function-overload.cu


Index: clang/test/SemaCUDA/function-overload.cu
===================================================================
--- clang/test/SemaCUDA/function-overload.cu
+++ clang/test/SemaCUDA/function-overload.cu
@@ -331,9 +331,6 @@
 // If we have a mix of HD and H-only or D-only candidates in the overload set,
 // normal C++ overload resolution rules apply first.
 template <typename T> TemplateReturnTy template_vs_hd_function(T arg)
-#ifdef __CUDA_ARCH__
-//expected-note at -2 {{declared here}}
-#endif
 {
   return TemplateReturnTy();
 }
@@ -342,11 +339,13 @@
 }
 
 __host__ __device__ void test_host_device_calls_hd_template() {
-  HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
-  TemplateReturnTy ret2 = template_vs_hd_function(1);
 #ifdef __CUDA_ARCH__
-  // expected-error at -2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}
+  typedef HostDeviceReturnTy ExpectedReturnTy;
+#else
+  typedef TemplateReturnTy ExpectedReturnTy;
 #endif
+  HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);
+  ExpectedReturnTy ret2 = template_vs_hd_function(1);
 }
 
 __host__ void test_host_calls_hd_template() {
Index: clang/lib/Sema/SemaOverload.cpp
===================================================================
--- clang/lib/Sema/SemaOverload.cpp
+++ clang/lib/Sema/SemaOverload.cpp
@@ -9821,8 +9821,10 @@
         llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
           // Check viable function only.
           return Cand->Viable && Cand->Function &&
-                 S.IdentifyCUDAPreference(Caller, Cand->Function) ==
-                     Sema::CFP_SameSide;
+                 (S.IdentifyCUDAPreference(Caller, Cand->Function) ==
+                      Sema::CFP_SameSide ||
+                  S.IdentifyCUDAPreference(Caller, Cand->Function) ==
+                      Sema::CFP_HostDevice);
         });
     if (ContainsSameSideCandidate) {
       auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D77954.256781.patch
Type: text/x-patch
Size: 1985 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200411/0d281f65/attachment.bin>