fix: fence_proxy for async global
This commit is contained in:
@@ -394,7 +394,8 @@ class FmhaKernel:
|
||||
gP_flat = cute.make_tensor(tCgP.iterator, cute.coalesce(tCgP.layout))
|
||||
for idx in cutlass.range(cute.size(rP_flat), vectorize=True):
|
||||
gP_flat[idx] = rP_flat[idx]
|
||||
cute.arch.fence_proxy("async.global", space="cta")
|
||||
# Ensure gP writes are visible before TMA reads
|
||||
cute.arch.fence_proxy("async", space="cta")
|
||||
cute.arch.fence_proxy("async.shared", space="cta")
|
||||
if kt > 0:
|
||||
for i in range(n_corr_tiles):
|
||||
|
||||
Reference in New Issue
Block a user