Skip to content

[fea&fix] reshape support && named barrier memalloc && tle.gpu.copy with mask && barrier arrival#840

Draft
Kafka-Hatsune wants to merge 6 commits into
flagos-ai:mainfrom
Kafka-Hatsune:audit/tle-hopper-flash-attention
Draft

[fea&fix] reshape support && named barrier memalloc && tle.gpu.copy with mask && barrier arrival#840
Kafka-Hatsune wants to merge 6 commits into
flagos-ai:mainfrom
Kafka-Hatsune:audit/tle-hopper-flash-attention

Conversation

@Kafka-Hatsune

Copy link
Copy Markdown
Contributor

为FlagGems hopper flash_attn_func_varlen PR进行的FlagTree功能补充与修复。包含以下四点:

1. 零拷贝 buffered_tensor.reshape

# 原共享内存视图:[32, 512]
q_smem = ...

# 重新解释为 [32, 4, 128],底层共享内存不发生复制
q_packed = q_smem.reshape([32, 4, 128])

适用于 PackGQA 等场景,将同一块 shared memory 以 WGMMA 所需形状访问。

2. warp specialization 中保留 named barrier

ready = tle.gpu.alloc_barrier()

# ready 被捕获到 warp-specialized consumer partition
with tle.gpu.warp_specialize(...):
    tle.gpu.barrier_wait(ready)
    consume_shared_tile()

修复后,ready 仍降低为硬件 named barrier,不会因为跨 partition capture 而额外分配并初始化一个无用的 shared-memory mbarrier。

3. 带 mask 的普通 pointer copy

rows = tl.arange(0, BLOCK_M)[:, None]
dims = tl.arange(0, HEAD_DIM_PADDED)[None, :]
mask = (rows < valid_rows) & (dims < HEAD_DIM)

tle.gpu.copy(
    q_ptrs,
    q_smem,
    [BLOCK_M, HEAD_DIM_PADDED],
    mask=mask,
    other=0.0,
)

越界行和 padding head dimension 会在 shared memory 中补零,同时仍可被现有 lowering 优化为 cp.async

4. participant barrier arrival

q_ready = tle.gpu.alloc_barrier(
    arrive_count=1,
    arrival_mode="participant",
)

# producer partition 中每个参与线程各到达一次
tle.gpu.barrier_arrive(q_ready, phaseIdx=phase)

# consumer 等待所有 producer participant 完成发布
tle.gpu.barrier_wait(q_ready, phaseIdx=phase)

例如 producer partition 有 128 个线程时,lowering 会将 mbarrier expected arrival count 初始化为 128,而不是由单个 elected thread 代表整个 partition 到达。

对照默认模式:

barrier = tle.gpu.alloc_barrier(
    arrival_mode="elected",  # 默认值
)

# 由 partition 中的 elected thread 执行逻辑到达
tle.gpu.barrier_arrive(barrier, phaseIdx=phase)

其中 participant/elected 只决定 direct mbarrier 的发布方式,不替代 cp.async.wait_group、线程 rendezvous 或 WGMMA 所需的 proxy fence。

Expose mask/other on the existing normal pointer-tensor copy path so tail and padded-head-dimension tiles are fully initialized in shared memory. This is required by the current Hopper attention source and preserves the load/store pattern consumed by the local-pointer async-copy rewrite; descriptor/TMA copies continue to reject these options.
Make the direct mbarrier arrival strategy explicit instead of treating every phase-indexed barrier as a participant barrier. Barrier allocations default to elected lowering, while opt-in participant allocations scale initialization counts by the arriving partition width and lower release-fenced unit arrivals from each participant lane.

Reject participant mode for named and TMA transaction barriers, propagate the mode through frontend types and the IR builder, and update the existing constructor helper required by the new type signature.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant