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