mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 18:17:44 +00:00

Files

Po Yen Chen c3f74a83b3 [CK_TILE] Add logits soft-capping & customization support to the FMHA forward kernel/pipelines (#2163 )

* hack for cap logits

* fix bug

* Re-format files

* Allow specifying logits_soft_cap through APIs

* Support turn on/off logits_soft_cap in async pipeline

* Do not generate non-verified kernels

* Align receipt used in Aiter

* Sync logits soft-capping across pipelines

* Re-enable some hdim pipelines

* fix perf

* Add attention variant for logits_soft_cap

* Add newline at end-of-file

* Fix performance

* Add comment to explain logits_soft_cap pre-processing

* Unify code

* Unify floating-point literal style

* Use class data member to slience the compilation error

* [CK_TILE] Update attention customizaton interface: add LogitsMask() (#2133)

* Send 'mask' along with variant params to the LogitsMask()

* Send block indices to the variant

* Add indices parameters in variant interface

* Fix fmha bwd codegen error

* Allow switch logits_soft_cap impl

* Eliminate register spills

* Fix compilation errors

* Fix wrong LSE

* Fix LSE for splitkv kernel

* Sync splitkv pipeline changes

* Add batch_prefill kernel/pipeline

* Fix codegen error

* Undo changes in CMakeLists.txt

* Merge pipeline filtering check

* Use different code path if kHasLogitsSoftCap=false

* Remove [[maybe_unused]] attribute

* Use pre-existing compile-time flag to instantiate templates

* Sync pipeline changes

* Update CHANGELOG.md

---------

Co-authored-by: Bernard <bernaliu@amd.com>
Co-authored-by: coderfeli <coderfeli@163.com>

[ROCm/composable_kernel commit: 2920604786]

2025-05-13 12:19:25 +08:00

algorithm

Solve the Static Encoding Pattern compile error when the tile size is too small (#2079 )

2025-04-13 20:09:30 -07:00

arch

[CK_TILE] optimize moe sorting kernel, boost large context case up to 20x (#2153 )

2025-05-06 17:32:07 +08:00

container

Add documentation for ck_tile::array<T,N> (#2078 )

2025-04-30 16:43:36 -07:00

numeric

[CK_TILE] Add logits soft-capping & customization support to the FMHA forward kernel/pipelines (#2163 )

2025-05-13 12:19:25 +08:00

tensor

[CK_TILE] Add logits soft-capping & customization support to the FMHA forward kernel/pipelines (#2163 )

2025-05-13 12:19:25 +08:00

utility

Split env.hpp header from the ck.hpp header. (#2049 )

2025-04-03 15:30:21 -07:00

config.hpp

[CK_TILE] optimize moe sorting kernel, boost large context case up to 20x (#2153 )

2025-05-06 17:32:07 +08:00

README.md

introducing ck_tile! (#1216 )

2024-04-15 19:27:12 -05:00

README.md

ck_tile/core

ck_tile/core contains every basic functions and structures to create a GPU kernel using ck_tile. User should only include ck_tile/core.hpp this single header to use all the functionality. Everything is under ck_tile namespace. The coding style under this folder should be similar to std (snake_case for structure/function, Camel for template types...)

algorithm/
    coordinate transform and some other reusable algorithm
arch/
    contains some basic device building block like mma, buffer addressing, etc...
container/
    contains basic container data structure, array/sequence/tuple/...
numeric/
    data type, and data type related math
tensor/
    tensor descriptors and tile level API
utility/
    other utility function for both host/device