[fea&fix] reshape support && named barrier memalloc && tle.gpu.copy with mask && barrier arrival#840
Draft
Kafka-Hatsune wants to merge 6 commits into
Draft
Conversation
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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
为FlagGems hopper flash_attn_func_varlen PR进行的FlagTree功能补充与修复。包含以下四点:
1. 零拷贝
buffered_tensor.reshape适用于 PackGQA 等场景,将同一块 shared memory 以 WGMMA 所需形状访问。
2. warp specialization 中保留 named barrier
修复后,
ready仍降低为硬件 named barrier,不会因为跨 partition capture 而额外分配并初始化一个无用的 shared-memory mbarrier。3. 带 mask 的普通 pointer copy
越界行和 padding head dimension 会在 shared memory 中补零,同时仍可被现有 lowering 优化为
cp.async。4. participant barrier arrival
例如 producer partition 有 128 个线程时,lowering 会将 mbarrier expected arrival count 初始化为 128,而不是由单个 elected thread 代表整个 partition 到达。
对照默认模式:
其中 participant/elected 只决定 direct mbarrier 的发布方式,不替代
cp.async.wait_group、线程 rendezvous 或 WGMMA 所需的 proxy fence。