[clang] [llvm] [AMDGPU] Extend __builtin_amdgcn_sched_group_barrier to support rules. (PR #85304)

[Matt Arsenault via llvm-commits]

================
@@ -2747,23 +2749,32 @@ void IGroupLPDAGMutation::initSchedGroupBarrierPipelineStage(
   int32_t SGMask = SGB.getOperand(0).getImm();
   int32_t Size = SGB.getOperand(1).getImm();
   int32_t SyncID = SGB.getOperand(2).getImm();
-  std::optional<int32_t> RuleID =
+  std::optional<uint64_t> RuleMask =
       (SGB.getOpcode() == AMDGPU::SCHED_GROUP_BARRIER_RULE)
           ? SGB.getOperand(3).getImm()
           : std::optional<int32_t>(std::nullopt);
 
-  // Sanitize the input
-  if (RuleID && (!SchedGroupBarrierRuleCallBacks.size() ||
-                 *RuleID > (int)(SchedGroupBarrierRuleCallBacks.size() - 1))) {
-    RuleID = std::nullopt;
-    llvm_unreachable("Bad rule ID!");
-  }
-
   auto SG = &SyncedSchedGroups[SyncID].emplace_back((SchedGroupMask)SGMask,
                                                     Size, SyncID, DAG, TII);
-  if (RuleID)
-    SG->addRule(SchedGroupBarrierRuleCallBacks[*RuleID](SG->getSGID()));
+  // Process the input mask
+  if (RuleMask) {
+    uint64_t TheMask = *RuleMask;
+    unsigned NextID = 0;
+    while (TheMask) {
+      if (!(TheMask & 0x1)) {
+        TheMask >>= 1;
+        ++NextID;
+        continue;
+      }
+      if ((!SchedGroupBarrierRuleCallBacks.size() ||
+           NextID > SchedGroupBarrierRuleCallBacks.size() - 1))
----------------
arsenm wrote:

the !size() check is redundant? 

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