r372356 - [CUDA][HIP] Re-apply part of r372318.

[Michael Liao via cfe-commits]

Author: hliao
Date: Thu Sep 19 14:26:18 2019
New Revision: 372356

URL: http://llvm.org/viewvc/llvm-project?rev=372356&view=rev
Log:
[CUDA][HIP] Re-apply part of r372318.

- r372318 causes violation of `use-of-uninitialized-value` detected by
  MemorySanitizer. Once `Viable` field is set to false, `FailureKind`
  needs setting as well as it will be checked during destruction if
  `Viable` is not true.
- Revert the part trying to skip `std::vector` erasing.

Modified:
    cfe/trunk/lib/Sema/SemaOverload.cpp
    cfe/trunk/test/SemaCUDA/function-overload.cu
    cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu

Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=372356&r1=372355&r2=372356&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Thu Sep 19 14:26:18 2019
@@ -9422,13 +9422,15 @@ OverloadCandidateSet::BestViableFunction
     const FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
     bool ContainsSameSideCandidate =
         llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
-          return Cand->Function &&
+          // Check 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 &&
+        // Check viable function only to avoid unnecessary data copying/moving.
+        return Cand->Viable && Cand->Function &&
                S.IdentifyCUDAPreference(Caller, Cand->Function) ==
                    Sema::CFP_WrongSide;
       };

Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=372356&r1=372355&r2=372356&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Thu Sep 19 14:26:18 2019
@@ -402,3 +402,20 @@ __host__ void test_host_template_overloa
 __device__ void test_device_template_overload() {
   template_overload(1); // OK. Attribute-based overloading picks __device__ variant.
 }
+
+// Two classes with `operator-` defined. One of them is device only.
+struct C1;
+struct C2;
+__device__
+int operator-(const C1 &x, const C1 &y);
+int operator-(const C2 &x, const C2 &y);
+
+template <typename T>
+__host__ __device__ 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);
+}

Modified: cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu?rev=372356&r1=372355&r2=372356&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu (original)
+++ cfe/trunk/test/SemaCUDA/implicit-member-target-collision-cxx11.cu Thu Sep 19 14:26:18 2019
@@ -74,11 +74,13 @@ struct B4_with_device_copy_ctor {
 struct C4_with_collision : A4_with_host_copy_ctor, B4_with_device_copy_ctor {
 };
 
-// expected-note at -3 {{copy constructor of 'C4_with_collision' is implicitly deleted because base class 'B4_with_device_copy_ctor' has no copy constructor}}
+// expected-note at -3 {{candidate constructor (the implicit copy constructor) not viable: call to invalid function from __host__ function}}
+// expected-note at -4 {{implicit copy constructor inferred target collision: call to both __host__ and __device__ members}}
+// expected-note at -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'}}
 }
 
 //------------------------------------------------------------------------------