diff options
| -rw-r--r-- | lib/Sema/SemaCUDA.cpp | 39 | ||||
| -rw-r--r-- | test/SemaCUDA/default-ctor.cu | 43 |
2 files changed, 70 insertions, 12 deletions
diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp index 203c09c571..e009dcb6f1 100644 --- a/lib/Sema/SemaCUDA.cpp +++ b/lib/Sema/SemaCUDA.cpp @@ -267,6 +267,18 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, 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 HasH = MemberDecl->hasAttr<CUDAHostAttr>(); + bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>(); + bool HasExplicitAttr = + (HasD && !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) || + (HasH && !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit()); + if (!InClass || HasExplicitAttr) + return false; + llvm::Optional<CUDAFunctionTarget> InferredTarget; // We're going to invoke special member lookup; mark that these special @@ -371,21 +383,24 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, } } + + // If no target was inferred, mark this member as __host__ __device__; + // it's the least restrictive option that can be invoked from any target. + bool NeedsH = true, NeedsD = true; 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; } diff --git a/test/SemaCUDA/default-ctor.cu b/test/SemaCUDA/default-ctor.cu new file mode 100644 index 0000000000..cbad7a1774 --- /dev/null +++ b/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; +} |