ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 10:37:44 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
9c6cca5bc5dde28359d33c6f9cd0a9a88396a845
composable_kernel/example
History
Po Yen Chen 9c6cca5bc5 [CK_TILE] FMHA FAv3 scheduling fine-tuning for performance (#2833) 
* Re-mapping thread block indices for causal=True kernels

* Use more intuitive remap_opt value

* Fallback to origin remapping if seqlen_q >= 64K

* Use GenericAttentionMask to reduce mask computation

* Avoid unnecessary boundary check for IsMasking=false case

* Fix wrong kernel entry specifier

* Add s_nop to prevent delay wave0-3

* Refine scheduling

* Remove unnecessary sched_group_barrier()

* Move sched_group_barrier() call to scheduler

* Replace inline asm s_setprio with intrinsics

* Rephrase comments

* Expend some o_acc rescaling insts to avoid SIMD idle

* Fix block idx special mapping logic

* Tune block index mapping for causal=False cases

* Tune block index mapping for causal=True cases

* Fix wrong vmcnt()

* Remove parameter name

* Use boolean option for turn on/off causal mask

* Update benchmark_fwd_v3.sh option usages

* Add option if compiler support it

[ROCm/composable_kernel commit: 7fbc9d6c97]
2025-09-16 11:32:38 +08:00
..
01_gemm
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07:00
02_gemm_bilinear
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
03_gemm_bias_relu
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -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
Revert "feature:tf32:add initial conv3d fwd kernel support (#2763)" (#2848)
2025-09-15 08:27:04 -07:00
10_convnd_fwd_multiple_d_multiple_reduce
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
12_reduce
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
13_pool2d_fwd
Refactoring cmake files to build data types separately. (#932)
2023-09-20 22:15:56 -07:00
14_gemm_quantization
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
15_grouped_gemm
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
16_gemm_multi_d_multi_reduces
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07: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
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07: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
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07: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
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
28_grouped_gemm_bias_e_permute
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -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
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
31_batched_gemm_gemm
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07:00
32_batched_gemm_scale_softmax_gemm
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
33_multiple_reduce
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
34_batchnorm
Fix build errors on windows (#2456)
2025-07-16 07:58:23 -07:00
35_splitK_gemm
WMMA support for GEMM reduce (#2823)
2025-09-12 21:36:43 +02:00
36_sparse_embedding
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
37_batched_gemm_add_add_relu_gemm_add
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
38_grouped_conv_bwd_data_multiple_d
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
39_permute
[CK][EXAMPLES] (#2826)
2025-09-11 12:33:00 -07:00
40_conv2d_fwd_quantization
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
41_grouped_conv_conv_fwd
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
42_groupnorm_fwd
[CK][EXAMPLES] (#2826)
2025-09-11 12:33:00 -07: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
[CK][EXAMPLES] (#2826)
2025-09-11 12:33:00 -07:00
45_elementwise_normalization
[CK][EXAMPLES] (#2826)
2025-09-11 12:33:00 -07:00
46_gemm_add_multiply
Code clean-up (#1285)
2024-05-10 09:41:39 -07:00
47_gemm_bias_softmax_gemm_permute
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07: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
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
61_contraction_multi_ABD
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
62_convnd_activ
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07:00
63_layernorm4d_fwd
upgrade from clang-format-12 to clang-format-18 (#2568)
2025-07-28 11:34:07 -07:00
64_fpAintB_gemm
Split the instances by architecture. (#1223)
2024-04-02 09:42:17 -07:00
65_gemm_multiply_multiply
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07:00
66_complex_contraction_bilinear
chore: unset executable permission (#2303)
2025-06-10 09:13:59 -07:00
67_gemm_microscaling
delete all slp compilation flag in CK Tile (#2625)
2025-08-06 00:34:39 -07:00
68_gemm_add
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
69_gemm_add_relu
Extend XDL kernel to Support RDNA3/4 - Part 4 (#2724)
2025-09-12 08:17:07 -07:00
ck_tile
[CK_TILE] FMHA FAv3 scheduling fine-tuning for performance (#2833)
2025-09-16 11:32:38 +08:00
CMakeLists.txt
Extend XDL kernel to Support RDNA3/4 - Part 5 (#2725)
2025-09-15 10:59:25 -07: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