ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-13 10:37:42 +00:00
Code Issues Packages Projects Releases Wiki Activity
2,307 Commits 778 Branches 40 Tags
vmcnt0issue
Commit Graph

135 Commits

Author SHA1 Message Date
aska-0096
7bdf6a7eef merge develop and solve conflicts 2025-08-22 03:15:51 +00:00
aska-0096
f21e916a8c Merge branch 'develop' of https://github.com/ROCm/composable_kernel into vmcnt0issue 2025-08-22 03:15:21 +00:00
Po Yen Chen
4a7ecce096 [CK_TILE][FMHA] Enable dwordx4 loading in async_load_tile_raw() (#2549) 
* Support async load dwordx4

* Enlarge load size on gfx950
 2025-08-22 10:13:47 +08:00
Yi DING
4cfa2c7158 [CK_TILE] FMHA BWD Fix Compilation with Bias (#2682) 
* [CK_TILE] FMHA BWD Fix Compilation with Bias

* Fix appendkv kApplyRoPE
 2025-08-22 10:01:10 +08:00
Ding, Yi
a7dd2a7d13 update bwd 2025-08-19 02:17:43 +00:00
linqunAMD
9fcc1ee9fd Support Wave32 in CK_TILE - Part 1 (#2594) 
* Support wave32/wave64 in CK_TILE - Part 1

* remove blocksize in kernel launch

* fix build error

* fix clang format

* fix clang format 2

* fix clang format 3

* fix fmha build error

* fix fmha build 2

* fix fmha build 3

* fix build error 4

* address review comment

* update change log

* replace KernelBlockSize with kBlockSize

* fix CI fail

* fix clang format

* address review comment and rebase code.

* fix universal test fail

---------

Co-authored-by: Lin, Qun <Quentin.Lin+amdeng@amd.com>
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
 2025-08-18 10:08:31 -07:00
Kevin Choi
75c4b9372f remove noinline attr as it causes a lot more s_waitcnt's 2025-08-16 08:15:23 +00:00
Tianyuan Wu
68134b60e4 [CK_TILE] CK_TILE GEMM WMMA Support for GFX11/GFX12 (#2466) 
* WMMA GEMM F16 Implementation

Signed-off-by: root <tianyuwu@amd.com>

* Self-review

Signed-off-by: root <tianyuwu@amd.com>

* ASIC check minor tweak

Signed-off-by: root <tianyuwu@amd.com>

* add missing include file

* Set GPU_TARGETS to gfx11/12 generic

Signed-off-by: root <tianyuwu@amd.com>

* INT8 GFX12

Signed-off-by: root <tianyuwu@amd.com>

* add int8x16 branch

* Fix CI script

Signed-off-by: root <tianyuwu@amd.com>

* Fix typo

Signed-off-by: root <tianyuwu@amd.com>

* Add CK_Tile WMMA example

Signed-off-by: Tianyuan Wu <tianyuwu@amd.com>

* Fix CI

Signed-off-by: Tianyuan Wu <tianyuwu@amd.com>

* fix clang format

* Set M/N_Warp Back to Constant

Signed-off-by: Tianyuan Wu <tianyuwu@amd.com>

* Use GemmConfigComputeV3 by default

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Enable CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT for gfx12

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Remove CK_Tile wmma gemm examples from the CI list

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Add atomic add fallback method for gfx11

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Fix typo

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Omit copyright year

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Support non-square cases

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Fix CI

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Add get_device_ip()

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Revert "Add atomic add fallback method for gfx11"

This reverts commit 07a79e797d.

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Revert "Enable CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT for gfx12"

This reverts commit ceee918007.

* Revise method name and typos

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* clang-format

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Try fix CI

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Revert "Try fix CI"

This reverts commit 7a7241085e.

* clang-format

Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>

* Fix typo caused by merge

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Fix typo caused by merging

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

---------

Signed-off-by: root <tianyuwu@amd.com>
Signed-off-by: Tianyuan Wu <tianyuwu@amd.com>
Signed-off-by: TianyuanWu <Tianyuan.Wu@amd.com>
Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>
Co-authored-by: joye <joye@amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>
 2025-08-15 16:22:27 -07:00
Kevin Choi
598e3fec41 remove innerloop, move restrict parameters to mainloop and add noinline attribute. 2025-08-14 12:11:17 +00:00
Kevin Choi
3340408537 Create inner lambda with restrict parameters, add restrict to some parameters 2025-08-14 07:06:51 +00:00
aska-0096
3bc45ecbc7 save for debug 2025-08-14 03:43:54 +00:00
SamiAario-AMD
28a97865f5 Cleanups (#2631) 
* Remove some duplicate code in fmha_fwd_appendkv_kernel.hpp

* Simplify two templated operator calls by having the templated types deduced automatically

* Simplify two GemmPipeline calls

* Fix GemmPipelineAgBgCrCompV4::GetName

* Refactor use of ArgParser in CK tile GEMM examples

* Update args in README.md to match the implementation in create_args

* Remove some unnecessary include statements

* Rename two variables

* Factor out common code

* Factor out do_verify

* Add and use type aliases for memory operation integral constants

* In gemm_basic.cpp, use kPadM, kPadN, kPadK, and kBlockPerCu from GemmConfig

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
 2025-08-13 10:12:08 +02:00
Haocong WANG
05a6e92705 Re-enable optimization for gfx950 fmha fwd (#2671) 
* Fix for fwd/bwd kernel build filter

* fix bwd code

* save an example for __bf16 type

* temp save, waiting for debug

* tempsave, fmha_decode

* temp save, change all instance to 1wave

* fix async copytest bug

* Add block_sync_lds_direct_load utility

* fix the s_waitcnt_imm calculation

* Improve s_waitcnt_imm calculation

* fix vmcnt shift

* add input validation and bug fix

* remove unnecessary output

* move test_copy into test

* temp save

* tempsave

* compile pass

* tempsave, trload+asyncload done

* tempsave. asynccopy+trload sanity checked

* remove unnecessary features

* fix the lds alignment caused performance regression

* enable prefill overload operator().

* remove all lds bankconflict with xor layouts

* enable larger tile size; upgrade xor pattern

* upgrade prefill pipeline; simple iglp; consistent data produce and consume order

* small refactor

* Load Q through lds, implement xor;

* add vmcnt guard before load ktile

* Add v_permlaneb32 for block_reduce. Disable it as it will cause un-coexecutable packed math in FA

* Add XOR fold strategy for hdim<128, but perf dropped; disable it by default; wait further perf debug

* add __restrict__ to tr load

* merge fa_decode pipeline into fmha_fwd api

* remove unnecessary files; rename some files

* Remove unnecessary changes

* bug fix, clang format;

* remove non-necessary change

* fix clangformat with 18.1.3

* fix bugs

* fix bug

* fix bug on non-gfx950

* fix bugs in gemm

* fix bug in pki4

* tempsave, update the blocksync functions

* change the warp setting for hdim32 fmha fwd

* clang format

* fix conflict. disable all v-col instance for fmha fwd

* Fix the bug

* clang format

* refactor blockgemm change, isolate to v2;

---------

Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
Co-authored-by: asleepzzz <hanwen.chang@amd.com>
 2025-08-13 14:57:43 +08:00
aska-0096
0810799e25 refactor blockgemm change, isolate to v2; 2025-08-12 14:25:50 +00:00
asleepzzz
5b39de4bb6 Revert "Optimize fmha fwd decode & prefill for gfx950 (#2641)" (#2670) 
This reverts commit b7322a521a.
 2025-08-12 20:27:10 +08:00
Haocong WANG
b7322a521a Optimize fmha fwd decode & prefill for gfx950 (#2641) 
* Fix for fwd/bwd kernel build filter

* fix bwd code

* save an example for __bf16 type

* temp save, waiting for debug

* tempsave, fmha_decode

* temp save, change all instance to 1wave

* fix async copytest bug

* Add block_sync_lds_direct_load utility

* fix the s_waitcnt_imm calculation

* Improve s_waitcnt_imm calculation

* fix vmcnt shift

* add input validation and bug fix

* remove unnecessary output

* move test_copy into test

* temp save

* tempsave

* compile pass

* tempsave, trload+asyncload done

* tempsave. asynccopy+trload sanity checked

* remove unnecessary features

* fix the lds alignment caused performance regression

* enable prefill overload operator().

* remove all lds bankconflict with xor layouts

* enable larger tile size; upgrade xor pattern

* upgrade prefill pipeline; simple iglp; consistent data produce and consume order

* small refactor

* Load Q through lds, implement xor;

* add vmcnt guard before load ktile

* Add v_permlaneb32 for block_reduce. Disable it as it will cause un-coexecutable packed math in FA

* Add XOR fold strategy for hdim<128, but perf dropped; disable it by default; wait further perf debug

* add __restrict__ to tr load

* merge fa_decode pipeline into fmha_fwd api

* remove unnecessary files; rename some files

* Remove unnecessary changes

* bug fix, clang format;

* remove non-necessary change

* fix clangformat with 18.1.3

* fix bugs

* fix bug

* fix bug on non-gfx950

* fix bugs in gemm

* fix bug in pki4

* tempsave, update the blocksync functions

* change the warp setting for hdim32 fmha fwd

* clang format

* fix conflict. disable all v-col instance for fmha fwd

* Fix the bug

* clang format

---------

Co-authored-by: Max Podkorytov <4273004+tenpercent@users.noreply.github.com>
 2025-08-12 19:43:14 +08:00
aska-0096
75f6f6bac4 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into wip-async-tr-fa 2025-08-12 09:04:41 +00:00
Yi DING
8e1eb0c1ee [CK_TILE] FMHA BWD Decode Pipeline (#2643) 
* Fix distr

* Duplicate block_fmha_bwd_dq_dk_dv_pipeline_trload_kr_ktr_vr

* decode 16x16 o2
 2025-08-12 17:02:52 +08:00
aska-0096
96d24497f5 fix conflict. disable all v-col instance for fmha fwd 2025-08-12 04:02:41 +00:00
aska-0096
1716171be4 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into wip-async-tr-fa 2025-08-12 03:52:34 +00:00
Yi DING
4fde1646e5 [CK_TILE] FMHA BWD Optimization For GFX950 (#2628) 
* simplify fmha_bwd_kernel MakeKargs & dq_dram_window

* simply duplicate

* trload pipeline

* Try two-stage

* add prefetch

* optimize & iglp
 2025-08-12 11:11:55 +08:00
aska-0096
efb8549279 fix bug 2025-08-08 17:53:19 +00:00
aska-0096
729e8785fb fix bugs 2025-08-08 15:42:15 +00:00
aska-0096
78edd7303b bug fix, clang format; 2025-08-08 09:04:02 +00:00
aska-0096
3b9fb6af38 Remove unnecessary changes 2025-08-08 08:08:03 +00:00
aska-0096
6bb57c2c57 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into wip-async-tr-fa 2025-08-08 07:50:12 +00:00
aska-0096
1ecee378d5 remove unnecessary files; rename some files 2025-08-08 06:19:31 +00:00
aska-0096
b4640a9de6 merge fa_decode pipeline into fmha_fwd api 2025-08-08 05:46:18 +00:00
Yi DING
b0a97498b0 [CK_TILE] FMHA BWD Remove Unnecessary Padding (#2550) 
* Remove unnecessary pssk

* Add BlockFmhaBwdDQDKDVPipeline wrapper

* Resolve copilot comments & Remove kpad & fix

* Remove spad
 2025-08-07 21:24:43 +08:00
aska-0096
414cad667b Add XOR fold strategy for hdim<128, but perf dropped; disable it by default; wait further perf debug 2025-08-05 07:23:51 +00:00
aska-0096
0d12fc944f Add v_permlaneb32 for block_reduce. Disable it as it will cause un-coexecutable packed math in FA 2025-08-04 10:27:42 +00:00
aska-0096
4f31847de1 add vmcnt guard before load ktile 2025-08-04 10:02:17 +00:00
aska-0096
746f4ccb99 Load Q through lds, implement xor; 2025-08-04 06:49:01 +00:00
aska-0096
2d4e73d2b4 small refactor 2025-08-01 10:44:54 +00:00
aska-0096
a28b6e67fe upgrade prefill pipeline; simple iglp; consistent data produce and consume order 2025-07-31 10:25:37 +00:00
aska-0096
75cba48682 enable larger tile size; upgrade xor pattern 2025-07-31 05:13:27 +00:00
aska-0096
69890afc98 remove all lds bankconflict with xor layouts 2025-07-30 12:25:33 +00:00
aska-0096
8dacc35c4c enable prefill overload operator(). 2025-07-30 03:51:06 +00:00
Illia Silin
49723e94bb fix the clang-format (#2578) 2025-07-28 20:49:55 -07:00
Yi DING
1926cd0cb8 [CK_TILE] FMHA bwd Support hdim as a Multiple of 32 (#2130) 
* Fix shuffle_tile

* Add fmha bwd d160

* CHANGELOG

* Use static_cast

* Update

---------

Co-authored-by: asleepzzz <hanwen.chang@amd.com>
 2025-07-29 09:31:14 +08:00
Illia Silin
504b101da3 upgrade from clang-format-12 to clang-format-18 (#2568) 
* upgrade to clang-format-18

* update to clang-format-18 in pre-commit-config
 2025-07-28 11:34:07 -07:00
shay-li77
8ae528a1b4 fix mha bwd dbias random mismatch (#2570) 
* fix mha bwd dbias random mismatch

* formatting code
 2025-07-28 14:39:31 +08:00
liang
d2459878cf reorder grid dim schedule (#2533) 
Co-authored-by: smallmou <liangshenghao.lsh@alibaba-inc.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
 2025-07-26 02:46:55 +08:00
aska-0096
13bcc913de fix the lds alignment caused performance regression 2025-07-25 07:10:01 +00:00
aska-0096
af28123cec remove unnecessary features 2025-07-23 09:05:57 +00:00
aska-0096
14e0ab70c6 tempsave. asynccopy+trload sanity checked 2025-07-22 08:04:05 +00:00
aska-0096
1b468bac0b tempsave, trload+asyncload done 2025-07-21 05:55:55 +00:00
aska-0096
afd96d8180 compile pass 2025-07-18 10:04:34 +00:00
aska-0096
5616551115 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into wip-async-tr-fa 2025-07-18 05:17:27 +00:00
aska-0096
ae39c84f55 tempsave 2025-07-18 05:16:39 +00:00