yaxunl updated this revision to Diff 220892.
yaxunl edited the summary of this revision.
yaxunl added a comment.
Skip inferring for explicit host/device attrs only. Adds checks for implicit
device and host attrs and avoid duplicates.
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,20 +382,31 @@
}
}
+ // The same special member may be inferred multiple times. Check to make sure
+ // each inference gets same result and not to add duplicate attributes.
+ auto addBothAttr = [=]() {
+ assert(MemberDecl->hasAttr<CUDAHostAttr>() ==
+ MemberDecl->hasAttr<CUDADeviceAttr>());
+ if (!MemberDecl->hasAttr<CUDADeviceAttr>())
+ MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ if (!MemberDecl->hasAttr<CUDAHostAttr>())
+ MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ };
if (InferredTarget.hasValue()) {
if (InferredTarget.getValue() == CFT_Device) {
- MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+ assert(!MemberDecl->hasAttr<CUDAHostAttr>());
+ if (!MemberDecl->hasAttr<CUDADeviceAttr>())
+ 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));
- }
+ assert(!MemberDecl->hasAttr<CUDADeviceAttr>());
+ if (!MemberDecl->hasAttr<CUDAHostAttr>())
+ MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ } else
+ addBothAttr();
} 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.
- MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
- MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ addBothAttr();
}
return false;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
