From 441d6acb1657a1f3f290ef7e0112202e6aa52c3b Mon Sep 17 00:00:00 2001 From: "denghaodong.dhd" Date: Mon, 17 Aug 2026 20:12:13 +0800 Subject: [PATCH] fix(attention): remove stale barrier_O arrive in empty-tile fast path MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit In BatchPrefillWithKVCacheKernel, when a work tile hits num_kv_tiles <= 0 and exits early via store_zero + continue, the epilogue barrier parity is now handled internally by store_zero, so the manual barrier_O arrive from warp NUM_WARPS-1 is no longer needed — and can cause a spurious wake when the epilogue protocol has already moved on. Co-Authored-By: Claude --- include/flashinfer/attention/hopper/prefill_sm90.cuh | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/include/flashinfer/attention/hopper/prefill_sm90.cuh b/include/flashinfer/attention/hopper/prefill_sm90.cuh index 242a712352..5b84a57376 100644 --- a/include/flashinfer/attention/hopper/prefill_sm90.cuh +++ b/include/flashinfer/attention/hopper/prefill_sm90.cuh @@ -243,17 +243,6 @@ __global__ void __launch_bounds__(Ktraits::NUM_WARPS* cutlass::NumThreadsPerWarp int num_kv_tiles = collective_mainloop.get_num_kv_tiles(mainloop_params, q_tile_idx, qo_len, kv_len, batch_idx); if (num_kv_tiles <= 0) { // We exit early and write 0 to gO and -inf to gLSE. - // Preserve producer/consumer barrier parity when skipping this tile. - // A following visible tile would otherwise wait on a missing arrival. - if (work_idx != 0) { - int lane_predicate = cute::elect_one_sync(); - if (cutlass::canonical_warp_idx_sync() == Ktraits::NUM_WARPS - 1 && lane_predicate) { -#pragma unroll - for (uint32_t cta_id = 0; cta_id < 1; ++cta_id) { - shared_storage.barrier_O.arrive(cta_id, lane_predicate); - } - } - } collective_epilogue.store_zero(epilogue_params, shared_storage, threadIdx.x - NUM_COPY_THREADS, block_coord); continue;