Po Yen Chen 791802b381 [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

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

..

add_rmsnorm2d_rdquant

[A8W8 GEMM] Optimized weight-preshuffled implementation & add quantization datatype for CK TILE rms_norm (#1862)

2025-02-20 14:00:27 -08:00

batched_transpose

Vectorized Transpose for Batched Transpose CK Tile Operator (#2131)

2025-05-12 00:41:45 -07:00

common

fix ck_tile/basic_gemm build error (#1988)

2025-03-20 22:01:14 -07:00

elementwise

[CK TILE] GEMM with packed i4 (#1885)

2025-02-20 09:59:49 +01:00

epilogue

fix for default epilogue (#2167)

2025-05-07 10:46:53 -07:00

flatmm

fix moe sorting build fail (#2190)

2025-05-14 09:31:26 +08:00

fmha

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

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

fused_moe

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

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

gemm

Adding validation for tile sizes in Tile Engine (#2189)

2025-05-15 10:28:31 -07:00

image_to_column

[CK_TILE] Image to Column kernel (#1532)

2024-09-27 22:57:38 +02:00

layernorm2d

[CK_TILE] Improve RMS/Layer Normalization 2 Pass Pipeline Performance (#1861)

2025-03-25 20:09:45 +08:00

norm_reduce

Ck tile/layernorm: implement naive reduce, opt performance (#1784)

2025-01-03 14:28:59 +08:00

permute

Remove virtual destructors from unary ops (#1610)

2024-10-30 17:42:50 +01:00

reduce/block

[Ck_tile] smoothquant (#1617)

2024-11-01 13:51:56 +08:00

rmsnorm2d

[CK_TILE] Improve RMS/Layer Normalization 2 Pass Pipeline Performance (#1861)

2025-03-25 20:09:45 +08:00

smoothquant

add fp8 as dst (#1830)

2025-01-22 17:34:27 +08:00

softmax/block

topk_softmax (#1592)

2024-10-26 23:52:49 +08:00

topk/block

topk_softmax (#1592)

2024-10-26 23:52:49 +08:00

topk_softmax

topk_softmax (#1592)

2024-10-26 23:52:49 +08:00

add_rmsnorm2d_rdquant.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

batched_transpose.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

common.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

elementwise.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

epilogue.hpp

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

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

flatmm.hpp

[flatmm] implement basic fp16 flatmm (#2089)

2025-04-16 16:51:17 +08:00

fmha.hpp

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

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

fused_moe.hpp

[CK_TILE] moe sorting ex kernel to support expert > 128 (#1840)

2025-02-11 17:49:17 +08:00

gemm.hpp

[CK_TILE] support gfx950 matrix core in 01_fmha fwd (#2110)

2025-04-23 12:40:18 -07:00

image_to_column.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

layernorm2d.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

norm_reduce.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

permute.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

reduce.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

rmsnorm2d.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

smoothquant.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

softmax.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

topk_softmax.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00

topk.hpp

[CK_TILE] Add GetName for GEMM kernels (#1791)

2025-02-11 09:49:48 +01:00