CXXRecordDecl *Sema::createLambdaClosureType(SourceRange IntroducerRange,
bool KnownDependent) {
DeclContext *DC = CurContext;
while (!(DC->isFunctionOrMethod() || DC->isRecord() || DC->isFileContext()))
DC = DC->getParent();
// Start constructing the lambda class.
CXXRecordDecl *Class = CXXRecordDecl::CreateLambda(Context, DC,
IntroducerRange.getBegin(),
KnownDependent);
DC->addDecl(Class);
return Class;
}
bool SymbolResolverCallback::LookupObject(LookupResult& R, Scope* S) {
if (!ShouldResolveAtRuntime(R, S))
return false;
if (m_IsRuntime) {
// We are currently parsing an EvaluateT() expression
if (!m_Resolve)
return false;
// Only for demo resolve all unknown objects to cling::test::Tester
if (!m_TesterDecl) {
clang::Sema& SemaR = m_Interpreter->getSema();
clang::NamespaceDecl* NSD = utils::Lookup::Namespace(&SemaR, "cling");
NSD = utils::Lookup::Namespace(&SemaR, "test", NSD);
m_TesterDecl = utils::Lookup::Named(&SemaR, "Tester", NSD);
}
assert (m_TesterDecl && "Tester not found!");
R.addDecl(m_TesterDecl);
return true; // Tell clang to continue.
}
// We are currently NOT parsing an EvaluateT() expression.
// Escape the expression into an EvaluateT() expression.
ASTContext& C = R.getSema().getASTContext();
DeclContext* DC = 0;
// For DeclContext-less scopes like if (dyn_expr) {}
while (!DC) {
DC = static_cast<DeclContext*>(S->getEntity());
S = S->getParent();
}
DeclarationName Name = R.getLookupName();
IdentifierInfo* II = Name.getAsIdentifierInfo();
SourceLocation Loc = R.getNameLoc();
VarDecl* Res = VarDecl::Create(C, DC, Loc, Loc, II, C.DependentTy,
/*TypeSourceInfo*/0, SC_None);
// Annotate the decl to give a hint in cling. FIXME: Current implementation
// is a gross hack, because TClingCallbacks shouldn't know about
// EvaluateTSynthesizer at all!
SourceRange invalidRange;
Res->addAttr(new (C) AnnotateAttr(invalidRange, C, "__ResolveAtRuntime", 0));
R.addDecl(Res);
DC->addDecl(Res);
// Say that we can handle the situation. Clang should try to recover
return true;
}
bool DeclExtractor::ExtractDecl(FunctionDecl* FD) {
llvm::SmallVector<NamedDecl*, 4> TouchedDecls;
CompoundStmt* CS = dyn_cast<CompoundStmt>(FD->getBody());
assert(CS && "Function body not a CompoundStmt?");
DeclContext* DC = FD->getTranslationUnitDecl();
Scope* TUScope = m_Sema->TUScope;
assert(TUScope == m_Sema->getScopeForContext(DC) && "TU scope from DC?");
llvm::SmallVector<Stmt*, 4> Stmts;
for (CompoundStmt::body_iterator I = CS->body_begin(), EI = CS->body_end();
I != EI; ++I) {
DeclStmt* DS = dyn_cast<DeclStmt>(*I);
if (!DS) {
Stmts.push_back(*I);
continue;
}
for (DeclStmt::decl_iterator J = DS->decl_begin();
J != DS->decl_end(); ++J) {
NamedDecl* ND = dyn_cast<NamedDecl>(*J);
if (isa<UsingDirectiveDecl>(*J))
continue; // FIXME: Here we should be more elegant.
if (ND) {
if (Stmts.size()) {
// We need to emit a new custom wrapper wrapping the stmts
EnforceInitOrder(Stmts);
assert(!Stmts.size() && "Stmt list must be flushed.");
}
// We know the transaction is closed, but it is safe.
getTransaction()->forceAppend(ND);
DeclContext* OldDC = ND->getDeclContext();
// Make sure the decl is not found at its old possition
ND->getLexicalDeclContext()->removeDecl(ND);
if (Scope* S = m_Sema->getScopeForContext(OldDC)) {
S->RemoveDecl(ND);
if (utils::Analyze::isOnScopeChains(ND, *m_Sema))
m_Sema->IdResolver.RemoveDecl(ND);
}
// For variable definitions causing var/function ambiguity such as:
// MyClass my();, C++ standard says it shall be resolved as a function
//
// In the particular context this definition is inside a function
// already, but clang thinks it as a lambda, so we need to ignore the
// check decl context vs lexical decl context.
if (ND->getDeclContext() == ND->getLexicalDeclContext()
|| isa<FunctionDecl>(ND))
ND->setLexicalDeclContext(DC);
else
assert(0 && "Not implemented: Decl with different lexical context");
ND->setDeclContext(DC);
if (VarDecl* VD = dyn_cast<VarDecl>(ND)) {
VD->setStorageClass(SC_None);
}
clearLinkage(ND);
TouchedDecls.push_back(ND);
}
}
}
bool hasNoErrors = !CheckForClashingNames(TouchedDecls, DC, TUScope);
if (hasNoErrors) {
for (size_t i = 0; i < TouchedDecls.size(); ++i) {
// We should skip the checks for annonymous decls and we should not
// register them in the lookup.
if (!TouchedDecls[i]->getDeclName())
continue;
m_Sema->PushOnScopeChains(TouchedDecls[i],
m_Sema->getScopeForContext(DC),
/*AddCurContext*/!isa<UsingDirectiveDecl>(TouchedDecls[i]));
// The transparent DeclContexts (eg. scopeless enum) doesn't have
// scopes. While extracting their contents we need to update the
// lookup tables and telling them to pick up the new possitions
// in the AST.
if (DeclContext* InnerDC = dyn_cast<DeclContext>(TouchedDecls[i])) {
if (InnerDC->isTransparentContext()) {
// We can't PushDeclContext, because we don't have scope.
Sema::ContextRAII pushedDC(*m_Sema, InnerDC);
for(DeclContext::decl_iterator DI = InnerDC->decls_begin(),
DE = InnerDC->decls_end(); DI != DE ; ++DI) {
if (NamedDecl* ND = dyn_cast<NamedDecl>(*DI))
InnerDC->makeDeclVisibleInContext(ND);
}
}
}
}
}
CS->setStmts(*m_Context, Stmts.data(), Stmts.size());
// The order matters, because when we extract decls from the wrapper we
// append them to the transaction. If the transaction gets unloaded it will
// introduce a fake dependency, so put the move last.
Transaction* T = getTransaction();
for (Transaction::iterator I = T->decls_begin(), E = T->decls_end();
I != E; ++I)
if (!I->m_DGR.isNull() && I->m_DGR.isSingleDecl()
&& I->m_DGR.getSingleDecl() == T->getWrapperFD()) {
T->erase(I);
break;
}
T->forceAppend(FD);
// Put the wrapper after its declarations. (Nice when AST dumping)
DC->removeDecl(FD);
DC->addDecl(FD);
return hasNoErrors;
}
// access image memory at given index
Expr *ASTTranslate::accessMemImgAt(DeclRefExpr *LHS, HipaccAccessor *Acc,
MemoryAccess mem_acc, Expr *idx_x, Expr *idx_y) {
Expr *result, *coord;
// mark image as being used within the kernel
Kernel->setUsed(LHS->getNameInfo().getAsString());
// construct coordinate: (int2)(gid_x, gid_y)
coord = createBinaryOperator(Ctx, idx_x, idx_y, BO_Comma, Ctx.IntTy);
coord = createParenExpr(Ctx, coord);
QualType QTcoord = simdTypes.getSIMDType(Ctx.IntTy, "int", SIMD2);
coord = createCStyleCastExpr(Ctx, QTcoord, CK_VectorSplat, coord, nullptr,
Ctx.getTrivialTypeSourceInfo(QTcoord));
FunctionDecl *image_function = getImageFunction(Acc, mem_acc);
// create function call for image objects in OpenCL
if (mem_acc == READ_ONLY) {
// parameters for read_image
SmallVector<Expr *, 16> args;
args.push_back(LHS);
args.push_back(kernelSamplerRef);
args.push_back(coord);
result = createFunctionCall(Ctx, image_function, args);
QualType QT = Acc->getImage()->getType();
if (QT->isVectorType()) {
SmallVector<Expr *, 16> args;
args.push_back(result);
result = createFunctionCall(Ctx, getConvertFunction(QT), args);
} else {
result = createExtVectorElementExpr(Ctx, QT, result, "x");
}
} else {
QualType QT;
// determine cast type for write_image functions
if (image_function == builtins.getBuiltinFunction(OPENCLBIwrite_imagei)) {
QT = simdTypes.getSIMDType(Ctx.IntTy, "int", SIMD4);
} else if (image_function ==
builtins.getBuiltinFunction(OPENCLBIwrite_imageui)) {
QT = simdTypes.getSIMDType(Ctx.UnsignedIntTy, "uint", SIMD4);
} else {
QT = simdTypes.getSIMDType(Ctx.FloatTy, "float", SIMD4);
}
// writeImageRHS is set by VisitBinaryOperator - side effect
if (!writeImageRHS->getType()->isVectorType()) {
// introduce temporary for propagating the RHS to a vector
std::string tmp_lit("_tmp" + std::to_string(literalCount++));
VarDecl *tmp_decl = createVarDecl(Ctx, kernelDecl, tmp_lit, QT,
writeImageRHS);
DeclContext *DC = FunctionDecl::castToDeclContext(kernelDecl);
DC->addDecl(tmp_decl);
DeclRefExpr *tmp_dre = createDeclRefExpr(Ctx, tmp_decl);
preStmts.push_back(createDeclStmt(Ctx, tmp_decl));
preCStmt.push_back(curCStmt);
writeImageRHS = tmp_dre;
}
if (writeImageRHS->getType() != QT) {
// convert to proper vector type
SmallVector<Expr *, 16> args;
args.push_back(writeImageRHS);
writeImageRHS = createFunctionCall(Ctx, getConvertFunction(QT), args);
}
// parameters for write_image
SmallVector<Expr *, 16> args;
args.push_back(LHS);
args.push_back(coord);
args.push_back(writeImageRHS);
result = createFunctionCall(Ctx, image_function, args);
}
return result;
}
// access linear texture memory at given index
Expr *ASTTranslate::accessMemTexAt(DeclRefExpr *LHS, HipaccAccessor *Acc,
MemoryAccess mem_acc, Expr *idx_x, Expr *idx_y) {
// mark image as being used within the kernel
Kernel->setUsed(LHS->getNameInfo().getAsString());
FunctionDecl *texture_function = getTextureFunction(Acc, mem_acc);
// clone Decl
TemplateArgumentListInfo templateArgs(LHS->getLAngleLoc(),
LHS->getRAngleLoc());
for (auto template_arg : LHS->template_arguments())
templateArgs.addArgument(template_arg);
assert(isa<ParmVarDecl>(LHS->getDecl()) && "texture variable must be a ParmVarDecl!");
ParmVarDecl *PVD = dyn_cast<ParmVarDecl>(LHS->getDecl());
DeclRefExpr *LHStex = DeclRefExpr::Create(Ctx,
LHS->getQualifierLoc(),
LHS->getTemplateKeywordLoc(),
CloneDeclTex(PVD, "_tex"),
LHS->refersToEnclosingVariableOrCapture(),
LHS->getLocation(),
LHS->getType(), LHS->getValueKind(),
LHS->getFoundDecl(),
LHS->getNumTemplateArgs()?&templateArgs:0);
setExprProps(LHS, LHStex);
// parameters for __ldg, tex1Dfetch, tex2D, or surf2Dwrite
SmallVector<Expr *, 16> args;
if (mem_acc == READ_ONLY) {
switch (Kernel->useTextureMemory(Acc)) {
case Texture::None:
assert(0 && "texture expected.");
case Texture::Linear1D:
args.push_back(LHStex);
args.push_back(createBinaryOperator(Ctx, createBinaryOperator(Ctx,
createParenExpr(Ctx, idx_y), getStrideDecl(Acc), BO_Mul,
Ctx.IntTy), idx_x, BO_Add, Ctx.IntTy));
break;
case Texture::Linear2D:
case Texture::Array2D:
args.push_back(LHStex);
args.push_back(idx_x);
args.push_back(idx_y);
break;
case Texture::Ldg:
// __ldg(&arr[idx])
args.push_back(createUnaryOperator(Ctx, accessMemArrAt(LHS,
getStrideDecl(Acc), idx_x, idx_y), UO_AddrOf, Ctx.IntTy));
break;
}
} else {
// writeImageRHS is set by VisitBinaryOperator - side effect
QualType QT = Acc->getImage()->getType();
if (writeImageRHS->IgnoreImpCasts()->getType() != QT) {
// introduce temporary for implicit casts
std::string tmp_lit("_tmp" + std::to_string(literalCount++));
VarDecl *tmp_decl = createVarDecl(Ctx, kernelDecl, tmp_lit, QT,
writeImageRHS);
DeclContext *DC = FunctionDecl::castToDeclContext(kernelDecl);
DC->addDecl(tmp_decl);
DeclRefExpr *tmp_dre = createDeclRefExpr(Ctx, tmp_decl);
preStmts.push_back(createDeclStmt(Ctx, tmp_decl));
preCStmt.push_back(curCStmt);
writeImageRHS = tmp_dre;
} else {
writeImageRHS = createParenExpr(Ctx, writeImageRHS);
}
args.push_back(writeImageRHS);
args.push_back(LHStex);
// byte addressing required for surf2Dwrite
args.push_back(createBinaryOperator(Ctx, idx_x, createIntegerLiteral(Ctx,
static_cast<int32_t>(Acc->getImage()->getPixelSize())), BO_Mul,
Ctx.IntTy));
args.push_back(idx_y);
}
return createFunctionCall(Ctx, texture_function, args);
}
bool DeclExtractor::ExtractDecl(Decl* D) {
FunctionDecl* FD = dyn_cast<FunctionDecl>(D);
if (FD) {
if (FD->getNameAsString().find("__cling_Un1Qu3"))
return true;
llvm::SmallVector<NamedDecl*, 4> TouchedDecls;
CompoundStmt* CS = dyn_cast<CompoundStmt>(FD->getBody());
assert(CS && "Function body not a CompoundStmt?");
DeclContext* DC = FD->getTranslationUnitDecl();
Scope* TUScope = m_Sema->TUScope;
assert(TUScope == m_Sema->getScopeForContext(DC) && "TU scope from DC?");
llvm::SmallVector<Stmt*, 4> Stmts;
for (CompoundStmt::body_iterator I = CS->body_begin(), EI = CS->body_end();
I != EI; ++I) {
DeclStmt* DS = dyn_cast<DeclStmt>(*I);
if (!DS) {
Stmts.push_back(*I);
continue;
}
for (DeclStmt::decl_iterator J = DS->decl_begin();
J != DS->decl_end(); ++J) {
NamedDecl* ND = dyn_cast<NamedDecl>(*J);
if (ND) {
DeclContext* OldDC = ND->getDeclContext();
// Make sure the decl is not found at its old possition
OldDC->removeDecl(ND);
if (Scope* S = m_Sema->getScopeForContext(OldDC)) {
S->RemoveDecl(ND);
m_Sema->IdResolver.RemoveDecl(ND);
}
if (ND->getDeclContext() == ND->getLexicalDeclContext())
ND->setLexicalDeclContext(DC);
else
assert("Not implemented: Decl with different lexical context");
ND->setDeclContext(DC);
if (VarDecl* VD = dyn_cast<VarDecl>(ND)) {
VD->setStorageClass(SC_None);
VD->setStorageClassAsWritten(SC_None);
// if we want to print the result of the initializer of int i = 5
// or the default initializer int i
if (I+1 == EI || !isa<NullStmt>(*(I+1))) {
QualType VDTy = VD->getType().getNonReferenceType();
Expr* DRE = m_Sema->BuildDeclRefExpr(VD, VDTy,VK_LValue,
SourceLocation()
).take();
Stmts.push_back(DRE);
}
}
// force recalc of the linkage (to external)
ND->ClearLinkageCache();
TouchedDecls.push_back(ND);
}
}
}
bool hasNoErrors = !CheckForClashingNames(TouchedDecls, DC, TUScope);
if (hasNoErrors) {
for (size_t i = 0; i < TouchedDecls.size(); ++i) {
m_Sema->PushOnScopeChains(TouchedDecls[i],
m_Sema->getScopeForContext(DC),
/*AddCurContext*/!isa<UsingDirectiveDecl>(TouchedDecls[i]));
// The transparent DeclContexts (eg. scopeless enum) doesn't have
// scopes. While extracting their contents we need to update the
// lookup tables and telling them to pick up the new possitions
// in the AST.
if (DeclContext* InnerDC = dyn_cast<DeclContext>(TouchedDecls[i])) {
if (InnerDC->isTransparentContext()) {
// We can't PushDeclContext, because we don't have scope.
Sema::ContextRAII pushedDC(*m_Sema, InnerDC);
for(DeclContext::decl_iterator DI = InnerDC->decls_begin(),
DE = InnerDC->decls_end(); DI != DE ; ++DI) {
if (NamedDecl* ND = dyn_cast<NamedDecl>(*DI))
InnerDC->makeDeclVisibleInContext(ND);
}
}
}
// Append the new top level decl to the current transaction.
getTransaction()->appendUnique(DeclGroupRef(TouchedDecls[i]));
}
}
CS->setStmts(*m_Context, Stmts.data(), Stmts.size());
// Put the wrapper after its declarations. (Nice when AST dumping)
DC->removeDecl(FD);
DC->addDecl(FD);
return hasNoErrors;
}
return true;
}
