diff --git a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_fp8_dpas.hpp b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_fp8_dpas.hpp
index 1638bf7af..6ebdf725a 100644
--- a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_fp8_dpas.hpp
+++ b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_fp8_dpas.hpp
@@ -733,14 +733,6 @@ CUTE_DEVICE void MoEGEMM(const ElementA* Activations, const ElementB* Weights,
   int group_range = item.get_group_range(1);
   int local_id = item.get_local_linear_id();
 
-  if (group_id == 0 && local_id == 0) {
-    auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
-                                sycl::memory_scope::device,
-                                sycl::access::address_space::global_space>(
-        atomic_buffer[0]);
-    atm.store(0);
-  }
-
   int pre_rows = 0;
   int pre_tiles = 0;
 
@@ -935,6 +927,8 @@ void moe_prefill_fp8_dpas_per_tensor_dispatch(
     throw std::runtime_error(
         "moe_prefill_fp8_dpas(per-tensor): failed to allocate atomic buffer");
   }
+  // Initialize the persistent scheduler counter before kernel launch.
+  q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();
 
 #define ARK_DPAS_PT_LAUNCH(policy)                                             \
   MoEGEMMLauncher<'R', 'R', policy, ScaleMode::kPerTensor>(                    \
@@ -993,6 +987,8 @@ void moe_prefill_fp8_dpas_per_group_dispatch(
     throw std::runtime_error(
         "moe_prefill_fp8_dpas(per-group): failed to allocate atomic buffer");
   }
+  // Initialize the persistent scheduler counter before kernel launch.
+  q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();
 
 #define ARK_DPAS_PG_LAUNCH(policy)                                             \
   MoEGEMMLauncher<'R', 'C', policy, ScaleMode::kPerGroup>(                     \
diff --git a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_int_dpas.hpp b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_int_dpas.hpp
index 2e05dfa44..f1be20c6f 100644
--- a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_int_dpas.hpp
+++ b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_int_dpas.hpp
@@ -584,14 +584,6 @@ CUTE_DEVICE void MoEGEMM_int(const ElementA* Activations,
   int group_range = item.get_group_range(1);
   int local_id = item.get_local_linear_id();
 
-  if (group_id == 0 && local_id == 0) {
-    auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
-                                sycl::memory_scope::device,
-                                sycl::access::address_space::global_space>(
-        atomic_buffer[0]);
-    atm.store(0);
-  }
-
   int pre_rows = 0;
   int pre_tiles = 0;
 
@@ -785,6 +777,8 @@ void moe_prefill_int_dpas_per_tensor_dispatch(
     throw std::runtime_error(
         "moe_prefill_int_dpas(per-tensor): failed to allocate atomic buffer");
   }
+  // Initialize the persistent scheduler counter before kernel launch.
+  q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();
 
 #define ARK_DPAS_INT_PT_LAUNCH(policy)                                         \
   MoEGEMMLauncher_int<'R', 'R', policy, ScaleMode::kPerTensor>(                \
@@ -848,6 +842,8 @@ void moe_prefill_int_dpas_per_group_dispatch(
     throw std::runtime_error(
         "moe_prefill_int_dpas(per-group): failed to allocate atomic buffer");
   }
+  // Initialize the persistent scheduler counter before kernel launch.
+  q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();
 
 #define ARK_DPAS_INT_PG_LAUNCH_SYM(policy)                                     \
   MoEGEMMLauncher_int<'R', 'C', policy, ScaleMode::kPerGroup>(                 \
diff --git a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_s4_dpas.hpp b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_s4_dpas.hpp
index fad8d87eb..f4befa896 100644
--- a/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_s4_dpas.hpp
+++ b/auto_round_extension/ark/auto_round_kernel/wrapper/include/sycl_tla_moe_prefill_s4_dpas.hpp
@@ -423,14 +423,6 @@ CUTE_DEVICE void MoEGEMM_s4(const ElementA* Activations,
   int group_range = item.get_group_range(1);
   int local_id = item.get_local_linear_id();
 
-  if (group_id == 0 && local_id == 0) {
-    auto atm = sycl::atomic_ref<int, sycl::memory_order::relaxed,
-                                sycl::memory_scope::device,
-                                sycl::access::address_space::global_space>(
-        atomic_buffer[0]);
-    atm.store(0);
-  }
-
   int pre_rows = 0;
   int pre_tiles = 0;
 
@@ -642,6 +634,10 @@ void moe_prefill_s4_dpas_per_group_dispatch(
     throw std::runtime_error(
         "moe_prefill_s4_dpas(per-group): failed to allocate atomic buffer");
   }
+  // Initialize the persistent scheduler counter on the host before launch.
+  // Doing this in-kernel without a grid-wide barrier can race with early
+  // work-groups that already execute atomicAdd on the same counter.
+  q->memset(atomic_buffer, 0, sizeof(int32_t)).wait();
 
 #define ARK_DPAS_S4_PG_LAUNCH_SYM(policy)                                      \
   MoEGEMMLauncher_s4<'R', 'C', policy>(                                        \