maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/lib/Sema/SemaCUDA.cpp

281 lines
10 KiB
C++
Raw Normal View History

Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
//===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
/// \file
/// \brief This file implements semantic analysis for CUDA constructs.
///
//===----------------------------------------------------------------------===//
#include "clang/Sema/Sema.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Decl.h"
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
#include "clang/Lex/Preprocessor.h"
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
#include "clang/Sema/SemaDiagnostic.h"
CUDA: Fix incorrect target inference for implicit members. As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
2014-09-30 04:38:29 +08:00
#include "llvm/ADT/Optional.h"
#include "llvm/ADT/SmallVector.h"
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
using namespace clang;
ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
MultiExprArg ExecConfig,
SourceLocation GGGLoc) {
FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
if (!ConfigDecl)
return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
<< "cudaConfigureCall");
QualType ConfigQTy = ConfigDecl->getType();
DeclRefExpr *ConfigDR = new (Context)
DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
MarkFunctionReferenced(LLLLoc, ConfigDecl);
return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
/*IsExecConfig=*/true);
}
/// IdentifyCUDATarget - Determine the CUDA compilation target for this function
Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
CUDA: Fix incorrect target inference for implicit members. As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
2014-09-30 04:38:29 +08:00
if (D->hasAttr<CUDAInvalidTargetAttr>())
return CFT_InvalidTarget;
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
if (D->hasAttr<CUDAGlobalAttr>())
return CFT_Global;
if (D->hasAttr<CUDADeviceAttr>()) {
if (D->hasAttr<CUDAHostAttr>())
return CFT_HostDevice;
return CFT_Device;
CUDA: mark the target of implicit intrinsics properly r218624 implemented target inference for implicit special members. However, other entities can be implicit - for example intrinsics. These can not have inference running on them, so they should be marked host device as before. This is the safest and most flexible setting, since by construction these functions don't invoke anything, and we'd like them to be invokable from both host and device code. LLVM's intrinsics definitions (where these intrinsics come from in the case of CUDA/NVPTX) have no notion of target, so both host and device intrinsics can be supported this way. llvm-svn: 218688
2014-10-01 01:38:34 +08:00
} else if (D->hasAttr<CUDAHostAttr>()) {
return CFT_Host;
} else if (D->isImplicit()) {
// Some implicit declarations (like intrinsic functions) are not marked.
// Set the most lenient target on them for maximal flexibility.
return CFT_HostDevice;
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
}
return CFT_Host;
}
CUDA: Fix incorrect target inference for implicit members. As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
2014-09-30 04:38:29 +08:00
bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
const FunctionDecl *Callee) {
Create a frontend flag to disable CUDA cross-target call checks For CUDA source, Sema checks that the targets of call expressions make sense (e.g. a host function can't call a device function). Adding a flag that lets us skip this check. Motivation: for source-to-source translation tools that have to accept code that's not strictly kosher CUDA but is still accepted by nvcc. The source-to-source translation tool can then fix the code and leave calls that are semantically valid for the actual compilation stage. Differential Revision: http://reviews.llvm.org/D9036 llvm-svn: 235049
2015-04-16 06:27:06 +08:00
// The CUDADisableTargetCallChecks short-circuits this check: we assume all
// cross-target calls are valid.
if (getLangOpts().CUDADisableTargetCallChecks)
return false;
Consider calls from implict host device functions as valid in SemaCUDA. In SemaCUDA all implicit functions were considered host device, this led to errors such as the following code snippet failing to compile: struct Copyable { const Copyable& operator=(const Copyable& x) { return *this; } }; struct Simple { Copyable b; }; void foo() { Simple a, b; a = b; } Above the implicit copy assignment operator was inferred as host device but there was only a host assignment copy defined which is an error in device compilation mode. Differential Revision: http://reviews.llvm.org/D6565 llvm-svn: 224358
2014-12-17 04:12:38 +08:00
CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
CalleeTarget = IdentifyCUDATarget(Callee);
CUDA: Fix incorrect target inference for implicit members. As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
2014-09-30 04:38:29 +08:00
// If one of the targets is invalid, the check always fails, no matter what
// the other target is.
if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
return true;
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
// CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
// Callable from the device only."
if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
return true;
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
// CUDA B.1.2 "The __global__ qualifier declares a function that is [...]
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
// Callable from the host only."
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
// CUDA B.1.3 "The __host__ qualifier declares a function that is [...]
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
// Callable from the host only."
if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
(CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
return true;
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
// CUDA B.1.3 "The __device__ and __host__ qualifiers can be used together
// however, in which case the function is compiled for both the host and the
// device. The __CUDA_ARCH__ macro [...] can be used to differentiate code
// paths between host and device."
if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
Consider calls from implict host device functions as valid in SemaCUDA. In SemaCUDA all implicit functions were considered host device, this led to errors such as the following code snippet failing to compile: struct Copyable { const Copyable& operator=(const Copyable& x) { return *this; } }; struct Simple { Copyable b; }; void foo() { Simple a, b; a = b; } Above the implicit copy assignment operator was inferred as host device but there was only a host assignment copy defined which is an error in device compilation mode. Differential Revision: http://reviews.llvm.org/D6565 llvm-svn: 224358
2014-12-17 04:12:38 +08:00
// If the caller is implicit then the check always passes.
if (Caller->isImplicit()) return false;
bool InDeviceMode = getLangOpts().CUDAIsDevice;
CUDA: Add option to allow host device functions to call host functions Commiting code from review http://reviews.llvm.org/D7841 llvm-svn: 230385
2015-02-25 05:45:33 +08:00
if (!InDeviceMode && CalleeTarget != CFT_Host)
return true;
if (InDeviceMode && CalleeTarget != CFT_Device) {
// Allow host device functions to call host functions if explicitly
// requested.
if (CalleeTarget == CFT_Host &&
getLangOpts().CUDAAllowHostCallsFromHostDevice) {
Diag(Caller->getLocation(),
diag::warn_host_calls_from_host_device)
<< Callee->getNameAsString() << Caller->getNameAsString();
return false;
}
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
return true;
CUDA: Add option to allow host device functions to call host functions Commiting code from review http://reviews.llvm.org/D7841 llvm-svn: 230385
2015-02-25 05:45:33 +08:00
}
CUDA host device code with two code paths Summary: Allow CUDA host device functions with two code paths using __CUDA_ARCH__ to differentiate between code path being compiled. For example: __host__ __device__ void host_device_function(void) { #ifdef __CUDA_ARCH__ device_only_function(); #else host_only_function(); #endif } Patch by Jacques Pienaar. Reviewed By: rnk Differential Revision: http://reviews.llvm.org/D6457 llvm-svn: 223271
2014-12-04 05:53:36 +08:00
}
Split off CUDA-specific Sema parts to a new file In line with SemaOpenMP.cpp, etc. CUDA-specific semantic analysis code goes into a separate file. This is in anticipation of adding extra functionality here in the near future. No change in functionality. Review: http://reviews.llvm.org/D5160 llvm-svn: 217043
2014-09-03 23:27:03 +08:00
return false;
}
CUDA: Fix incorrect target inference for implicit members. As PR20495 demonstrates, Clang currenlty infers the CUDA target (host/device, etc) for implicit members (constructors, etc.) incorrectly. This causes errors and even assertions in Clang when compiling code (assertions in C++11 mode where implicit move constructors are added into the mix). Fix the problem by inferring the target from the methods the implicit member should call (depending on its base classes and fields). llvm-svn: 218624
2014-09-30 04:38:29 +08:00
/// When an implicitly-declared special member has to invoke more than one
/// base/field special member, conflicts may occur in the targets of these
/// members. For example, if one base's member __host__ and another's is
/// __device__, it's a conflict.
/// This function figures out if the given targets \param Target1 and
/// \param Target2 conflict, and if they do not it fills in
/// \param ResolvedTarget with a target that resolves for both calls.
/// \return true if there's a conflict, false otherwise.
static bool
resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
Sema::CUDAFunctionTarget Target2,
Sema::CUDAFunctionTarget *ResolvedTarget) {
if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
// TODO: this shouldn't happen, really. Methods cannot be marked __global__.
// Clang should detect this earlier and produce an error. Then this
// condition can be changed to an assertion.
return true;
}
if (Target1 == Sema::CFT_HostDevice) {
*ResolvedTarget = Target2;
} else if (Target2 == Sema::CFT_HostDevice) {
*ResolvedTarget = Target1;
} else if (Target1 != Target2) {
return true;
} else {
*ResolvedTarget = Target1;
}
return false;
}
bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
CXXSpecialMember CSM,
CXXMethodDecl *MemberDecl,
bool ConstRHS,
bool Diagnose) {
llvm::Optional<CUDAFunctionTarget> InferredTarget;
// We're going to invoke special member lookup; mark that these special
// members are called from this one, and not from its caller.
ContextRAII MethodContext(*this, MemberDecl);
// Look for special members in base classes that should be invoked from here.
// Infer the target of this member base on the ones it should call.
// Skip direct and indirect virtual bases for abstract classes.
llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
for (const auto &B : ClassDecl->bases()) {
if (!B.isVirtual()) {
Bases.push_back(&B);
}
}
if (!ClassDecl->isAbstract()) {
for (const auto &VB : ClassDecl->vbases()) {
Bases.push_back(&VB);
}
}
for (const auto *B : Bases) {
const RecordType *BaseType = B->getType()->getAs<RecordType>();
if (!BaseType) {
continue;
}
CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
Sema::SpecialMemberOverloadResult *SMOR =
LookupSpecialMember(BaseClassDecl, CSM,
/* ConstArg */ ConstRHS,
/* VolatileArg */ false,
/* RValueThis */ false,
/* ConstThis */ false,
/* VolatileThis */ false);
if (!SMOR || !SMOR->getMethod()) {
continue;
}
CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
if (!InferredTarget.hasValue()) {
InferredTarget = BaseMethodTarget;
} else {
bool ResolutionError = resolveCalleeCUDATargetConflict(
InferredTarget.getValue(), BaseMethodTarget,
InferredTarget.getPointer());
if (ResolutionError) {
if (Diagnose) {
Diag(ClassDecl->getLocation(),
diag::note_implicit_member_target_infer_collision)
<< (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
}
MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
return true;
}
}
}
// Same as for bases, but now for special members of fields.
for (const auto *F : ClassDecl->fields()) {
if (F->isInvalidDecl()) {
continue;
}
const RecordType *FieldType =
Context.getBaseElementType(F->getType())->getAs<RecordType>();
if (!FieldType) {
continue;
}
CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
Sema::SpecialMemberOverloadResult *SMOR =
LookupSpecialMember(FieldRecDecl, CSM,
/* ConstArg */ ConstRHS && !F->isMutable(),
/* VolatileArg */ false,
/* RValueThis */ false,
/* ConstThis */ false,
/* VolatileThis */ false);
if (!SMOR || !SMOR->getMethod()) {
continue;
}
CUDAFunctionTarget FieldMethodTarget =
IdentifyCUDATarget(SMOR->getMethod());
if (!InferredTarget.hasValue()) {
InferredTarget = FieldMethodTarget;
} else {
bool ResolutionError = resolveCalleeCUDATargetConflict(
InferredTarget.getValue(), FieldMethodTarget,
InferredTarget.getPointer());
if (ResolutionError) {
if (Diagnose) {
Diag(ClassDecl->getLocation(),
diag::note_implicit_member_target_infer_collision)
<< (unsigned)CSM << InferredTarget.getValue()
<< FieldMethodTarget;
}
MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
return 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));
}
return false;
}