summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorMichael Liao <michael.hliao@gmail.com>2019-09-19 13:14:03 +0000
committerMichael Liao <michael.hliao@gmail.com>2019-09-19 13:14:03 +0000
commitf2fcfddc47faf4ca932ace58ca0df87ca4ef0310 (patch)
treed78bc96c4e478507dd76550540315848d38a30e2
parent23d1a459594f2fcf33a00b355fec6a75b349f38d (diff)
downloadclang-f2fcfddc47faf4ca932ace58ca0df87ca4ef0310.tar.gz
[CUDA][HIP] Fix typo in `BestViableFunction`
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. Reviewers: tra Subscribers: cfe-commits, yaxunl Tags: #clang Differential Revision: https://reviews.llvm.org/D67730 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@372318 91177308-0d34-0410-b5e6-96231b3b80d8
-rw-r--r--lib/Sema/SemaOverload.cpp16
-rw-r--r--test/SemaCUDA/function-overload.cu17
-rw-r--r--test/SemaCUDA/implicit-member-target-collision-cxx11.cu6
3 files changed, 30 insertions, 9 deletions
diff --git a/lib/Sema/SemaOverload.cpp b/lib/Sema/SemaOverload.cpp
index 9cf778f934..3e07bd2181 100644
--- a/lib/Sema/SemaOverload.cpp
+++ b/lib/Sema/SemaOverload.cpp
@@ -9422,17 +9422,19 @@ OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
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 &&
+ S.IdentifyCUDAPreference(Caller, Cand->Function) ==
+ Sema::CFP_WrongSide)
+ Cand->Viable = false;
+ });
}
}
diff --git a/test/SemaCUDA/function-overload.cu b/test/SemaCUDA/function-overload.cu
index 1d78636a70..b9efd1c09e 100644
--- a/test/SemaCUDA/function-overload.cu
+++ b/test/SemaCUDA/function-overload.cu
@@ -402,3 +402,20 @@ __host__ void test_host_template_overload() {
__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);
+}
diff --git a/test/SemaCUDA/implicit-member-target-collision-cxx11.cu b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
index 7aa1dd3f20..06015ed0d6 100644
--- a/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
+++ b/test/SemaCUDA/implicit-member-target-collision-cxx11.cu
@@ -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'}}
}
//------------------------------------------------------------------------------