maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[CUDA][HIP] Fix hostness of defaulted constructor 

Clang does not respect the explicit device host attributes of defaulted special members.
Also clang does not respect the hostness of special members determined by their
first declarations.
Clang also adds duplicate implicit device or host attributes in certain cases.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D67509

llvm-svn: 372394
This commit is contained in:
Yaxun Liu 2019-09-20 14:28:09 +00:00
parent 4896f7243d
commit e5d17c511f
2 changed files with 71 additions and 13 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

41
clang/lib/Sema/SemaCUDA.cpp
View File

@ -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,22 +383,25 @@ 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.
MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
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;
}

43
clang/test/SemaCUDA/default-ctor.cu Normal file
View File

@ -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;
}