ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 19:28:33 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
db98d0e5aa339274fa39e613f7b33f3dd97da1c3
composable_kernel/example
History
Amir Ghamarian d93efe1b61 Add fused topk_softmax_decode kernel for M=1 MoE decode 
New CK tile kernel variant that fuses topk_softmax and moe_sorting into
a single kernel launch for the decode case (M=1, single token). The
pipeline inlines the topk loop with results in shared memory (no global
scratch), then thread 0 emits moe_sorting-compatible packed output.

Includes CMake target tile_example_topk_softmax_decode with built-in
comparison benchmark against the separate topk+sorting baseline.

Validated on gfx950, E=8..1024, k=1..8, bf16/fp16.

Made-with: Cursor
2026-03-29 18:06:03 +00:00
..
01_gemm
Implement device_gemm_universal_preshuffle_instance for RDNA4 (#3429)
2026-01-15 07:19:31 -08:00
02_gemm_bilinear
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
03_gemm_bias_relu
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
04_gemm_add_add_fastgelu
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
09_convnd_fwd
[CK] Integrate GPU reference into ckProfiler for convolutions (#3379)
2025-12-18 07:59:45 +01:00
10_convnd_fwd_multiple_d_multiple_reduce
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
11_convnd_fwd_bias
[DOCS] Documentation Addition (Readme updates) (#2495)
2025-10-16 03:10:57 -07:00
12_reduce
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
13_pool2d_fwd
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
14_gemm_quantization
[CK][Examples] Fixing stride issues in ck examples 14/65/68/69 by workaround - Bypassing hostTensor validation
2026-01-15 16:43:02 +01:00
15_grouped_gemm
Implement grouped gemm tile loop for RDNA4 (#3304)
2026-01-13 07:14:23 +01:00
16_gemm_multi_d_multi_reduces
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
17_convnd_bwd_data
[CK] Integrate GPU reference into ckProfiler for convolutions (#3379)
2025-12-18 07:59:45 +01:00
18_batched_gemm_reduce
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
19_binary_elementwise
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
20_grouped_conv_bwd_weight
[CK] Integrate GPU reference into ckProfiler for convolutions (#3379)
2025-12-18 07:59:45 +01:00
21_gemm_layernorm
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
22_cgemm
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
23_softmax
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
24_batched_gemm
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
25_gemm_bias_e_permute
Implement batched gemm bias permute for RDNA4 (#3534)
2026-01-17 08:30:27 +01:00
26_contraction
[ck][gfx12] support contraction on gfx12 (#3421)
2025-12-15 07:16:01 -08:00
27_layernorm2d_fwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
28_grouped_gemm_bias_e_permute
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
29_batched_gemm_bias_e_permute
Implement batched gemm bias permute for RDNA4 (#3534)
2026-01-17 08:30:27 +01:00
30_grouped_conv_fwd_multiple_d
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
31_batched_gemm_gemm
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
32_batched_gemm_scale_softmax_gemm
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
33_multiple_reduce
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
34_batchnorm
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
35_splitK_gemm
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
36_sparse_embedding
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
37_batched_gemm_add_add_relu_gemm_add
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
38_grouped_conv_bwd_data_multiple_d
Grouped convolution backward data WMMA v3 implementation (#3460)
2025-12-30 16:25:08 +01:00
39_permute
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
40_conv2d_fwd_quantization
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
41_grouped_conv_conv_fwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
42_groupnorm_fwd
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
43_splitk_gemm_bias_e_permute
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
44_elementwise_permute
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
45_elementwise_normalization
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
46_gemm_add_multiply
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
47_gemm_bias_softmax_gemm_permute
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
48_pool3d_fwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
49_maxpool2d_bwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
50_put_element
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
51_avgpool3d_bwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
52_im2col_col2im
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
53_layernorm2d_bwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
54_groupnorm_bwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
59_grouped_gemm_multi_ABD
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
60_gemm_multi_ABD
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
61_contraction_multi_ABD
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
62_convnd_activ
Grouped convolution forward device implementation and base flavors for RDNA3/4 (#2964)
2025-12-18 13:12:15 -07:00
63_layernorm4d_fwd
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
64_fpAintB_gemm
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
65_gemm_multiply_multiply
[CK][Examples] Fixing stride issues in ck examples 14/65/68/69 by workaround - Bypassing hostTensor validation
2026-01-15 16:43:02 +01:00
66_complex_contraction_bilinear
chore(copyright) update library wide CMakeLists.txt copyright header template (#3313)
2025-11-28 13:49:54 -08:00
67_gemm_microscaling
[CI, CK examples] Disable time_kernel for CI tests and examples (#3464)
2026-01-07 16:30:57 +01:00
68_gemm_add
[CK][Examples] Fixing stride issues in ck examples 14/65/68/69 by workaround - Bypassing hostTensor validation
2026-01-15 16:43:02 +01:00
69_gemm_add_relu
[CK][Examples] Fixing stride issues in ck examples 14/65/68/69 by workaround - Bypassing hostTensor validation
2026-01-15 16:43:02 +01:00
ck_tile
Add fused topk_softmax_decode kernel for M=1 MoE decode
2026-03-29 18:06:03 +00:00
CMakeLists.txt
Build CK on Windows (#3458)
2026-01-14 07:31:45 -08:00
README.md
Add basic documentation structure (#1715)
2024-12-04 00:46:47 +01:00

README.md

Back to the main page

Composable Kernel examples

Reference in New Issue View Git Blame Copy Permalink