[BACKEND] Add workaround for TMA device side descriptor race condition (#7293)

A fence is currently missing after TMA update. Adding extra instruction
to force the fence to be materialized.

Patch provided by @mbrookhart

Co-authored-by: Matthew Brookhart <[email protected]>

Author: ThomasRaoux

test/Conversion/tritonnvidiagpu_to_llvm.mlir

@@ -182,6 +182,9 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.targ
 module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, ttg.target = "cuda:90", "ttg.threads-per-warp" = 32 : i32} {
   tt.func public @tensormap_fenceproxy_acquire(%arg0: !tt.ptr<i8> {tt.divisibility = 16 : i32}) attributes {noinline = false} {
     // CHECK: fence.proxy.tensormap::generic.acquire.gpu [ $0 + 0 ], 0x80;
+    // ptxas missing fence workaround:
+    // CHECK: cp.async.bulk.commit_group
+    // CHECK: cp.async.bulk.wait_group.read 0
     ttng.tensormap_fenceproxy_acquire %arg0 : !tt.ptr<i8>
     tt.return
   }

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/TMAToLLVM.cpp

@@ -206,6 +206,13 @@ struct TensormapFenceproxyAcquireOpConversion
         *ptxBuilder.create<>("fence.proxy.tensormap::generic.acquire.gpu");
     fence(descAddrOpr, sizeOpr).predicate(pred);
 
+    // Workaround for a ptxas bug missing a fence after generic.acquire.gpu.
+    // TODO: remove the workaround once ptxas is fixed.
+    auto &commit = *ptxBuilder.create<>("cp.async.bulk.commit_group");
+    commit().predicate(pred);
+    auto &wait = *ptxBuilder.create<>("cp.async.bulk.wait_group.read 0");
+    wait().predicate(pred);
+
     ptxBuilder.launch(rewriter, loc, getVoidType());
 
     // We run the fence on a single warp, then use a barrier to synchronize the