[clang] [llvm] [AMDGPU][clang][CodeGen][opt] Add late-resolved feature identifying predicates (PR #134016)

[via llvm-commits]

github-actions[bot] wrote:

<!--LLVM CODE FORMAT COMMENT: {clang-format}-->


:warning: C/C++ code formatter, clang-format found issues in your code. :warning:

<details>
<summary>
You can test this locally with the following command:
</summary>

``````````bash
git-clang-format --diff HEAD~1 HEAD --extensions h,c,cpp -- clang/test/CodeGen/amdgpu-builtin-cpu-is.c clang/test/CodeGen/amdgpu-builtin-is-invocable.c clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp llvm/lib/Target/AMDGPU/AMDGPUExpandPseudoIntrinsics.cpp clang/lib/Basic/Targets/SPIR.cpp clang/lib/Basic/Targets/SPIR.h clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp clang/lib/Sema/SemaExpr.cpp llvm/lib/Target/AMDGPU/AMDGPU.h llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
``````````

</details>

<details>
<summary>
View the diff from clang-format here.
</summary>

``````````diff
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 7b1a38151..8ad1ab74f 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -292,8 +292,8 @@ static Value *GetOrInsertAMDGPUPredicate(CodeGenFunction &CGF, Twine Name) {
   P->setConstant(true);
   P->setExternallyInitialized(true);
 
-  return CGF.Builder.CreateLoad(RawAddress(P, PTy, CharUnits::One(),
-                                           KnownNonNull));
+  return CGF.Builder.CreateLoad(
+      RawAddress(P, PTy, CharUnits::One(), KnownNonNull));
 }
 
 Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
@@ -600,7 +600,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_processor_is: {
     assert(CGM.getTriple().isSPIRV() &&
            "__builtin_amdgcn_processor_is should never reach CodeGen for "
-             "concrete targets!");
+           "concrete targets!");
     StringRef Proc = cast<clang::StringLiteral>(E->getArg(0))->getString();
     return GetOrInsertAMDGPUPredicate(*this, "llvm.amdgcn.is." + Proc);
   }
@@ -609,7 +609,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
            "__builtin_amdgcn_is_invocable should never reach CodeGen for "
            "concrete targets!");
     auto FD = cast<FunctionDecl>(
-      cast<DeclRefExpr>(E->getArg(0))->getReferencedDeclOfCallee());
+        cast<DeclRefExpr>(E->getArg(0))->getReferencedDeclOfCallee());
     StringRef RF =
         getContext().BuiltinInfo.getRequiredFeatures(FD->getBuiltinID());
     return GetOrInsertAMDGPUPredicate(*this, "llvm.amdgcn.has." + RF);
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 24f5262ab..bd0183ae4 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -20549,14 +20549,16 @@ static bool ValidateAMDGPUPredicateBI(Sema &Sema, CallExpr *CE) {
         (!Sema.getASTContext().getAuxTargetInfo() ||
          !Sema.getASTContext().getAuxTargetInfo()->isValidCPUName(N))) {
       Sema.Diag(CE->getExprLoc(),
-                diag::err_amdgcn_processor_is_arg_invalid_value) << N;
+                diag::err_amdgcn_processor_is_arg_invalid_value)
+          << N;
       return false;
     }
   } else {
     auto Arg = CE->getArg(0);
     if (!Arg || Arg->getType() != Sema.getASTContext().BuiltinFnTy) {
       Sema.Diag(CE->getExprLoc(),
-                diag::err_amdgcn_is_invocable_arg_invalid_value) << Arg;
+                diag::err_amdgcn_is_invocable_arg_invalid_value)
+          << Arg;
       return false;
     }
   }
@@ -20568,10 +20570,9 @@ static Expr *MaybeHandleAMDGPUPredicateBI(Sema &Sema, Expr *E, bool &Invalid) {
   if (auto UO = dyn_cast<UnaryOperator>(E)) {
     auto SE = dyn_cast<CallExpr>(UO->getSubExpr());
     if (IsAMDGPUPredicateBI(SE)) {
-      assert(
-        UO->getOpcode() == UnaryOperator::Opcode::UO_LNot &&
-        "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
-          "can only be used as operands of logical ops!");
+      assert(UO->getOpcode() == UnaryOperator::Opcode::UO_LNot &&
+             "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
+             "can only be used as operands of logical ops!");
 
       if (!ValidateAMDGPUPredicateBI(Sema, SE)) {
         Invalid = true;
@@ -20588,10 +20589,9 @@ static Expr *MaybeHandleAMDGPUPredicateBI(Sema &Sema, Expr *E, bool &Invalid) {
     auto LHS = dyn_cast<CallExpr>(BO->getLHS());
     auto RHS = dyn_cast<CallExpr>(BO->getRHS());
     if (IsAMDGPUPredicateBI(LHS) && IsAMDGPUPredicateBI(RHS)) {
-      assert(
-          BO->isLogicalOp() &&
-          "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
-            "can only be used as operands of logical ops!");
+      assert(BO->isLogicalOp() &&
+             "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
+             "can only be used as operands of logical ops!");
 
       if (!ValidateAMDGPUPredicateBI(Sema, LHS) ||
           !ValidateAMDGPUPredicateBI(Sema, RHS)) {
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUExpandPseudoIntrinsics.cpp b/llvm/lib/Target/AMDGPU/AMDGPUExpandPseudoIntrinsics.cpp
index 125051c6a..77084de61 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUExpandPseudoIntrinsics.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUExpandPseudoIntrinsics.cpp
@@ -108,7 +108,7 @@ inline void collectUsers(Value *V, ModulePassManager &AlwaysInliner,
 
 std::pair<PreservedAnalyses, bool>
 handlePredicate(const GCNSubtarget &ST, ModuleAnalysisManager &MAM,
-                SmallDenseMap<Function *, Function *>& InlinableClones,
+                SmallDenseMap<Function *, Function *> &InlinableClones,
                 GlobalVariable *P) {
   auto PV = P->getName().substr(P->getName().rfind('.') + 1).str();
   auto Dx = PV.find(',');

``````````

</details>


https://github.com/llvm/llvm-project/pull/134016