DeclContext::lookup_result
DeclContext::lookup(DeclarationName Name) {
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this)
return PrimaryContext->lookup(Name);
if (hasExternalVisibleStorage())
LoadVisibleDeclsFromExternalStorage();
/// If there is no lookup data structure, build one now by walking
/// all of the linked DeclContexts (in declaration order!) and
/// inserting their values.
if (!LookupPtr) {
buildLookup(this);
if (!LookupPtr)
return lookup_result(0, 0);
}
StoredDeclsMap *Map = static_cast<StoredDeclsMap*>(LookupPtr);
StoredDeclsMap::iterator Pos = Map->find(Name);
if (Pos == Map->end())
return lookup_result(0, 0);
return Pos->second.getLookupResult(getParentASTContext());
}
/// DeclareGlobalAllocationFunction - Declares a single implicit global
/// allocation function if it doesn't already exist.
void Sema::DeclareGlobalAllocationFunction(DeclarationName Name,
QualType Return, QualType Argument)
{
DeclContext *GlobalCtx = Context.getTranslationUnitDecl();
// Check if this function is already declared.
{
DeclContext::lookup_iterator Alloc, AllocEnd;
for (llvm::tie(Alloc, AllocEnd) = GlobalCtx->lookup(Context, Name);
Alloc != AllocEnd; ++Alloc) {
// FIXME: Do we need to check for default arguments here?
FunctionDecl *Func = cast<FunctionDecl>(*Alloc);
if (Func->getNumParams() == 1 &&
Context.getCanonicalType(Func->getParamDecl(0)->getType())==Argument)
return;
}
}
QualType FnType = Context.getFunctionType(Return, &Argument, 1, false, 0);
FunctionDecl *Alloc =
FunctionDecl::Create(Context, GlobalCtx, SourceLocation(), Name,
FnType, FunctionDecl::None, false, true,
SourceLocation());
Alloc->setImplicit();
ParmVarDecl *Param = ParmVarDecl::Create(Context, Alloc, SourceLocation(),
0, Argument, VarDecl::None, 0);
Alloc->setParams(Context, &Param, 1);
// FIXME: Also add this declaration to the IdentifierResolver, but
// make sure it is at the end of the chain to coincide with the
// global scope.
((DeclContext *)TUScope->getEntity())->addDecl(Context, Alloc);
}
DeclContext::lookup_result
DeclContext::lookup(DeclarationName Name) {
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this)
return PrimaryContext->lookup(Name);
if (hasExternalVisibleStorage()) {
// Check to see if we've already cached the lookup results.
if (LookupPtr) {
StoredDeclsMap::iterator I = LookupPtr->find(Name);
if (I != LookupPtr->end())
return I->second.getLookupResult(getParentASTContext());
}
ExternalASTSource *Source = getParentASTContext().getExternalSource();
return Source->FindExternalVisibleDeclsByName(this, Name);
}
/// If there is no lookup data structure, build one now by walking
/// all of the linked DeclContexts (in declaration order!) and
/// inserting their values.
if (!LookupPtr) {
buildLookup(this);
if (!LookupPtr)
return lookup_result(lookup_iterator(0), lookup_iterator(0));
}
StoredDeclsMap::iterator Pos = LookupPtr->find(Name);
if (Pos == LookupPtr->end())
return lookup_result(lookup_iterator(0), lookup_iterator(0));
return Pos->second.getLookupResult(getParentASTContext());
}
DeclContext::lookup_result
DeclContext::lookup(DeclarationName Name) {
assert(DeclKind != Decl::LinkageSpec &&
"Should not perform lookups into linkage specs!");
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this)
return PrimaryContext->lookup(Name);
if (hasExternalVisibleStorage()) {
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
else if (NeedToReconcileExternalVisibleStorage)
reconcileExternalVisibleStorage();
if (!Map)
Map = CreateStoredDeclsMap(getParentASTContext());
// If a PCH/module has a result for this name, and we have a local
// declaration, we will have imported the PCH/module result when adding the
// local declaration or when reconciling the module.
std::pair<StoredDeclsMap::iterator, bool> R =
Map->insert(std::make_pair(Name, StoredDeclsList()));
if (!R.second)
return R.first->second.getLookupResult();
ExternalASTSource *Source = getParentASTContext().getExternalSource();
if (Source->FindExternalVisibleDeclsByName(this, Name)) {
if (StoredDeclsMap *Map = LookupPtr.getPointer()) {
StoredDeclsMap::iterator I = Map->find(Name);
if (I != Map->end())
return I->second.getLookupResult();
}
}
return lookup_result(lookup_iterator(0), lookup_iterator(0));
}
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
if (!Map)
return lookup_result(lookup_iterator(0), lookup_iterator(0));
StoredDeclsMap::iterator I = Map->find(Name);
if (I == Map->end())
return lookup_result(lookup_iterator(0), lookup_iterator(0));
return I->second.getLookupResult();
}
///\brief This is the most important function of the class ASTImportSource
/// since from here initiates the lookup and import part of the missing
/// Decl(s) (Contexts).
///
bool ASTImportSource::FindExternalVisibleDeclsByName(
const DeclContext *childCurrentDeclContext, DeclarationName childDeclName) {
assert(childCurrentDeclContext->hasExternalVisibleStorage() &&
"DeclContext has no visible decls in storage");
//Check if we have already found this declaration Name before
DeclarationName parentDeclName;
std::map<clang::DeclarationName,
clang::DeclarationName>::iterator II = m_DeclName_map.find(childDeclName);
if (II != m_DeclName_map.end()) {
parentDeclName = II->second;
} else {
// Get the identifier info from the parent interpreter
// for this Name.
llvm::StringRef name(childDeclName.getAsString());
IdentifierTable &parentIdentifierTable =
m_parent_Interp->getCI()->getASTContext().Idents;
IdentifierInfo &parentIdentifierInfo = parentIdentifierTable.get(name);
DeclarationName parentDeclNameTemp(&parentIdentifierInfo);
parentDeclName = parentDeclNameTemp;
}
// Search in the map of the stored Decl Contexts for this
// Decl Context.
std::map<const clang::DeclContext *, clang::DeclContext *>::iterator I;
if ((I = m_DeclContexts_map.find(childCurrentDeclContext))
!= m_DeclContexts_map.end()) {
// If childCurrentDeclContext was found before and is already in the map,
// then do the lookup using the stored pointer.
DeclContext *parentDeclContext = I->second;
Decl *fromDeclContext = Decl::castFromDeclContext(parentDeclContext);
ASTContext &from_ASTContext = fromDeclContext->getASTContext();
Decl *toDeclContext = Decl::castFromDeclContext(childCurrentDeclContext);
ASTContext &to_ASTContext = toDeclContext->getASTContext();
DeclContext::lookup_result lookup_result =
parentDeclContext->lookup(parentDeclName);
// Check if we found this Name in the parent interpreter
if (!lookup_result.empty()) {
// Do the import
if (Import(lookup_result, from_ASTContext, to_ASTContext,
childCurrentDeclContext, childDeclName, parentDeclName))
return true;
}
}
return false;
}
DeclContext::lookup_result
DeclContext::lookup(DeclarationName Name) {
assert(DeclKind != Decl::LinkageSpec &&
"Should not perform lookups into linkage specs!");
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this)
return PrimaryContext->lookup(Name);
if (hasExternalVisibleStorage()) {
// If a PCH has a result for this name, and we have a local declaration, we
// will have imported the PCH result when adding the local declaration.
// FIXME: For modules, we could have had more declarations added by module
// imoprts since we saw the declaration of the local name.
if (StoredDeclsMap *Map = LookupPtr.getPointer()) {
StoredDeclsMap::iterator I = Map->find(Name);
if (I != Map->end())
return I->second.getLookupResult();
}
ExternalASTSource *Source = getParentASTContext().getExternalSource();
return Source->FindExternalVisibleDeclsByName(this, Name);
}
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
if (!Map)
return lookup_result(lookup_iterator(0), lookup_iterator(0));
StoredDeclsMap::iterator I = Map->find(Name);
if (I == Map->end())
return lookup_result(lookup_iterator(0), lookup_iterator(0));
return I->second.getLookupResult();
}
DeclContext::lookup_result
DeclContext::lookup(IdentifierInfo &Name) {
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this) // FIXME: not needed?
return PrimaryContext->lookup(Name);
#if 0 // FIXME: modules. eventually. resync this part from clang.
if (hasExternalVisibleStorage()) {
// If a PCH has a result for this name, and we have a local declaration, we
// will have imported the PCH result when adding the local declaration.
// FIXME: For modules, we could have had more declarations added by module
// imoprts since we saw the declaration of the local name.
if (StoredDeclsMap *Map = LookupPtr.getPointer()) {
StoredDeclsMap::iterator I = Map->find(&Name);
if (I != Map->end())
return I->second.getLookupResult();
}
ExternalASTSource *Source = getParentASTContext().getExternalSource();
return Source->FindExternalVisibleDeclsByName(this, Name);
}
#endif
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
if (!Map)
return lookup_result(lookup_iterator(0), lookup_iterator(0));
StoredDeclsMap::iterator I = Map->find(&Name);
if (I == Map->end())
return lookup_result(lookup_iterator(0), lookup_iterator(0));
return I->second.getLookupResult();
}
DeclContext::lookup_result
DeclContext::lookup(DeclarationName Name) {
assert(DeclKind != Decl::LinkageSpec &&
"Should not perform lookups into linkage specs!");
DeclContext *PrimaryContext = getPrimaryContext();
if (PrimaryContext != this)
return PrimaryContext->lookup(Name);
#if AXEL_LOOKUP_CHANGES
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
#endif
if (hasExternalVisibleStorage()) {
// If a PCH has a result for this name, and we have a local declaration, we
// will have imported the PCH result when adding the local declaration.
// FIXME: For modules, we could have had more declarations added by module
// imoprts since we saw the declaration of the local name.
#if AXEL_LOOKUP_CHANGES
if (Map) {
#else
if (StoredDeclsMap *Map = LookupPtr.getPointer()) {
#endif
StoredDeclsMap::iterator I = Map->find(Name);
if (I != Map->end())
return I->second.getLookupResult();
}
ExternalASTSource *Source = getParentASTContext().getExternalSource();
return Source->FindExternalVisibleDeclsByName(this, Name);
}
#ifndef AXEL_LOOKUP_CHANGES
StoredDeclsMap *Map = LookupPtr.getPointer();
if (LookupPtr.getInt())
Map = buildLookup();
#endif
if (!Map)
return lookup_result(lookup_iterator(0), lookup_iterator(0));
StoredDeclsMap::iterator I = Map->find(Name);
if (I == Map->end())
return lookup_result(lookup_iterator(0), lookup_iterator(0));
return I->second.getLookupResult();
}
void DeclContext::localUncachedLookup(DeclarationName Name,
llvm::SmallVectorImpl<NamedDecl *> &Results) {
Results.clear();
// If there's no external storage, just perform a normal lookup and copy
// the results.
if (!hasExternalVisibleStorage() && !hasExternalLexicalStorage() && Name) {
lookup_result LookupResults = lookup(Name);
Results.insert(Results.end(), LookupResults.first, LookupResults.second);
return;
}
// If we have a lookup table, check there first. Maybe we'll get lucky.
if (Name) {
if (StoredDeclsMap *Map = LookupPtr.getPointer()) {
StoredDeclsMap::iterator Pos = Map->find(Name);
if (Pos != Map->end()) {
Results.insert(Results.end(),
Pos->second.getLookupResult().first,
Pos->second.getLookupResult().second);
return;
}
}
}
// Slow case: grovel through the declarations in our chain looking for
// matches.
for (Decl *D = FirstDecl; D; D = D->getNextDeclInContext()) {
if (NamedDecl *ND = dyn_cast<NamedDecl>(D))
if (ND->getDeclName() == Name)
Results.push_back(ND);
}
}
// CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
// array and kernels are launched using cudaLaunchKernel().
void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
FunctionArgList &Args) {
// Build the shadow stack entry at the very start of the function.
// Calculate amount of space we will need for all arguments. If we have no
// args, allocate a single pointer so we still have a valid pointer to the
// argument array that we can pass to runtime, even if it will be unused.
Address KernelArgs = CGF.CreateTempAlloca(
VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args",
llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
// Store pointers to the arguments in a locally allocated launch_args.
for (unsigned i = 0; i < Args.size(); ++i) {
llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, VoidPtrTy);
CGF.Builder.CreateDefaultAlignedStore(
VoidVarPtr, CGF.Builder.CreateConstGEP1_32(KernelArgs.getPointer(), i));
}
llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
// Lookup cudaLaunchKernel function.
// cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
// void **args, size_t sharedMem,
// cudaStream_t stream);
TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
IdentifierInfo &cudaLaunchKernelII =
CGM.getContext().Idents.get("cudaLaunchKernel");
FunctionDecl *cudaLaunchKernelFD = nullptr;
for (const auto &Result : DC->lookup(&cudaLaunchKernelII)) {
if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
cudaLaunchKernelFD = FD;
}
if (cudaLaunchKernelFD == nullptr) {
CGM.Error(CGF.CurFuncDecl->getLocation(),
"Can't find declaration for cudaLaunchKernel()");
return;
}
// Create temporary dim3 grid_dim, block_dim.
ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1);
QualType Dim3Ty = GridDimParam->getType();
Address GridDim =
CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim");
Address BlockDim =
CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
Address ShmemSize =
CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
Address Stream =
CGF.CreateTempAlloca(VoidPtrTy, CGM.getPointerAlign(), "stream");
llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
llvm::FunctionType::get(IntTy,
{/*gridDim=*/GridDim.getType(),
/*blockDim=*/BlockDim.getType(),
/*ShmemSize=*/ShmemSize.getType(),
/*Stream=*/Stream.getType()},
/*isVarArg=*/false),
"__cudaPopCallConfiguration");
CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn,
{GridDim.getPointer(), BlockDim.getPointer(),
ShmemSize.getPointer(), Stream.getPointer()});
// Emit the call to cudaLaunch
llvm::Value *Kernel = CGF.Builder.CreatePointerCast(CGF.CurFn, VoidPtrTy);
CallArgList LaunchKernelArgs;
LaunchKernelArgs.add(RValue::get(Kernel),
cudaLaunchKernelFD->getParamDecl(0)->getType());
LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty);
LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty);
LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()),
cudaLaunchKernelFD->getParamDecl(3)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)),
cudaLaunchKernelFD->getParamDecl(4)->getType());
LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)),
cudaLaunchKernelFD->getParamDecl(5)->getType());
QualType QT = cudaLaunchKernelFD->getType();
QualType CQT = QT.getCanonicalType();
llvm::Type *Ty = CGM.getTypes().ConvertType(CQT);
llvm::FunctionType *FTy = dyn_cast<llvm::FunctionType>(Ty);
const CGFunctionInfo &FI =
CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD);
llvm::FunctionCallee cudaLaunchKernelFn =
CGM.CreateRuntimeFunction(FTy, "cudaLaunchKernel");
CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(),
LaunchKernelArgs);
CGF.EmitBranch(EndBlock);
CGF.EmitBlock(EndBlock);
}