ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-13 01:36:06 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
d814fefe1898971b2c3eb97b986bdafc450f18b5
composable_kernel/example
History
Po Yen Chen ad9863fe05 [CK_TILE] Low CU utilization optimization for fMHA fwd kernels (#2402) 
* Wrap tile size mapping as class method

* Warp pipeline generating as class method

* Add constraint as kernel dispatching criteria

* Support mutltiple tile size for a (hdim, hdim_v) combination

* Use smaller tile size if CU utilization is low

* Use integar as the key of the tile size map

* Fix type error

* Simply override parent class method return value

* Add attribute to eliminate warnging

* Allow using environment variables to turn on/off custom factory

* Unify param naming style

* Add missing HIP runtime include directive

* Fix os.environ.get() usage
2025-07-09 22:01:33 +08:00
..
01_gemm
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +08:00
02_gemm_bilinear
Fix issue with multiple targets and remove smfmac tests from unsupported test targets (#1372)
2024-07-03 23:34:38 -07:00
03_gemm_bias_relu
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
04_gemm_add_add_fastgelu
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
09_convnd_fwd
Add a daily CI build on gfx908. (#1987)
2025-03-17 18:08:53 -07:00
10_convnd_fwd_multiple_d_multiple_reduce
Fix example_convnd_fwd_max_xdl_int8 failures on MI300 (#1666)
2024-11-14 08:40:50 -08:00
12_reduce
Support large: 12d tensor size for reduction kenrel (#1465)
2024-08-13 16:15:47 +02:00
13_pool2d_fwd
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
14_gemm_quantization
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
15_grouped_gemm
Fix example_grouped_gemm_multiple_d_xdl_fp16 on gfx950 (#2203)
2025-05-19 14:25:50 -07:00
16_gemm_multi_d_multi_reduces
Fix example_convnd_fwd_max_xdl_int8 failures on MI300 (#1666)
2024-11-14 08:40:50 -08:00
17_convnd_bwd_data
Add Grouped Conv Fwd Large Tensor kernel (#1432)
2024-08-06 10:06:10 +02:00
18_batched_gemm_reduce
remove support for gfx940 and gfx941 targets (#1944)
2025-03-05 11:07:33 -08:00
19_binary_elementwise
Refactor elementwise kernels (#1222)
2024-04-19 13:31:17 +02:00
20_grouped_conv_bwd_weight
Apply universal gemm to bwd_weight_cshuffle operator (#1873)
2025-02-18 10:10:22 +01:00
21_gemm_layernorm
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
22_cgemm
Clean DTYPES conditions in CMake (#974)
2023-10-18 11:14:14 -05:00
23_softmax
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
24_batched_gemm
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
25_gemm_bias_e_permute
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
26_contraction
Code clean-up (#1285)
2024-05-10 09:41:39 -07:00
27_layernorm2d_fwd
Split the static library into several files. (#1044)
2023-11-28 11:17:37 -08:00
28_grouped_gemm_bias_e_permute
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
29_batched_gemm_bias_e_permute
Merging the gfx12 code into public repo. (#1362)
2024-06-27 00:33:34 -07:00
30_grouped_conv_fwd_multiple_d
Remove CK_USE_AMD_MFMA_GFX950 (#1935)
2025-03-04 10:32:25 -08:00
31_batched_gemm_gemm
Merge from internal (#1857)
2025-02-07 15:05:05 -07:00
32_batched_gemm_scale_softmax_gemm
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
33_multiple_reduce
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
34_batchnorm
Refactor elementwise kernels (#1222)
2024-04-19 13:31:17 +02:00
35_splitK_gemm
Optimized GEMMs for MX FP4/8 (#2294)
2025-06-05 13:54:15 -06:00
36_sparse_embedding
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
37_batched_gemm_add_add_relu_gemm_add
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
38_grouped_conv_bwd_data_multiple_d
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
39_permute
Clean DTYPES conditions in CMake (#974)
2023-10-18 11:14:14 -05:00
40_conv2d_fwd_quantization
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
41_grouped_conv_conv_fwd
Merge from internal (#1857)
2025-02-07 15:05:05 -07:00
42_groupnorm_fwd
Split the static library into several files. (#1044)
2023-11-28 11:17:37 -08:00
43_splitk_gemm_bias_e_permute
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
44_elementwise_permute
Prevent instantiation of undefined FP8 operators. (#1639)
2024-11-05 13:58:29 -08:00
45_elementwise_normalization
Layernorm and groupnorm support to save mean and inverse std in forward (#929)
2023-10-19 07:36:29 +08:00
46_gemm_add_multiply
Code clean-up (#1285)
2024-05-10 09:41:39 -07:00
47_gemm_bias_softmax_gemm_permute
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
48_pool3d_fwd
Fixing most of the cppcheck errors. (#1142)
2024-01-24 13:47:48 -08:00
49_maxpool2d_bwd
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
50_put_element
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
51_avgpool3d_bwd
Fixing most of the cppcheck errors. (#1142)
2024-01-24 13:47:48 -08:00
52_im2col_col2im
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
53_layernorm2d_bwd
layernorm and groupnorm backward data (#1083)
2023-12-19 04:23:11 +08:00
54_groupnorm_bwd
layernorm and groupnorm backward data (#1083)
2023-12-19 04:23:11 +08:00
59_grouped_gemm_multi_ABD
OCP FP8 support for gfx12. (#1710)
2024-12-03 08:42:55 -08:00
60_gemm_multi_ABD
bf16A_Int8B with fastgelu/bias (#1264)
2024-04-26 07:26:30 -05:00
61_contraction_multi_ABD
add gemm_bias_add example (#1361)
2024-07-11 18:08:07 -07:00
62_convnd_activ
remove support for gfx940 and gfx941 targets (#1944)
2025-03-05 11:07:33 -08:00
63_layernorm4d_fwd
Split the static library into several files. (#1044)
2023-11-28 11:17:37 -08:00
64_fpAintB_gemm
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
65_gemm_multiply_multiply
Huaiguxu/moe fp8 pertoken scale fix (#2391)
2025-06-27 10:24:34 +08:00
66_complex_contraction_bilinear
chore: unset executable permission (#2303)
2025-06-10 09:13:59 -07:00
67_gemm_microscaling
MX GEMM - FP6 Example (#2419)
2025-07-07 10:33:26 -06:00
ck_tile
[CK_TILE] Low CU utilization optimization for fMHA fwd kernels (#2402)
2025-07-09 22:01:33 +08:00
CMakeLists.txt
Add MoE & FP8 Blockscale WP Kernels for GFX950 (#2297)
2025-06-12 09:25:59 +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