// With -fcuda-host-device-constexpr, an unattributed constexpr function is
// treated as implicitly __host__ __device__, unless:
// * it is a variadic function (device-side variadic functions are not
// allowed), or
// * a __device__ function with this signature was already declared, in which
// case in which case we output an error, unless the __device__ decl is in a
// system header, in which case we leave the constexpr function unattributed.
//
// In addition, all function decls are treated as __host__ __device__ when
// ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
// #pragma clang force_cuda_host_device_begin/end
// pair).
void Sema::maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *NewD,
const LookupResult &Previous) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
if (ForceCUDAHostDeviceDepth > 0) {
if (!NewD->hasAttr<CUDAHostAttr>())
NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
if (!NewD->hasAttr<CUDADeviceAttr>())
NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
return;
}
if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
return;
// Is D a __device__ function with the same signature as NewD, ignoring CUDA
// attributes?
auto IsMatchingDeviceFn = [&](NamedDecl *D) {
if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(D))
D = Using->getTargetDecl();
FunctionDecl *OldD = D->getAsFunction();
return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
!OldD->hasAttr<CUDAHostAttr>() &&
!IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
/* ConsiderCudaAttrs = */ false);
};
auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
if (It != Previous.end()) {
// We found a __device__ function with the same name and signature as NewD
// (ignoring CUDA attrs). This is an error unless that function is defined
// in a system header, in which case we simply return without making NewD
// host+device.
NamedDecl *Match = *It;
if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
Diag(NewD->getLocation(),
diag::err_cuda_unattributed_constexpr_cannot_overload_device)
<< NewD->getName();
Diag(Match->getLocation(),
diag::note_cuda_conflicting_device_function_declared_here);
}
return;
}
NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
}
/// Checks access to all the declarations in the given result set.
void Sema::CheckLookupAccess(const LookupResult &R) {
assert(getLangOptions().AccessControl
&& "performing access check without access control");
assert(R.getNamingClass() && "performing access check without naming class");
for (LookupResult::iterator I = R.begin(), E = R.end(); I != E; ++I) {
if (I.getAccess() != AS_public) {
AccessTarget Entity(Context, AccessedEntity::Member,
R.getNamingClass(), I.getPair(),
R.getBaseObjectType());
Entity.setDiag(diag::err_access);
CheckAccess(*this, R.getNameLoc(), Entity);
}
}
}
