================
@@ -697,4 +707,299 @@ void SemaAMDGPU::handleAMDGPUMaxNumWorkGroupsAttr(Decl *D,
addAMDGPUMaxNumWorkGroupsAttr(D, AL, AL.getArgAsExpr(0), YExpr, ZExpr);
}
+Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
+ ASTContext &Ctx = getASTContext();
+ QualType BoolTy = Ctx.getLogicalOperationType();
+ llvm::APInt False = llvm::APInt::getZero(Ctx.getIntWidth(BoolTy));
+ llvm::APInt True = llvm::APInt::getAllOnes(Ctx.getIntWidth(BoolTy));
+ SourceLocation Loc = CE->getExprLoc();
+
+ if (!CE->getBuiltinCallee())
+ return *ExpandedPredicates
+ .insert(IntegerLiteral::Create(Ctx, False, BoolTy, Loc))
+ .first;
+
+ bool P = false;
+ unsigned BI = CE->getBuiltinCallee();
+ if (Ctx.BuiltinInfo.isAuxBuiltinID(BI))
+ BI = Ctx.BuiltinInfo.getAuxBuiltinID(BI);
+
+ if (BI == AMDGPU::BI__builtin_amdgcn_processor_is) {
+ auto *GFX = dyn_cast<StringLiteral>(CE->getArg(0)->IgnoreParenCasts());
+ if (!GFX) {
+ Diag(Loc, diag::err_amdgcn_processor_is_arg_not_literal);
+ return nullptr;
+ }
+
+ StringRef N = GFX->getString();
+ const TargetInfo &TI = Ctx.getTargetInfo();
+ const TargetInfo *AuxTI = Ctx.getAuxTargetInfo();
+ if (!TI.isValidCPUName(N) && (!AuxTI || !AuxTI->isValidCPUName(N))) {
+ Diag(Loc, diag::err_amdgcn_processor_is_arg_invalid_value) << N;
+ SmallVector<StringRef, 32> ValidList;
+ if (TI.getTriple().getVendor() == llvm::Triple::VendorType::AMD)
+ TI.fillValidCPUList(ValidList);
+ else if (AuxTI) // Since the BI is present it must be an AMDGPU triple.
+ AuxTI->fillValidCPUList(ValidList);
+ if (!ValidList.empty())
+ Diag(Loc, diag::note_amdgcn_processor_is_valid_options)
+ << llvm::join(ValidList, ", ");
+ return nullptr;
+ }
+ if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
+ CE->setType(BoolTy);
+ return *ExpandedPredicates.insert(CE).first;
+ }
+
+ if (auto TID = Ctx.getTargetInfo().getTargetID())
+ P = TID->find(N) == 0;
+ } else {
+ Expr *Arg = CE->getArg(0);
+ if (!Arg || Arg->getType() != Ctx.BuiltinFnTy) {
+ Diag(Loc, diag::err_amdgcn_is_invocable_arg_invalid_value) << Arg;
+ return nullptr;
+ }
+
+ if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
+ CE->setType(BoolTy);
+ return *ExpandedPredicates.insert(CE).first;
+ }
+
+ auto *FD = cast<FunctionDecl>(Arg->getReferencedDeclOfCallee());
+
+ StringRef RF = Ctx.BuiltinInfo.getRequiredFeatures(FD->getBuiltinID());
+ llvm::StringMap<bool> CF;
+ Ctx.getFunctionFeatureMap(CF, FD);
+
+ P = Builtin::evaluateRequiredTargetFeatures(RF, CF);
+ }
+
+ return *ExpandedPredicates
+ .insert(
+ IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc))
+ .first;
+}
+
+bool SemaAMDGPU::IsPredicate(Expr *E) const {
+ return ExpandedPredicates.contains(E);
+}
+
+void SemaAMDGPU::AddPotentiallyUnguardedBuiltinUser(FunctionDecl *FD) {
+ PotentiallyUnguardedBuiltinUsers.insert(FD);
+}
+
+bool SemaAMDGPU::HasPotentiallyUnguardedBuiltinUsage(FunctionDecl *FD) const {
+ return PotentiallyUnguardedBuiltinUsers.contains(FD);
+}
+
+namespace {
+/// This class implements -Wamdgpu-unguarded-builtin-usage.
+///
+/// This is done with a traversal of the AST of a function that includes a
+/// call to a target specific builtin. Whenever we encounter an \c if of the
+/// form: \c if(__builtin_amdgcn_is_invocable), we consider the then statement
+/// guarded.
+class DiagnoseUnguardedBuiltins : public DynamicRecursiveASTVisitor {
+ // TODO: this could eventually be extended to consider what happens when
there
+ // are multiple target architectures specified via
target("arch=gfxXXX")
+ // target("arch=gfxyyy") etc., as well as feature disabling via "-XXX".
+ Sema &SemaRef;
+
+ SmallVector<StringRef> TargetFeatures;
+ SmallVector<std::pair<SourceLocation, StringRef>> CurrentGFXIP;
+ SmallVector<unsigned> GuardedBuiltins;
+
+ static Expr *FindPredicate(Expr *Cond) {
+ if (auto *CE = dyn_cast<CallExpr>(Cond)) {
+ if (CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_is_invocable ||
+ CE->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is)
+ return Cond;
+ } else if (auto *UO = dyn_cast<UnaryOperator>(Cond)) {
+ return FindPredicate(UO->getSubExpr());
+ } else if (auto *BO = dyn_cast<BinaryOperator>(Cond)) {
+ if ((Cond = FindPredicate(BO->getLHS())))
+ return Cond;
+ return FindPredicate(BO->getRHS());
+ }
+ return nullptr;
+ }
+
+ bool EnterPredicateGuardedContext(CallExpr *P);
+ void ExitPredicateGuardedContext(bool WasProcessorCheck);
+ bool TraverseGuardedStmt(Stmt *S, CallExpr *P);
+
+public:
+ DiagnoseUnguardedBuiltins(Sema &SemaRef) : SemaRef(SemaRef) {
+ if (auto *TAT = SemaRef.getCurFunctionDecl(true)->getAttr<TargetAttr>()) {
+ // We use the somewhat misnamed x86 accessors because they provide
exactly
+ // what we require.
+ TAT->getX86AddedFeatures(TargetFeatures);
+ if (auto GFXIP = TAT->getX86Architecture())
+ CurrentGFXIP.emplace_back(TAT->getLocation(), *GFXIP);
+ }
+ }
+
+ bool TraverseLambdaExpr(LambdaExpr *LE) override {
+ if (SemaRef.AMDGPU().HasPotentiallyUnguardedBuiltinUsage(
+ LE->getCallOperator()))
+ return true; // We have already handled this.
+ return DynamicRecursiveASTVisitor::TraverseLambdaExpr(LE);
+ }
+
+ bool TraverseStmt(Stmt *S) override {
+ if (!S)
+ return true;
+ return DynamicRecursiveASTVisitor::TraverseStmt(S);
+ }
+
+ void IssueDiagnostics(Stmt *S) { TraverseStmt(S); }
+
+ bool TraverseIfStmt(IfStmt *If) override {
+ if (auto *CE = dyn_cast_or_null<CallExpr>(FindPredicate(If->getCond())))
+ return TraverseGuardedStmt(If, CE);
+ return DynamicRecursiveASTVisitor::TraverseIfStmt(If);
+ }
+
+ bool TraverseCaseStmt(CaseStmt *CS) override {
+ return TraverseStmt(CS->getSubStmt());
+ }
+
+ bool TraverseConditionalOperator(ConditionalOperator *CO) override {
+ if (auto *CE = dyn_cast_or_null<CallExpr>(FindPredicate(CO->getCond())))
+ return TraverseGuardedStmt(CO, CE);
+ return DynamicRecursiveASTVisitor::TraverseConditionalOperator(CO);
+ }
+
+ bool VisitAsmStmt(AsmStmt *ASM) override;
+ bool VisitCallExpr(CallExpr *CE) override;
+};
+
+bool DiagnoseUnguardedBuiltins::EnterPredicateGuardedContext(CallExpr *P) {
+ bool IsProcessorCheck =
+ P->getBuiltinCallee() == AMDGPU::BI__builtin_amdgcn_processor_is;
+
+ if (IsProcessorCheck) {
+ StringRef G = cast<clang::StringLiteral>(P->getArg(0))->getString();
+ // TODO: handle generic ISAs.
+ if (!CurrentGFXIP.empty() && G != CurrentGFXIP.back().second) {
+ SemaRef.Diag(P->getExprLoc(),
+ diag::err_amdgcn_conflicting_is_processor_options)
+ << P;
+ SemaRef.Diag(CurrentGFXIP.back().first,
+ diag::note_amdgcn_previous_is_processor_guard);
+ }
+ CurrentGFXIP.emplace_back(P->getExprLoc(), G);
+ } else {
+ auto *FD = cast<FunctionDecl>(
+ cast<DeclRefExpr>(P->getArg(0))->getReferencedDeclOfCallee());
+ unsigned ID = FD->getBuiltinID();
+ StringRef F = SemaRef.getASTContext().BuiltinInfo.getRequiredFeatures(ID);
+ GuardedBuiltins.push_back(ID);
+ }
+
+ return IsProcessorCheck;
+}
+
+void DiagnoseUnguardedBuiltins::ExitPredicateGuardedContext(bool WasProcCheck)
{
+ if (WasProcCheck)
+ CurrentGFXIP.pop_back();
+ else
+ GuardedBuiltins.pop_back();
+}
+
+inline std::pair<Stmt *, Stmt *> GetTraversalOrder(Stmt *S) {
+ std::pair<Stmt *, Stmt *> Ordered;
+ Expr *Condition = nullptr;
+
+ if (auto *CO = dyn_cast<ConditionalOperator>(S)) {
+ Condition = CO->getCond();
+ Ordered = {CO->getTrueExpr(), CO->getFalseExpr()};
+ } else if (auto *If = dyn_cast<IfStmt>(S)) {
+ Condition = If->getCond();
+ Ordered = {If->getThen(), If->getElse()};
+ }
+
+ if (auto *UO = dyn_cast<UnaryOperator>(Condition))
+ if (UO->getOpcode() == UnaryOperatorKind::UO_LNot)
+ std::swap(Ordered.first, Ordered.second);
+
+ return Ordered;
+}
+
+bool DiagnoseUnguardedBuiltins::TraverseGuardedStmt(Stmt *S, CallExpr *P) {
+ assert(S && "Unexpected missing Statement!");
+ assert(P && "Unexpected missing Predicate!");
+
+ auto [Guarded, Unguarded] = GetTraversalOrder(S);
+
+ bool WasProcessorCheck = EnterPredicateGuardedContext(P);
+
+ bool Continue = TraverseStmt(Guarded);
+
+ ExitPredicateGuardedContext(WasProcessorCheck);
+
+ return Continue && TraverseStmt(Unguarded);
+}
+
+bool DiagnoseUnguardedBuiltins::VisitAsmStmt(AsmStmt *ASM) {
+ // TODO: should we check if the ASM is valid for the target? Can we?
+ if (!CurrentGFXIP.empty())
+ return true;
+
+ std::string S = ASM->generateAsmString(SemaRef.getASTContext());
+ SemaRef.Diag(ASM->getAsmLoc(), diag::warn_amdgcn_unguarded_asm_stmt) << S;
+ SemaRef.Diag(ASM->getAsmLoc(), diag::note_amdgcn_unguarded_asm_silence) << S;
+
+ return true;
+}
+
+bool DiagnoseUnguardedBuiltins::VisitCallExpr(CallExpr *CE) {
+ unsigned ID = CE->getBuiltinCallee();
+ Builtin::Context &BInfo = SemaRef.getASTContext().BuiltinInfo;
+
+ if (!ID)
+ return true;
+ if (!BInfo.isTSBuiltin(ID))
+ return true;
+ if (ID == AMDGPU::BI__builtin_amdgcn_processor_is ||
+ ID == AMDGPU::BI__builtin_amdgcn_is_invocable)
+ return true;
+ if (llvm::find(GuardedBuiltins, ID) != GuardedBuiltins.end())
+ return true;
+
+ StringRef FL(BInfo.getRequiredFeatures(ID));
+ llvm::StringMap<bool> FeatureMap;
+ if (CurrentGFXIP.empty()) {
+ for (auto &&F : TargetFeatures)
+ FeatureMap[F] = true;
+ for (auto &&GID : GuardedBuiltins)
+ for (auto &&F : llvm::split(BInfo.getRequiredFeatures(GID), ','))
+ FeatureMap[F] = true;
+ } else {
+ static const llvm::Triple AMDGCN("amdgcn-amd-amdhsa");
+ llvm::AMDGPU::fillAMDGPUFeatureMap(CurrentGFXIP.back().second, AMDGCN,
+ FeatureMap);
+ }
+
+ FunctionDecl *BI = CE->getDirectCallee();
+ SourceLocation BICallLoc = CE->getExprLoc();
+ if (Builtin::evaluateRequiredTargetFeatures(FL, FeatureMap)) {
----------------
AlexVlx wrote:
> Thank you for circling back. In what regards the mechanics, that is correct:
> the assumption is that if you used a predicate, or set the target attribute,
> you have an expectation, so if those are wrong or yield contradictions, it
> should be an error.
>
> > There seems to be one part missing here, though. I think there should be
> > some function attribute which marks a function as "guarded", i.e. the
> > caller is responsible for providing an appropriate guard. On the callee
> > side, it suppresses the warning. On the caller side, it propagates the
> > error: it's treated like a "builtin function". Without this, you end up
> > with weird situations where the user can't easily refactor their code: you
> > can't move code from the body of an guard if statement into a separate
> > function.
>
> This is a good idea, thank you, will definitely add it. There will probably
> be a few updates / refinements here over the coming months.
https://github.com/llvm/llvm-project/pull/134016
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
