From 0dbdc4f865f3f11ad94fb6f1c6a1a219ff2f996d Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 23 May 2026 03:52:47 +0000 Subject: [PATCH] fix: CopyAtomUniversalOp --- dsv4/kernels/attention/fmha.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dsv4/kernels/attention/fmha.py b/dsv4/kernels/attention/fmha.py index 39e73ebe..f3d5572b 100644 --- a/dsv4/kernels/attention/fmha.py +++ b/dsv4/kernels/attention/fmha.py @@ -200,8 +200,8 @@ class FmhaKernel: tScS = qk_thr.partition_C(cS) tTMEM_LOADcS = thr_load.partition_D(tScS) - # P → SMEM copy setup - p_copy_atom = cute.make_copy_atom(cute.nvgpu.copy.AutoCopyOp(), self.q_dtype) + # P → SMEM copy setup (use SMEM copy atom for register→SMEM) + p_copy_atom = cute.make_copy_atom(cute.CopyAtomUniversalOp(), self.q_dtype) tiled_p_copy = cute.make_tiled_copy(p_copy_atom, tCrP_smem.layout, tCrP_reg.layout, tidx) tPS_sP = tiled_p_copy.get_slice(sfw_idx).partition_D(tCrP_smem) tPS_rP = tiled_p_copy.get_slice(sfw_idx).partition_S(tCrP_reg)