diff --git a/‎src/transform/lower_hopper_intrin.cc‎
Lines changed: 8 additions & 3 deletions b/‎src/transform/lower_hopper_intrin.cc‎
Lines changed: 8 additions & 3 deletions
diff --git a/‎src/transform/lower_shared_barrier.cc‎
Lines changed: 5 additions & 2 deletions b/‎src/transform/lower_shared_barrier.cc‎
Lines changed: 5 additions & 2 deletions
@@ -142,6 +142,14 @@ class LowerHopperIntrin : public StmtExprMutator {
           auto stmts = prefetch_calls_;
           stmts.insert(stmts.end(), init_mbarrier_calls_.begin(),
                        init_mbarrier_calls_.end());
+          // Add fence_barrier_init inside the if block, after the barrier init
+          // calls. This ensures the fence is only executed by the thread that
+          // initializes the barriers, matching the pattern in DeepGEMM/cutlass
+          if (!init_mbarrier_calls_.empty()) {
+            stmts.push_back(Evaluate(Call(DataType::Handle(),
+                                          tvm::tl::ptx_fence_barrier_init(),
+                                          {})));
+          }
           PrimExpr condition;
           if (!disable_shuffle_elect_) {
             condition = Call(DataType::Bool(), tl_shuffle_elect(), {0});
@@ -159,9 +167,6 @@ class LowerHopperIntrin : public StmtExprMutator {
             // with an appropriate sync instruction with the right scope to
             // ensure visibility eg. __syncthreads() or a cluster_arrive() +
             // cluster_wait()
-            Stmt mem_fence = Evaluate(Call(
-                DataType::Handle(), tvm::tl::ptx_fence_barrier_init(), {}));
-            stmt_seq.push_back(mem_fence);
             Stmt mem_sync =
                 Evaluate(Call(DataType::Handle(), builtin::tvm_storage_sync(),
                               {StringImm("shared")}));
 
@@ -137,14 +137,17 @@ class SharedBarrierRewriter : public StmtExprMutator {
     } else {
       condition = EQ(thread_var_->var, 0);
     }
+    // Add fence_barrier_init inside the if block, after the barrier init calls
+    // This ensures the fence is only executed by the thread that initializes
+    // the barriers, matching the pattern in DeepGEMM/cutlass
+    init_mbarrier_calls_.push_back(
+        Evaluate(Call(DataType::Handle(), ptx_fence_barrier_init(), {})));
     new_body.push_back(IfThenElse(condition,
                                   init_mbarrier_calls_.size() == 1
                                       ? init_mbarrier_calls_.back()
                                       : SeqStmt(init_mbarrier_calls_),
                                   Stmt()));
 
-    new_body.push_back(
-        Evaluate(Call(DataType::Handle(), ptx_fence_barrier_init(), {})));
     new_body.push_back(Evaluate(
         Call(DataType::Handle(), builtin::tvm_storage_sync(),
              {StringImm(has_cluster_barrier_ ? "cluster" : "shared")})));