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@-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'}} } //------------------------------------------------------------------------------ _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits