Stage C: integrate example3 multi-tile fixes into unit test
- Combined K+V barrier (one acquire per kt, kvh.count == kt) - O rescale for kt > 0 (online softmax O correction) - final_o_bar sync (MMA signals before producer_tail) - s_k as constructor param (compile-time for V layout) - kv_tx_bytes covers both K and V transfers - Test covers n=128, 256, 512, 1024
This commit is contained in:
@@ -1,8 +1,26 @@
|
||||
"""
|
||||
FMHA v3 Stage-C: Real softmax + O normalization.
|
||||
Builds on the 12w identity-softmax test by replacing identity softmax with
|
||||
online softmax (row_max, exp2 scaling, P store) and adding O normalization
|
||||
by row_sum before the epilogue writes to GMEM.
|
||||
FMHA v3 Stage-C: Real online softmax with full multi-tile support.
|
||||
|
||||
Integrated from:
|
||||
- test_fmha_v3.py (Stage A+B: QK→identity softmax→PV)
|
||||
- test_fmha_v3_stage_c_full.py (Stage C: single-tile real softmax)
|
||||
- fmha_v3_stage_c_example3.py (Multi-tile: combined K+V barrier, O rescale, final_o_bar)
|
||||
|
||||
Architecture (6-warp, 192 threads):
|
||||
Warps 0-3: Softmax + Epilogue (row_max, exp2, P store, O rescale, O normalize, TMA store)
|
||||
Warp 4: MMA (QK GEMM, PV GEMM)
|
||||
Warp 5: TMA (Q, K, V load)
|
||||
|
||||
Multi-tile key changes vs single-tile:
|
||||
1. Combined K+V barrier: one acquire_and_advance per kt, both K and V share kvh.barrier.
|
||||
kvh.count == kt naturally — no interleaving, no Python int in TMA coordinates.
|
||||
2. kv_tx_bytes covers BOTH K and V transfers per stage.
|
||||
3. V FMHA layout uses s_k as the sequence dimension (compile-time, not hardcoded 128).
|
||||
4. MMA: same slot index (kvh.index) for both K (QK) and V (PV). Release after PV.
|
||||
5. O rescale for kt > 0: exp2((old_max - new_max) * scale_log2) on O in TMEM.
|
||||
6. final_o_bar: MMA arrives between producer_commit and producer_tail;
|
||||
softmax arrives_and_wait before reading O for final normalize.
|
||||
7. s_k is a constructor parameter — each sequence length requires its own compiled kernel.
|
||||
"""
|
||||
import torch, cutlass, cutlass.cute as cute, cutlass.utils as utils, cutlass.pipeline as pipeline
|
||||
from cutlass.cute.nvgpu import cpasync, tcgen05
|
||||
@@ -15,8 +33,10 @@ import math
|
||||
|
||||
HEAD_DIM = 64
|
||||
|
||||
|
||||
class FmhaV3StageC:
|
||||
def __init__(self, s_k=128, scale_softmax=None):
|
||||
# s_k MUST equal actual sequence length n (compile-time constant for V layout).
|
||||
self.s_k = s_k
|
||||
self.acc_dtype = Float32; self.qk_acc_dtype = Float32
|
||||
self.q_dtype = BFloat16; self.o_dtype = BFloat16; self.c_dtype = BFloat16
|
||||
@@ -49,24 +69,24 @@ class FmhaV3StageC:
|
||||
pv_thr = pv_mma.get_slice(0); pv_as = pv_thr.partition_shape_C(self.pv_mma_tiler[:2])
|
||||
tOtO = pv_thr.make_fragment_C(pv_as)
|
||||
self.tmem_s0_offset = 0; self.tmem_p0_offset = 32
|
||||
# P occupies [tmem_p0_offset, tmem_p0_offset + p_cols_fp32)
|
||||
# S occupies [0, qk_mma_tiler[1]) = [0, 128)
|
||||
# O must NOT overlap P. Place O after max(S end, P end), aligned to 32.
|
||||
p_cols_fp32 = self.pv_mma_tiler[2] * self.q_dtype.width // self.qk_acc_dtype.width
|
||||
p_end = self.tmem_p0_offset + p_cols_fp32 # 32 + 64 = 96
|
||||
s_cols = self.qk_mma_tiler[1] # 128
|
||||
o_after = max(s_cols, p_end) # 128
|
||||
self.tmem_o0_offset = ((o_after + 31) // 32) * 32 # align to 32 = 128
|
||||
o_cols = find_tmem_tensor_col_offset(tOtO) # footprint of O
|
||||
p_end = self.tmem_p0_offset + p_cols_fp32
|
||||
s_cols = self.qk_mma_tiler[1]
|
||||
o_after = max(s_cols, p_end)
|
||||
self.tmem_o0_offset = ((o_after + 31) // 32) * 32
|
||||
o_cols = find_tmem_tensor_col_offset(tOtO)
|
||||
total = self.tmem_o0_offset + o_cols
|
||||
# Must be multiple of 32 AND power of 2
|
||||
self.num_tmem_alloc_cols = 1
|
||||
while self.num_tmem_alloc_cols < total:
|
||||
self.num_tmem_alloc_cols *= 2
|
||||
cta = cute.size(qk_mma.thr_id.shape)
|
||||
q_s = cute.slice_(self.q_smem_s,(None,None,None,0)); k_s = cute.slice_(self.k_smem_s,(None,None,None,0))
|
||||
q_s = cute.slice_(self.q_smem_s,(None,None,None,0))
|
||||
k_s = cute.slice_(self.k_smem_s,(None,None,None,0))
|
||||
v_s = cute.slice_(self.v_smem_s,(None,None,None,0))
|
||||
self.q_tx_bytes = cute.size_in_bytes(self.q_dtype, q_s) * cta
|
||||
self.kv_tx_bytes = cute.size_in_bytes(self.q_dtype, k_s) * cta
|
||||
# Combined barrier: tx_count covers BOTH K and V transfers per acquire.
|
||||
self.kv_tx_bytes = (cute.size_in_bytes(self.q_dtype, k_s) +
|
||||
cute.size_in_bytes(self.q_dtype, v_s)) * cta
|
||||
|
||||
@cute.jit
|
||||
def __call__(self, q, k, v, c, stream):
|
||||
@@ -74,6 +94,7 @@ class FmhaV3StageC:
|
||||
self.a_major = LayoutEnum.from_tensor(q).mma_major_mode()
|
||||
self.b_major = LayoutEnum.from_tensor(k).mma_major_mode()
|
||||
# FMHA-style V: reconstruct as (HEAD_DIM, s_k, 1) MN-major
|
||||
# s_k is compile-time — must match actual sequence length n.
|
||||
v_fmha = cute.make_tensor(
|
||||
v.iterator,
|
||||
cute.make_layout(
|
||||
@@ -111,9 +132,17 @@ class FmhaV3StageC:
|
||||
smem = utils.SmemAllocator(); st = smem.allocate(SS)
|
||||
|
||||
qp,qc = pipeline.PipelineTmaUmma.create(barrier_storage=st.q_bar.data_ptr(),num_stages=self.q_stage,producer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread),consumer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread,1),tx_count=self.q_tx_bytes,cta_layout_vmnk=cl_vmnk,defer_sync=True).make_participants()
|
||||
# Combined K+V pipeline: each stage carries BOTH K and V loaded together.
|
||||
# One acquire per kt → kvh.count == kt (pipeline state value, accepted by cute.copy).
|
||||
# No interleaving problem. No Python int in TMA coordinates.
|
||||
kvp,kvc = pipeline.PipelineTmaUmma.create(barrier_storage=st.kv_bar.data_ptr(),num_stages=self.kv_stage,producer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread),consumer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread,1),tx_count=self.kv_tx_bytes,cta_layout_vmnk=cl_vmnk,defer_sync=True).make_participants()
|
||||
s_prod,s_cons = pipeline.PipelineUmmaAsync.create(barrier_storage=st.s_bar.data_ptr(),num_stages=1,producer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread),consumer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread,32*len(self.epilogue_warp_id))).make_participants()
|
||||
softmax_done_bar = pipeline.NamedBarrier(barrier_id=3, num_threads=32 + 32*len(self.epilogue_warp_id))
|
||||
# Final-O sync: MMA arrives between producer_commit and producer_tail;
|
||||
# softmax arrives_and_waits before reading O for the final normalize.
|
||||
# This prevents softmax from racing MMA's PV[N-1] and dividing a
|
||||
# partially-accumulated O by row_sum.
|
||||
final_o_bar = pipeline.NamedBarrier(barrier_id=4, num_threads=32 + 32*len(self.epilogue_warp_id))
|
||||
acc_pipe = pipeline.PipelineUmmaAsync.create(barrier_storage=st.acc_bar.data_ptr(),num_stages=self.num_acc_stage,producer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread),consumer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread,len(self.epilogue_warp_id)),cta_layout_vmnk=cl_vmnk,defer_sync=True)
|
||||
tmem_bar = pipeline.NamedBarrier(barrier_id=2,num_threads=32*len((self.mma_warp_id,*self.epilogue_warp_id)))
|
||||
tmem = utils.TmemAllocator(st.holding.ptr,barrier_for_retrieve=tmem_bar,allocator_warp_id=self.epilogue_warp_id[0],is_two_cta=cute.size(qk_mma.thr_id.shape)==2,two_cta_tmem_dealloc_mbar_ptr=st.tmem_dealloc.ptr)
|
||||
@@ -162,22 +191,25 @@ class FmhaV3StageC:
|
||||
tCtO_fake = pv_mma.make_fragment_C(cute.append(pv_as, self.num_acc_stage))
|
||||
pipeline.pipeline_init_wait(cluster_shape_mn=cl_vmnk)
|
||||
|
||||
# TMA LOAD
|
||||
# ===== TMA LOAD warp =====
|
||||
# One acquire per kt; K and V both target kvh.barrier. kvh.count == kt.
|
||||
if warp_idx == self.tma_warp_id:
|
||||
qp.reset(); qh = qp.acquire_and_advance()
|
||||
cute.copy(tma_q,tAgQ[(None,qh.count)],tAsQ[(None,qh.index)],tma_bar_ptr=qh.barrier)
|
||||
cute.copy(tma_q, tAgQ[(None, qh.count)], tAsQ[(None, qh.index)], tma_bar_ptr=qh.barrier)
|
||||
qp.tail()
|
||||
kvp.reset(); pk = kvp.try_acquire()
|
||||
for kt in cutlass.range(n_kv_tiles,unroll=1):
|
||||
kh = kvp.acquire_and_advance(pk)
|
||||
cute.copy(tma_k,tBgK[(None,kh.count)],tBsK[(None,kh.index)],tma_bar_ptr=kh.barrier)
|
||||
pk = cutlass.Boolean(1)
|
||||
vh = kvp.acquire_and_advance(pk)
|
||||
cute.copy(tma_v,tVgV[(None,vh.count)],tVsV[(None,vh.index)],tma_bar_ptr=vh.barrier)
|
||||
for kt in cutlass.range(n_kv_tiles, unroll=1):
|
||||
kvh = kvp.acquire_and_advance(pk)
|
||||
# Both transfers decrement the same barrier's tx_count.
|
||||
# kvh.count is a pipeline-state Int32 (the form cute.copy accepts).
|
||||
cute.copy(tma_k, tBgK[(None, kvh.count)], tBsK[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
|
||||
cute.copy(tma_v, tVgV[(None, kvh.count)], tVsV[(None, kvh.index)], tma_bar_ptr=kvh.barrier)
|
||||
pk = cutlass.Boolean(1)
|
||||
kvp.tail()
|
||||
|
||||
# MMA
|
||||
# ===== MMA warp =====
|
||||
# One wait per kt; same slot index used for both K (QK) and V (PV).
|
||||
# Release happens AFTER PV — combined slot stays held across QK+PV.
|
||||
if warp_idx == self.mma_warp_id:
|
||||
tmem.wait_for_alloc()
|
||||
qc.reset(); qh = qc.wait_and_advance(); qh.release()
|
||||
@@ -185,46 +217,48 @@ class FmhaV3StageC:
|
||||
acc_st = pipeline.make_pipeline_state(pipeline.PipelineUserType.Producer, self.num_acc_stage)
|
||||
acc_pipe.producer_acquire(acc_st)
|
||||
for kt in range(n_kv_tiles):
|
||||
kh = kvc.wait_and_advance(pk); pk = cutlass.Boolean(1)
|
||||
kvh = kvc.wait_and_advance(pk); pk = cutlass.Boolean(1)
|
||||
sh = s_prod.acquire_and_advance()
|
||||
qk_mma.set(tcgen05.Field.ACCUMULATE, False)
|
||||
for kb in cutlass.range(cute.size(tCrQ,mode=[2]), unroll_full=True):
|
||||
cute.gemm(qk_mma, tStS0, tCrQ[(None,None,kb,0)], tCrK[(None,None,kb,kh.index)], tStS0)
|
||||
for kb in cutlass.range(cute.size(tCrQ, mode=[2]), unroll_full=True):
|
||||
cute.gemm(qk_mma, tStS0, tCrQ[(None,None,kb,0)], tCrK[(None,None,kb,kvh.index)], tStS0)
|
||||
qk_mma.set(tcgen05.Field.ACCUMULATE, True)
|
||||
cute.arch.fence_view_async_tmem_store()
|
||||
sh.commit(); kh.release()
|
||||
sh.commit()
|
||||
softmax_done_bar.arrive_and_wait()
|
||||
vh = kvc.wait_and_advance(pk); pk = cutlass.Boolean(1)
|
||||
pv_mma.set(tcgen05.Field.ACCUMULATE, kt != 0)
|
||||
for kb in cutlass.range(cute.size(tOrP0,mode=[2]), unroll_full=True):
|
||||
cute.gemm(pv_mma, tOtO0, tOrP0[(None,None,kb)], tCrV[(None,None,kb,vh.index)], tOtO0)
|
||||
for kb in cutlass.range(cute.size(tOrP0, mode=[2]), unroll_full=True):
|
||||
cute.gemm(pv_mma, tOtO0, tOrP0[(None,None,kb)], tCrV[(None,None,kb,kvh.index)], tOtO0)
|
||||
pv_mma.set(tcgen05.Field.ACCUMULATE, True)
|
||||
cute.arch.fence_view_async_tmem_store()
|
||||
vh.release()
|
||||
kvh.release()
|
||||
acc_pipe.producer_commit(acc_st); acc_st.advance()
|
||||
# Signal softmax FIRST so it can run normalize + epilogue. Then
|
||||
# wait for the epilogue's consumer-release in producer_tail.
|
||||
# Reverse order deadlocks: producer_tail blocks waiting for
|
||||
# consumer release; softmax blocks at final_o_bar waiting for
|
||||
# MMA arrive; the epilogue (which does the release) is gated
|
||||
# behind softmax's final_o_bar wait. Cycle.
|
||||
final_o_bar.arrive()
|
||||
acc_pipe.producer_tail(acc_st)
|
||||
|
||||
# SOFTMAX + EPILOGUE (warps 0-3)
|
||||
# Step 1: Real online softmax: load S, compute row_max, exp2 scale, store P
|
||||
# Step 2: After all KV tiles, normalize O in TMEM by row_sum
|
||||
# Step 3: Epilogue writes normalized O from TMEM to GMEM
|
||||
# ===== SOFTMAX + EPILOGUE warps =====
|
||||
if warp_idx < self.mma_warp_id:
|
||||
tmem.allocate(self.num_tmem_alloc_cols)
|
||||
tmem.wait_for_alloc()
|
||||
tmem_ptr = tmem.retrieve_ptr(self.qk_acc_dtype)
|
||||
sfw_idx = tidx % (32 * len(self.epilogue_warp_id))
|
||||
|
||||
# --- S load (QK C-fragment layout) ---
|
||||
# S load (QK C-fragment layout)
|
||||
tmem_load_atom = cute.make_copy_atom(tcgen05.copy.Ld32x32bOp(tcgen05.copy.Repetition(32)), self.qk_acc_dtype)
|
||||
tiled_tmem_load = tcgen05.make_tmem_copy(tmem_load_atom, tStS0)
|
||||
thr_load = tiled_tmem_load.get_slice(sfw_idx)
|
||||
tTMEM_LOADtS = thr_load.partition_S(tStS0)
|
||||
|
||||
cS = cute.make_identity_tensor((self.qk_mma_tiler[0], self.qk_mma_tiler[1]))
|
||||
tScS = qk_thr.partition_C(cS)
|
||||
tTMEM_LOADcS = thr_load.partition_D(tScS)
|
||||
|
||||
# --- P store (QK C-fragment layout composition, FMHA pattern) ---
|
||||
# P store (QK C-fragment layout composition, FMHA pattern)
|
||||
p_cols_fp32 = self.pv_mma_tiler[2] * self.q_dtype.width // self.qk_acc_dtype.width
|
||||
tStP_layout = cute.composition(tStS.layout, cute.make_layout((self.pv_mma_tiler[0], p_cols_fp32)))
|
||||
tStP0 = cute.make_tensor(tStS.iterator + self.tmem_p0_offset, tStP_layout)
|
||||
@@ -236,8 +270,7 @@ class FmhaV3StageC:
|
||||
tScP = cute.make_tensor(tScS.iterator, tScP_layout)
|
||||
tTMEM_STOREcP = thr_store.partition_S(tScP)
|
||||
|
||||
# --- Online softmax loop ---
|
||||
# --- O rescale + normalize setup (before softmax loop) ---
|
||||
# O rescale / normalize path
|
||||
cO = cute.make_identity_tensor((self.pv_mma_tiler[0], self.pv_mma_tiler[1]))
|
||||
tOcO = pv_thr.partition_C(cO)
|
||||
corr_tile_size = 16
|
||||
@@ -255,18 +288,21 @@ class FmhaV3StageC:
|
||||
tTMEM_LOAD_OcO = thr_load_o.partition_D(tOcO_i)
|
||||
tTMEM_STORE_OtO = thr_store_o.partition_D(tOtO_i)
|
||||
|
||||
row_max = -Float32.inf; row_sum = Float32(0.0)
|
||||
o_col_tiles = self.pv_mma_tiler[1] // corr_tile_size
|
||||
|
||||
row_max = -Float32.inf
|
||||
row_sum = Float32(0.0)
|
||||
scale_log2 = Float32(self.scale_softmax_log2)
|
||||
|
||||
for kt in range(n_kv_tiles):
|
||||
si_handle = s_cons.wait_and_advance()
|
||||
|
||||
# Load S from TMEM
|
||||
# Load S[kt]
|
||||
tTMEM_LOADrS = cute.make_rmem_tensor(tTMEM_LOADcS.shape, self.qk_acc_dtype)
|
||||
cute.copy(tiled_tmem_load, tTMEM_LOADtS, tTMEM_LOADrS)
|
||||
cute.arch.fence_view_async_tmem_load()
|
||||
|
||||
# Compute row_max from S (element-wise, before exp2)
|
||||
# Pass 1: update row_max
|
||||
old_row_max = row_max
|
||||
frg_cnt = 4
|
||||
frg_tile = cute.size(tTMEM_LOADrS) // frg_cnt
|
||||
@@ -276,15 +312,18 @@ class FmhaV3StageC:
|
||||
row_max = cute.arch.fmax(row_max, tTMEM_LOADrS_frg[k, j] * scale_log2)
|
||||
|
||||
row_max_safe = row_max
|
||||
if row_max == -cutlass.Float32.inf: row_max_safe = Float32(0.0)
|
||||
if row_max == -cutlass.Float32.inf:
|
||||
row_max_safe = Float32(0.0)
|
||||
|
||||
# Scale existing row_sum
|
||||
# acc_scale for both row_sum rescale and O rescale.
|
||||
acc_scale_ = scale_log2 * (old_row_max - row_max_safe)
|
||||
acc_scale = cute.math.exp2(acc_scale_, fastmath=True)
|
||||
if old_row_max == -cutlass.Float32.inf: acc_scale = Float32(0.0)
|
||||
if old_row_max == -cutlass.Float32.inf:
|
||||
acc_scale = Float32(0.0)
|
||||
row_sum *= acc_scale
|
||||
|
||||
# Second pass: compute P and accumulate row_sum
|
||||
# Pass 2: P = exp2((S - new_max) * log2), accumulate row_sum,
|
||||
# store BF16 P through the FP32-backed register bridge.
|
||||
rP_words = cute.make_rmem_tensor(tTMEM_STOREcP.shape, self.qk_acc_dtype)
|
||||
rP_bf16 = cute.make_tensor(cute.recast_ptr(rP_words.iterator, dtype=self.q_dtype), tTMEM_LOADrS.layout)
|
||||
minus_row_max_scale = (Float32(0.0) - row_max_safe) * scale_log2
|
||||
@@ -301,38 +340,54 @@ class FmhaV3StageC:
|
||||
cute.copy(tiled_tmem_store, rP_words, tTMEM_STOREtP)
|
||||
cute.arch.fence_view_async_tmem_store()
|
||||
|
||||
# O rescale: if kt > 0, rescale O in TMEM by exp2((old_max - new_max) * scale_log2)
|
||||
# TEMPORARILY DISABLED FOR DEBUGGING
|
||||
# if kt > 0:
|
||||
# corr_scale = cute.math.exp2(old_row_max - row_max_safe, fastmath=True)
|
||||
# o_col_tiles = self.pv_mma_tiler[1] // corr_tile_size
|
||||
# for i in range(o_col_tiles):
|
||||
# tTMEM_LOAD_O_i = cute.make_tensor(tTMEM_LOAD_OtO.iterator + i * corr_tile_size, tTMEM_LOAD_OtO.layout)
|
||||
# tTMEM_STORE_O_i = cute.make_tensor(tTMEM_STORE_OtO.iterator + i * corr_tile_size, tTMEM_STORE_OtO.layout)
|
||||
# tTMrO = cute.make_rmem_tensor(tTMEM_LOAD_OcO.shape, self.acc_dtype)
|
||||
# cute.copy(tiled_tmem_load_o, tTMEM_LOAD_O_i, tTMrO)
|
||||
# for k in cutlass.range(cute.size(tTMrO), vectorize=True):
|
||||
# tTMrO[k] = tTMrO[k] * corr_scale
|
||||
# cute.copy(tiled_tmem_store_o, tTMrO, tTMEM_STORE_O_i)
|
||||
# cute.arch.fence_view_async_tmem_store()
|
||||
# O rescale for kt > 0. Reads O written by MMA's PV[kt-1];
|
||||
# visibility is provided by s_cons.wait_and_advance above
|
||||
# (acquires on MMA's S[kt] commit, which orders PV[kt-1] before).
|
||||
if kt > 0:
|
||||
for i in range(o_col_tiles):
|
||||
tTMEM_LOAD_O_i = cute.make_tensor(
|
||||
tTMEM_LOAD_OtO.iterator + i * corr_tile_size,
|
||||
tTMEM_LOAD_OtO.layout,
|
||||
)
|
||||
tTMEM_STORE_O_i = cute.make_tensor(
|
||||
tTMEM_STORE_OtO.iterator + i * corr_tile_size,
|
||||
tTMEM_STORE_OtO.layout,
|
||||
)
|
||||
tTMrO = cute.make_rmem_tensor(tTMEM_LOAD_OcO.shape, self.acc_dtype)
|
||||
cute.copy(tiled_tmem_load_o, tTMEM_LOAD_O_i, tTMrO)
|
||||
cute.arch.fence_view_async_tmem_load()
|
||||
for k in cutlass.range(cute.size(tTMrO), vectorize=True):
|
||||
tTMrO[k] = tTMrO[k] * acc_scale
|
||||
cute.copy(tiled_tmem_store_o, tTMrO, tTMEM_STORE_O_i)
|
||||
cute.arch.fence_view_async_tmem_store()
|
||||
|
||||
si_handle.release()
|
||||
softmax_done_bar.arrive()
|
||||
|
||||
# --- Final O normalization by row_sum ---
|
||||
# Wait for MMA's last PV to commit before reading O for normalize.
|
||||
# Without this barrier softmax can race MMA's PV[N-1].
|
||||
final_o_bar.arrive_and_wait()
|
||||
|
||||
# Final O = O / row_sum
|
||||
inv_row_sum = Float32(1.0) / row_sum
|
||||
o_col_tiles = self.pv_mma_tiler[1] // corr_tile_size
|
||||
for i in range(o_col_tiles):
|
||||
tTMEM_LOAD_O_i = cute.make_tensor(tTMEM_LOAD_OtO.iterator + i * corr_tile_size, tTMEM_LOAD_OtO.layout)
|
||||
tTMEM_STORE_O_i = cute.make_tensor(tTMEM_STORE_OtO.iterator + i * corr_tile_size, tTMEM_STORE_OtO.layout)
|
||||
tTMEM_LOAD_O_i = cute.make_tensor(
|
||||
tTMEM_LOAD_OtO.iterator + i * corr_tile_size,
|
||||
tTMEM_LOAD_OtO.layout,
|
||||
)
|
||||
tTMEM_STORE_O_i = cute.make_tensor(
|
||||
tTMEM_STORE_OtO.iterator + i * corr_tile_size,
|
||||
tTMEM_STORE_OtO.layout,
|
||||
)
|
||||
tTMrO = cute.make_rmem_tensor(tTMEM_LOAD_OcO.shape, self.acc_dtype)
|
||||
cute.copy(tiled_tmem_load_o, tTMEM_LOAD_O_i, tTMrO)
|
||||
cute.arch.fence_view_async_tmem_load()
|
||||
for k in cutlass.range(cute.size(tTMrO), vectorize=True):
|
||||
tTMrO[k] = tTMrO[k] * inv_row_sum
|
||||
cute.copy(tiled_tmem_store_o, tTMrO, tTMEM_STORE_O_i)
|
||||
cute.arch.fence_view_async_tmem_store()
|
||||
|
||||
# --- Epilogue: write O from TMEM to GMEM ---
|
||||
# Epilogue: TMEM -> SMEM -> GMEM via TMA store
|
||||
tCtO_base = cute.make_tensor(tmem_ptr + self.tmem_o0_offset, tCtO_fake.layout)
|
||||
acc_cons_st = pipeline.make_pipeline_state(pipeline.PipelineUserType.Consumer, self.num_acc_stage)
|
||||
c_grp = pipeline.CooperativeGroup(pipeline.Agent.Thread, 32 * len(self.epilogue_warp_id))
|
||||
@@ -345,39 +400,52 @@ class FmhaV3StageC:
|
||||
|
||||
def test():
|
||||
torch.manual_seed(42)
|
||||
# Test single tile (n=128) and multi-tile (n=256)
|
||||
for n in [128, 256]:
|
||||
for seed in [42]:
|
||||
torch.manual_seed(seed)
|
||||
m, hd = 128, HEAD_DIM
|
||||
q = torch.randn(m, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
k = torch.randn(n, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
v = torch.randn(n, hd, dtype=torch.bfloat16, device='cuda')
|
||||
v_kernel = v.unsqueeze(-1)
|
||||
c = torch.zeros(m, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
qf = q[:,:,0].float(); kf = k[:,:,0].float()
|
||||
scale = 1.0 / math.sqrt(hd)
|
||||
attn = qf @ kf.T * scale
|
||||
attn = torch.softmax(attn, dim=-1)
|
||||
ref = attn @ v.float()
|
||||
mQ = ct.from_dlpack(q).mark_layout_dynamic(leading_dim=ct.get_leading_dim(q))
|
||||
mK = ct.from_dlpack(k).mark_layout_dynamic(leading_dim=ct.get_leading_dim(k))
|
||||
mV = ct.from_dlpack(v_kernel).mark_layout_dynamic(leading_dim=ct.get_leading_dim(v_kernel))
|
||||
mC = ct.from_dlpack(c).mark_layout_dynamic(leading_dim=ct.get_leading_dim(c))
|
||||
stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)
|
||||
kernel = FmhaV3StageC(s_k=128)
|
||||
if seed == 42:
|
||||
print(f'seed={seed}: Compiling...', flush=True)
|
||||
compiled = cute.compile(kernel, mQ, mK, mV, mC, stream)
|
||||
if seed == 42:
|
||||
print(f'tmem_offsets: s0={kernel.tmem_s0_offset} p0={kernel.tmem_p0_offset} o0={kernel.tmem_o0_offset} alloc={kernel.num_tmem_alloc_cols}', flush=True)
|
||||
compiled(mQ, mK, mV, mC, stream)
|
||||
torch.cuda.synchronize()
|
||||
out = c[:,:,0].float()
|
||||
cos = torch.nn.functional.cosine_similarity(out.flatten().unsqueeze(0), ref.flatten().unsqueeze(0)).item()
|
||||
print(f'FMHA Stage-C n={n} seed={seed}: cosine {cos:.6f} {"PASS" if cos >= 0.99 else "FAIL"}')
|
||||
if cos < 0.99:
|
||||
print(f' out[0,:4]={out[0,:4].tolist()} ref[0,:4]={ref[0,:4].tolist()}')
|
||||
for n in [128, 256, 512, 1024]:
|
||||
torch.manual_seed(42)
|
||||
m, hd = 128, HEAD_DIM
|
||||
q = torch.randn(m, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
k = torch.randn(n, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
v = torch.randn(n, hd, dtype=torch.bfloat16, device='cuda')
|
||||
v_kernel = v.unsqueeze(-1)
|
||||
c = torch.zeros(m, hd, 1, dtype=torch.bfloat16, device='cuda')
|
||||
|
||||
# PyTorch reference
|
||||
qf = q[:, :, 0].float()
|
||||
kf = k[:, :, 0].float()
|
||||
scale = 1.0 / math.sqrt(hd)
|
||||
attn = qf @ kf.T * scale
|
||||
attn = torch.softmax(attn, dim=-1)
|
||||
ref = attn @ v.float()
|
||||
|
||||
mQ = ct.from_dlpack(q).mark_layout_dynamic(leading_dim=ct.get_leading_dim(q))
|
||||
mK = ct.from_dlpack(k).mark_layout_dynamic(leading_dim=ct.get_leading_dim(k))
|
||||
mV = ct.from_dlpack(v_kernel).mark_layout_dynamic(leading_dim=ct.get_leading_dim(v_kernel))
|
||||
mC = ct.from_dlpack(c).mark_layout_dynamic(leading_dim=ct.get_leading_dim(c))
|
||||
stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)
|
||||
|
||||
# Each n requires its own compiled kernel (s_k is compile-time).
|
||||
kernel = FmhaV3StageC(s_k=n)
|
||||
print(f'n={n}: Compiling...', flush=True)
|
||||
compiled = cute.compile(kernel, mQ, mK, mV, mC, stream)
|
||||
print(f'n={n}: tmem s0={kernel.tmem_s0_offset} p0={kernel.tmem_p0_offset} '
|
||||
f'o0={kernel.tmem_o0_offset} alloc={kernel.num_tmem_alloc_cols} '
|
||||
f'kv_tx_bytes={kernel.kv_tx_bytes}', flush=True)
|
||||
compiled(mQ, mK, mV, mC, stream)
|
||||
torch.cuda.synchronize()
|
||||
|
||||
out = c[:, :, 0].float()
|
||||
cos = torch.nn.functional.cosine_similarity(
|
||||
out.flatten().unsqueeze(0), ref.flatten().unsqueeze(0)
|
||||
).item()
|
||||
max_abs = (out - ref).abs().max().item()
|
||||
n_tiles = n // 128
|
||||
print(f'FMHA Stage-C n={n} ({n_tiles} kv tiles): '
|
||||
f'cos {cos:.6f} max_abs {max_abs:.4f} '
|
||||
f'{"PASS" if cos >= 0.99 else "FAIL"}')
|
||||
if cos < 0.99:
|
||||
print(f' out[0,:4]={out[0,:4].tolist()}')
|
||||
print(f' ref[0,:4]={ref[0,:4].tolist()}')
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
test()
|
||||
|
||||
Reference in New Issue
Block a user