Po Yen Chen 3ca1b9df66 [CK_TILE] fMHA batch_prefill block index & logits soft-capping optimizations (#2198) ...

* Write soft-sign in inline asm

* Change tile idx computation

* Add macro to turn off soft-sign asm opt

* Use simple for loop to avoid register spill

* Only do block id transform for masking cases

[ROCm/composable_kernel commit: 791802b381]

2025-05-16 15:14:46 +08:00

..

block_fmha_batch_prefill_pipeline_qr_ks_vs_async_default_policy.hpp

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

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

block_fmha_batch_prefill_pipeline_qr_ks_vs_async.hpp

[CK_TILE] fMHA batch_prefill block index & logits soft-capping optimizations (#2198)

2025-05-16 15:14:46 +08:00

block_fmha_bwd_convert_dq.hpp

[CK_TILE] FA bwd kernels optimization (#1397)

2024-08-16 13:40:10 -07:00

block_fmha_bwd_dot_do_o.hpp

[CK_TILE] FA bwd kernels optimization (#1397)

2024-08-16 13:40:10 -07:00

block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr_iglp.hpp

[CK_TILE] Improve headdim96 performance for fmha-bwd (#1573)

2024-10-16 18:14:32 +08:00

block_fmha_bwd_dq_dk_dv_pipeline_kr_ktr_vr.hpp

[CK_TILE] Improve headdim96 performance for fmha-bwd (#1573)

2024-10-16 18:14:32 +08:00

block_fmha_bwd_pipeline_default_policy.hpp

[CK_TILE] Improve headdim96 performance for fmha-bwd (#1573)

2024-10-16 18:14:32 +08:00

block_fmha_bwd_pipeline_enum.hpp

[CK_TILE] FA bwd kernels optimization (#1397)

2024-08-16 13:40:10 -07:00

block_fmha_bwd_pipeline_problem.hpp

[CK_TILE] FA bwd kernels optimization (#1397)

2024-08-16 13:40:10 -07:00

block_fmha_fwd_appendkv_pipeline_default_policy.hpp

[CK_TILE] Add PagedAttention kernels (#1387)

2024-08-28 20:50:43 +08:00

block_fmha_fwd_appendkv_pipeline.hpp

[CK_TILE] Add PagedAttention kernels (#1387)

2024-08-28 20:50:43 +08:00

block_fmha_fwd_splitkv_combine_pipeline_default_policy.hpp

[CK_TILE] Add fmha fwd N-Warp S-Shuffle pipeline (fmha fwd splitkv pipeline variant) (#1705)

2024-12-20 14:41:01 +08:00

block_fmha_fwd_splitkv_combine_pipeline.hpp

roll back splitkv

2025-01-08 10:09:54 -08:00

block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_fwd_splitkv_pipeline_nwarp_sshuffle_qr_ks_vs.hpp

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

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

block_fmha_fwd_splitkv_pipeline_qr_ks_vs_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_fwd_splitkv_pipeline_qr_ks_vs.hpp

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

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

block_fmha_pipeline_enum.hpp

[CK_TILE] support alibi (#1269)

2024-05-07 22:32:54 +08:00

block_fmha_pipeline_problem.hpp

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

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

block_fmha_pipeline_qr_ks_vs_async_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_pipeline_qr_ks_vs_async.hpp

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

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

block_fmha_pipeline_qr_ks_vs_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_pipeline_qr_ks_vs_fp8.hpp

[CK_TILE] Add fmha fwd headdim96 support (#1608)

2024-10-30 14:03:16 +08:00

block_fmha_pipeline_qr_ks_vs_whole_k_prefetch_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_pipeline_qr_ks_vs.hpp

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

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

block_fmha_pipeline_qs_ks_vs_default_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

block_fmha_pipeline_qs_ks_vs.hpp

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

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

block_fmha_pipeline_qx_ks_vs_custom_policy.hpp

Ck tile/complete k prefetch (#1941)

2025-03-07 14:19:51 +08:00

tile_fmha_shape.hpp

Reapply "[CK_TILE] support hdim=192/128 pair for deepseekv3 (#1961)" … (#1971)

2025-03-13 11:41:39 +08:00

tile_fmha_traits.hpp

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

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