[Backend] Check if wait_barrier is a constant true (#7508)

When the `wait_barrier` predicate ends up being true, either generated
from Gluon or for some other optimization, generate the non-predicated
version of the instruction, which is slightly more efficient due to 1
less branch. I observed a speedup of like 3 TFLOPS on the attention
kernel because of this. It was stable and reproducible and not just
noise.

Author: Mogball

test/Conversion/tritonnvidiagpu_to_llvm.mlir

@@ -17,11 +17,29 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
 #smem = #ttg.shared_memory
 module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
   // CHECK-LABEL: wait_barrier
-  tt.func @wait_barrier(%alloc: !ttg.memdesc<1xi64, #shared0, #smem>, %phase: i32) {
+  tt.func @wait_barrier(%alloc: !ttg.memdesc<1xi64, #shared0, #smem>, %phase: i32, %pred: i1) {
     // CHECK: waitLoop:
     // CHECK: mbarrier.try_wait.parity.shared.b64
     // CHECK: @!complete bra.uni waitLoop
+    // CHECK-NOT: skipWait
+    // CHECK: %{{[0-9]+}}, %arg1 :
     ttng.wait_barrier %alloc, %phase : !ttg.memdesc<1xi64, #shared0, #smem>
+    %true = arith.constant true
+
+    // CHECK: waitLoop:
+    // CHECK: mbarrier.try_wait.parity.shared.b64
+    // CHECK: @!complete bra.uni waitLoop
+    // CHECK-NOT: skipWait
+    // CHECK: %{{[0-9]+}}, %arg1 :
+    ttng.wait_barrier %alloc, %phase, %true : !ttg.memdesc<1xi64, #shared0, #smem>
+
+    // CHECK: @!$2 bra.uni skipWait
+    // CHECK: waitLoop:
+    // CHECK: mbarrier.try_wait.parity.shared.b64
+    // CHECK: @!complete bra.uni waitLoop
+    // CHECK: skipWait:
+    // CHECK: %{{[0-9]+}}, %arg1, %arg2 :
+    ttng.wait_barrier %alloc, %phase, %pred : !ttg.memdesc<1xi64, #shared0, #smem>
     tt.return
   }
 

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/BarrierOpToLLVM.cpp

@@ -160,7 +160,8 @@ struct WaitBarrierOpConversion
         typeConverter->convertType(op.getAlloc().getType().getElementType()),
         rewriter);
     auto loc = op.getLoc();
-    bool predicated = adaptor.getPred() != nullptr;
+    bool predicated =
+        adaptor.getPred() && !matchPattern(op.getPred(), m_NonZero());
     std::string ptx;
     if (targetInfo->getComputeCapability() < 90) {
       if (!predicated) {