[Hopper] Move dummy moves before wgmma.fence (#9530)

0a0a57c introduced these moves before `wgmma.mma_async`, but producer
and consumer of register inputs need to be separated by a `wgmma.fence`.

See also: #9433

Author: chsigg

test/Conversion/tritongpu_to_llvm_hopper.mlir

@@ -140,6 +140,7 @@ module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} {
   // CHECK: %[[A_MOV1:.*]] = llvm.inline_asm has_side_effects asm_dialect = att operand_attrs = [] "mov.b32 $0, $1;", "=r,r" %{{.*}} : (i32) -> i32
   // CHECK: %[[A_MOV2:.*]] = llvm.inline_asm has_side_effects asm_dialect = att operand_attrs = [] "mov.b32 $0, $1;", "=r,r" %{{.*}} : (i32) -> i32
   // CHECK: %[[A_MOV3:.*]] = llvm.inline_asm has_side_effects asm_dialect = att operand_attrs = [] "mov.b32 $0, $1;", "=r,r" %{{.*}} : (i32) -> i32
+  // CHECK: nvvm.wgmma.fence.aligned
   // CHECK: %[[A_PACK0:.*]] = llvm.insertvalue %[[A_MOV0]], %{{.*}}[0] : !llvm.struct<(i32, i32, i32, i32)>
   // CHECK: %[[A_PACK1:.*]] = llvm.insertvalue %[[A_MOV1]], %[[A_PACK0]][1] : !llvm.struct<(i32, i32, i32, i32)>
   // CHECK: %[[A_PACK2:.*]] = llvm.insertvalue %[[A_MOV2]], %[[A_PACK1]][2] : !llvm.struct<(i32, i32, i32, i32)>

third_party/nvidia/lib/TritonNVIDIAGPUToLLVM/DotOpToLLVM/WGMMA.cpp

@@ -302,6 +302,8 @@ LogicalResult convertDot(const LLVMTypeConverter *typeConverter,
           // WGMMA, we can safely overlap transformations on the A operand with
           // the previous-iteration WGMMA still in flight.
           for (Value &regAVal : regA) {
+            OpBuilder::InsertionGuard guard(rewriter);
+            rewriter.setInsertionPoint(startSequence);
             Type regTy = regAVal.getType();
             if (!regTy.isIntOrFloat() || regTy.getIntOrFloatBitWidth() != 32) {
               return mlir::emitError(loc, "unsupported WGMMA A register type ")