[CK_TILE][FMHA] Add sparse attention VSA#3341

Merged

poyenc merged 34 commits intodevelopfrom

sparse_attention_VSA

Jan 30, 2026

Contributor

jiangyon-amd commented Dec 3, 2025 •

edited

Loading

the unit test is :example/ck_tile/50_sparse_attn/ ,

the bin is：./bin/tile_example_vsa_sparse_attn, ./bin/tile_example_jenga_sparse_attn

run this cmd to verify the correctness


          add sparse attention VSA

0607d31

jiangyon-amd requested review from ThomasNing, afagaj, andriy-ca, aosewski, asleepzzz, bartekxk, carlushuang, cgmillette, coderfeli, geyyer, illsilin, poyenc, qianfengz, shumway, tenpercent and vidyasagar-amd as code owners

December 3, 2025 02:48


          Merge branch 'develop' into sparse_attention_VSA

4fc61d9

Collaborator

ltqin commented Dec 3, 2025

Add kernel and pipeline directories to the spare_attn directory, and place the kernel file and pipeline file separately

jiangyon-amd added 11 commits

December 4, 2025 08:51


          Merge branch 'develop' into sparse_attention_VSA

1dc15fb


          fix the pre-commit

8ff98b8


          Add jenga test and pre-commit

3b00e40


          add bf16 for vsa

997ec8f


          add jenga support bf16

29d96a9


          remove lse arg

5e8a010


          split kernel code to block & kernel

d2278ab


          fix the pre-commit

faff9ab


          fix the pre-commit

55d9a8e


          fix the copyrights

a86fc80


          fix the copyright

12420cd

poyenc reviewed

View reviewed changes

example/ck_tile/50_sparse_attn/codegen/ops/fmha_fwd_vsa.py Outdated Show resolved Hide resolved

jiangyon-amd added 2 commits

January 28, 2026 23:32


          remove the useless code & resolved the comments

83bed30


          remove useless code

3eef3ad

poyenc reviewed

View reviewed changes

example/ck_tile/50_sparse_attn/fmha_fwd_trek.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_jenga.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_jenga.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_jenga.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/pipeline/block_fmha_pipeline_qr_ks_vs_async_jenga.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

poyenc reviewed

View reviewed changes

include/ck_tile/ops/sparse_attn/kernel/fmha_fwd_vsa_kernel.hpp Outdated Show resolved Hide resolved

jiangyon-amd and others added 3 commits

January 29, 2026 15:48


          remove useless code

a9456ba


          Merge branch 'develop' into sparse_attention_VSA

56a3ab3


          Clean up code

4693a5c

jiangyon-amd marked this pull request as draft

January 29, 2026 09:41

poyenc added 3 commits

January 29, 2026 03:43


          Remove more unused code

685842c


          Re-format .hpp

368d050


          Refactor codegen scripts

565ab74

jiangyon-amd marked this pull request as ready for review

January 29, 2026 10:21

poyenc and others added 3 commits

January 30, 2026 01:22


          Merge branch 'develop' into sparse_attention_VSA

b375f94


          Merge branch 'develop' into sparse_attention_VSA

d885111


          Merge branch 'develop' into sparse_attention_VSA

6dd398d

poyenc approved these changes

View reviewed changes

Contributor

poyenc left a comment •

edited

Loading

We should add more unit tests for the new kernel/pipelines.

asleepzzz approved these changes

View reviewed changes

poyenc merged commit 4d2f8c1 into develop

19 of 20 checks passed

poyenc deleted the sparse_attention_VSA branch

January 30, 2026 16:59

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

ltqin ltqin left review comments

Copilot code review Copilot Copilot left review comments

poyenc poyenc approved these changes

asleepzzz asleepzzz approved these changes

illsilin Awaiting requested review from illsilin illsilin is a code owner

carlushuang Awaiting requested review from carlushuang carlushuang is a code owner

qianfengz Awaiting requested review from qianfengz qianfengz is a code owner

aosewski Awaiting requested review from aosewski aosewski is a code owner

geyyer Awaiting requested review from geyyer geyyer is a code owner

bartekxk Awaiting requested review from bartekxk bartekxk is a code owner

andriy-ca Awaiting requested review from andriy-ca andriy-ca is a code owner

afagaj Awaiting requested review from afagaj afagaj is a code owner

tenpercent Awaiting requested review from tenpercent tenpercent is a code owner

ThomasNing Awaiting requested review from ThomasNing ThomasNing is a code owner

coderfeli Awaiting requested review from coderfeli coderfeli is a code owner

cgmillette Awaiting requested review from cgmillette cgmillette is a code owner

shumway Awaiting requested review from shumway shumway is a code owner

vidyasagar-amd Awaiting requested review from vidyasagar-amd vidyasagar-amd is a code owner

vpietila-amd Awaiting requested review from vpietila-amd vpietila-amd is a code owner

Snektron Awaiting requested review from Snektron Snektron is a code owner

Labels

None yet