[CUDA][HIP] Fix typo in `BestViableFunction`

[hliao]

Summary

• Should consider viable ones only when checking SameSide candidates.
• Replace erasing with clearing viable flag to reduce data moving/copying.
• Add one and revise another one as the diagnostic message are more relevant compared to previous one.

[tra]

LGTM. You may want to wait a bit for Justin's feedback, in case he has some concerns.

clang/test/SemaCUDA/function-overload.cu

406

I'd remove irrelevant.

407–411

We don't need definitions for struct and operator-;

This should do:

struct C1;
struct C2;
__device__ int operator-(const C1 &x, const C1 &y);
int operator-(const C2 &x, const C2 &y);
...
int test_constexpr_overload(C2 &x, C2 &y) {
    return constexpr_overload(x, y);
}

414

Is constexpr necessary here? If not, then we would not need -std=c++11 either.
I think you really want __host__ __device__ here instead. constexpr is just applying the attributes implicitly.

[hliao]

r372318 with test case revised following suggestion.

clang/lib/Sema/SemaOverload.cpp

OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
   // exclude non-viable overload candidates from consideration based
   // only on their host/device attributes. Specifically, if one
   // candidate call is WrongSide and the other is SameSide, we ignore
   // the WrongSide candidate.
   if (S.getLangOpts().CUDA) {
     const FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
     bool ContainsSameSideCandidate =
         llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
-          return Cand->Function &&
+          // Consider viable function only.
+          return Cand->Viable && Cand->Function &&
                  S.IdentifyCUDAPreference(Caller, Cand->Function) ==
                      Sema::CFP_SameSide;
         });
     if (ContainsSameSideCandidate) {
-      auto IsWrongSideCandidate = [&](OverloadCandidate *Cand) {
-        return Cand->Function &&
-               S.IdentifyCUDAPreference(Caller, Cand->Function) ==
-                   Sema::CFP_WrongSide;
-      };
-      llvm::erase_if(Candidates, IsWrongSideCandidate);
+      // Clear viable flag for WrongSide varible candidates.
+      llvm::for_each(Candidates, [&](OverloadCandidate *Cand) {
+        if (Cand->Viable && Cand->Function &&
+                Sema::CFP_WrongSide)
+          Cand->Viable = false;
+      });
     }
   }
 
   // Find the best viable function.
   Best = end();
   for (auto *Cand : Candidates)
     if (Cand->Viable)
       if (Best == end() ||

clang/test/SemaCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
-// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
 // Opaque return types used to check that we pick the right overloads.
 struct HostReturnTy {};
 struct HostReturnTy2 {};
 struct DeviceReturnTy {};
 struct DeviceReturnTy2 {};

 template <typename T> __device__ T template_overload(const T &a) { return a; };
 
 __host__ void test_host_template_overload() {
   template_overload(1); // OK. Attribute-based overloading picks __host__ variant.
 }
 __device__ void test_device_template_overload() {
   template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
 }
+
+// Two irrelevant classes with `operator-` defined. One of them is device only.
+struct C1 { int m; };
+struct C2 { int *m; };
+__device__
+int operator-(const C1 &x, const C1 &y) { return x.m - y.m; }
+int operator-(const C2 &x, const C2 &y) { return x.m - y.m; }
+
+template <typename T>
+constexpr int constexpr_overload(const T &x, const T &y) {
+  return x - y;
+}
+
+// Verify that function overloading doesn't prune candidate wrongly.
+int test_constexpr_overload(C2 x, C2 y) {
+  return constexpr_overload(x, y);
+}

clang/test/SemaCUDA/implicit-member-target-collision-cxx11.cu

 struct B4_with_device_copy_ctor {
   B4_with_device_copy_ctor() {}
   __device__ B4_with_device_copy_ctor(const B4_with_device_copy_ctor&) {}
 };
 
 struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
 };
 
-// expected-note@-3 {{copy constructor of 'C4_with_collision' is implicitly deleted because base class 'B4_with_device_copy_ctor' has no copy constructor}}
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable: call to invalid function from __host__ function}}
+// expected-note@-4 {{implicit copy constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note@-5 {{candidate constructor (the implicit default constructor) not viable: requires 0 arguments, but 1 was provided}}
 
 void hostfoo4() {
   C4_with_collision c;
-  C4_with_collision c2 = c; // expected-error {{call to implicitly-deleted copy constructor of 'C4_with_collision'}}
+  C4_with_collision c2 = c; // expected-error {{no matching constructor for initialization of 'C4_with_collision'}}
 }
 
 //------------------------------------------------------------------------------
 // Test 5: collision on resolving a move ctor
 
 struct A5_with_host_move_ctor {
   A5_with_host_move_ctor() {}
   A5_with_host_move_ctor(A5_with_host_move_ctor&&) {}