[Bwd,Sm100] Add fence_view_async_shared before LSE release

tridao · tridao · commit d91ea94fb745 · 2026-03-06T00:00:39.000+07:00
diff --git a/flash_attn/cute/flash_bwd_sm100.py b/flash_attn/cute/flash_bwd_sm100.py
@@ -3133,6 +3133,7 @@ def compute_loop(
                     )
 
                 cute.arch.fence_view_async_tmem_store()
+                cute.arch.fence_view_async_shared()
                 self.compute_sync_barrier.arrive_and_wait()
                 if const_expr(not self.tile_hdim == 192):
                     # Signal tmem store P completion with pipeline_S_P

Original file line number	Diff line number	Diff line change
`@@ -3133,6 +3133,7 @@ def compute_loop(`
`3133`	`3133`	`)`
`3134`	`3134`
`3135`	`3135`	`cute.arch.fence_view_async_tmem_store()`
	`3136`	`+ cute.arch.fence_view_async_shared()`
`3136`	`3137`	`self.compute_sync_barrier.arrive_and_wait()`
`3137`	`3138`	`if const_expr(not self.tile_hdim == 192):`
`3138`	`3139`	`# Signal tmem store P completion with pipeline_S_P`