Ali Nouri 91317bdfe9 Add atomic-free MOE GEMM implementation ...

- Add FusedMoeGemmTilePartitioner_NoAtomic: Forces single workgroup per expert
- Add FusedMoeGemmPipelineFlatmmPolicy_NoAtomic: Fixes alignment consistency
- Update API to use no-atomic approach when intermediate_size <= Block_N0

Eliminates atomic operations by ensuring each workgroup handles complete
expert computation without K-dimension splitting.

2025-09-26 22:50:37 +00:00

..

add_rmsnorm2d_rdquant

Support Wave32 in CK_TILE - Part 1 (#2594)

2025-08-18 10:08:31 -07:00

batched_transpose

refactor: use snake_case naming in ck_tile/core components (#2766)

2025-09-03 09:34:11 +02:00

common

[CK_TILE] Refine Generic2dBlockShape to fix ck_tile example 2,10,11,14 on rdna3 and 4 (#2795)

2025-09-10 08:29:20 +08:00

elementwise

[CK_TILE] Fix example batched_gemm, grouped_gemm, gemm_multi_d, convolution on gfx11 & gfx12 (#2808)

2025-09-11 07:27:33 -07:00

epilogue

[CK_TILE] Add permuteN optimization to remove lds operation in c_shuffle (#2764)

2025-09-08 22:02:48 -07:00

flatmm

[CK_TILE] Fix flatmm on gfx11 and gfx12 (#2790)

2025-09-10 08:28:00 +08:00

fmha

[CK_TILE] Fix kname & typo in FMHA BWD (#2809)

2025-09-09 15:08:00 -07:00

fused_moe

Add atomic-free MOE GEMM implementation

2025-09-26 22:50:37 +00:00

gemm

[CK_TILE] Fix flatmm on gfx11 and gfx12 (#2790)

2025-09-10 08:28:00 +08:00

gemm_group_quant

[CK TILE GEMM] Fixed the regression issue with transpose C in Quant Gemm (#2819)

2025-09-11 11:38:16 -07:00

grouped_convolution

[CK_TILE] Fix example batched_gemm, grouped_gemm, gemm_multi_d, convolution on gfx11 & gfx12 (#2808)

2025-09-11 07:27:33 -07:00

image_to_column

Support Wave32 in CK_TILE - Part 1 (#2594)

2025-08-18 10:08:31 -07:00

layernorm2d

[CK-TILE] Default2DEpilogue, example and adding nullptr_t type for D (#2752)

2025-08-28 12:45:50 -07:00

norm_reduce

[CK_TILE] Fix compilation errors introduced in #2320, #2219 and #2214 (#2388)

2025-06-23 12:29:15 +08:00

permute

Support Wave32 in CK_TILE - Part 1 (#2594)

2025-08-18 10:08:31 -07:00

reduce

[CK_TILE] fix example reduces, permute and elementwise on gfx11 & gfx12 (#2810)

2025-09-11 12:41:20 +08:00

rmsnorm2d

[CK-TILE] Default2DEpilogue, example and adding nullptr_t type for D (#2752)

2025-08-28 12:45:50 -07:00

smoothquant

Support Wave32 in CK_TILE - Part 1 (#2594)

2025-08-18 10:08:31 -07: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

Support Wave32 in CK_TILE - Part 1 (#2594)

2025-08-18 10:08:31 -07: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] Merge transpose examples (#2450)

2025-07-26 21:51:54 -07:00

common.hpp

[CK TILE] Stream-K tile partitioner (#2708)

2025-09-03 13:38:17 -07:00

elementwise.hpp

[CK_TILE] Support for elementwise kernel (#2246)

2025-07-24 11:21:45 +02: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 FAv3 fwd pipeline (#2731)

2025-09-01 09:16:45 +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_group_quant.hpp

[CK_TILE] Implement Row/Col quant grouped gemm (#2786)

2025-09-08 10:25:57 -07:00

gemm.hpp

[Regression] Fix CK_TILE build error in grouped_convolution, copy_basic and fused_moegemm_kernel (#2728)

2025-08-28 20:30:30 +08:00

grouped_convolution.hpp

[CK Tile] Grouped convolution backward data (#2652)

2025-08-20 05:29:57 -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

GEMM Multi D for CK Tile Engine (#2660)

2025-08-12 16:05:05 -07:00

rmsnorm2d.hpp

[CK_TILE] Enhance RMSNorm Accuracy: New Pipeline Pass for Selectable Implementation (#2409)

2025-07-16 14:05:26 +08: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