yaxunl updated this revision to Diff 220907.
yaxunl added a comment.
simplify logic by Artem's comments.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D67509/new/
https://reviews.llvm.org/D67509
Files:
lib/Sema/SemaCUDA.cpp
test/SemaCUDA/default-ctor.cu
Index: test/SemaCUDA/default-ctor.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN: -fcuda-is-device -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN: -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+struct In { In() = default; };
+struct InD { __device__ InD() = default; };
+struct InH { __host__ InH() = default; };
+struct InHD { __host__ __device__ InHD() = default; };
+
+struct Out { Out(); };
+struct OutD { __device__ OutD(); };
+struct OutH { __host__ OutH(); };
+struct OutHD { __host__ __device__ OutHD(); };
+
+Out::Out() = default;
+__device__ OutD::OutD() = default;
+__host__ OutH::OutH() = default;
+__host__ __device__ OutHD::OutHD() = default;
+
+__device__ void fd() {
+ In in;
+ InD ind;
+ InH inh; // expected-error{{no matching constructor for initialization of 'InH'}}
+ InHD inhd;
+ Out out; // expected-error{{no matching constructor for initialization of 'Out'}}
+ OutD outd;
+ OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
+ OutHD outhd;
+}
+
+__host__ void fh() {
+ In in;
+ InD ind; // expected-error{{no matching constructor for initialization of 'InD'}}
+ InH inh;
+ InHD inhd;
+ Out out;
+ OutD outd; // expected-error{{no matching constructor for initialization of 'OutD'}}
+ OutH outh;
+ OutHD outhd;
+}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -267,6 +267,17 @@
CXXMethodDecl *MemberDecl,
bool ConstRHS,
bool Diagnose) {
+ // If the defaulted special member is defined lexically outside of its
+ // owning class, or the special member already has explicit device or host
+ // attributes, do not infer.
+ bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
+ bool HasExpAttr = (MemberDecl->hasAttr<CUDADeviceAttr>() &&
+ !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) ||
+ (MemberDecl->hasAttr<CUDAHostAttr>() &&
+ !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit());
+ if (!InClass || HasExpAttr)
+ return false;
+
llvm::Optional<CUDAFunctionTarget> InferredTarget;
// We're going to invoke special member lookup; mark that these special
@@ -371,21 +382,26 @@
}
}
+ bool NeedsH = true, NeedsD = true;
+ bool HasH = MemberDecl->hasAttr<CUDAHostAttr>();
+ bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>();
+
+ // If no target was inferred, mark this member as __host__ __device__;
+ // it's the least restrictive option that can be invoked from any target.
if (InferredTarget.hasValue()) {
- if (InferredTarget.getValue() == CFT_Device) {
- MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
- } else if (InferredTarget.getValue() == CFT_Host) {
- MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
- } else {
- MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
- MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
- }
- } else {
- // If no target was inferred, mark this member as __host__ __device__;
- // it's the least restrictive option that can be invoked from any target.
+ if (InferredTarget.getValue() == CFT_Device)
+ NeedsH = false;
+ else if (InferredTarget.getValue() == CFT_Host)
+ NeedsD = false;
+ }
+
+ // We either setting attributes first time, or the inferred ones must match
+ // previously set ones.
+ assert(!(HasD || HasH) || (NeedsD == HasD && NeedsH == HasH));
+ if (NeedsD && !HasD)
MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ if (NeedsH && !HasH)
MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
- }
return false;
}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
