Add missing async proxy fence

lw · facebook-github-bot · commit 47e629a122ef · 2026-04-15T02:43:49.000-07:00
Differential Revision: D100956644
diff --git a/fbgemm_gpu/experimental/hstu/src/hstu_hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp b/fbgemm_gpu/experimental/hstu/src/hstu_hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp
@@ -633,6 +633,7 @@ struct CollectiveMainloopFwd {
               flatten(sVt_divide(_, i, j, stage)));
         }
       }
+      cutlass::arch::fence_view_async_shared();
       cutlass::arch::NamedBarrier::sync(
           cutlass::NumThreadsPerWarpGroup,
           static_cast<int>(FwdNamedBarriers::ProducerWG));

Original file line number	Diff line number	Diff line change
`@@ -633,6 +633,7 @@ struct CollectiveMainloopFwd {`
`633`	`633`	`flatten(sVt_divide(_, i, j, stage)));`
`634`	`634`	`}`
`635`	`635`	`}`
	`636`	`+ cutlass::arch::fence_view_async_shared();`
`636`	`637`	`cutlass::arch::NamedBarrier::sync(`
`637`	`638`	`cutlass::NumThreadsPerWarpGroup,`
`638`	`639`	`static_cast<int>(FwdNamedBarriers::ProducerWG));`