ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-09 00:27:01 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,969 Commits 777 Branches 40 Tags
75550ff770cd3b49eaa6ddb81d22473caf5a6281
Commit Graph

1969 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Zoltan Lakatos
75550ff770 added test for gemm_add_multiply 2025-06-03 13:49:46 +00:00
Zoltan Lakatos
07f75d9c1d fixed template arguments for test instances 2025-06-03 13:49:46 +00:00
Zoltan Lakatos
63513c372b improved multiply_add profiler layout help 2025-06-03 13:49:45 +00:00
Zoltan Lakatos
8727762e74 fixed profiler help message for gemm_add_multiply 2025-06-03 13:49:45 +00:00
Zoltan Lakatos
538fa87141 initial gemm_add_multiply instance implementations 2025-06-03 13:49:45 +00:00
Zoltan Lakatos
b8e45c7dbe removed unnecessary ck parts from compilation 2025-06-03 13:49:42 +00:00
Anton Gorenko
959defbb22 Add a special wrapper to use DeviceGemmMultipleD_Wmma_CShuffleV3 with old API 
ckProfiler uses DeviceGemmMultipleD (tests also call its functions), the wrapper allows to use
DeviceGemmMultipleDSplitK instances there.
 2025-06-02 15:51:12 +05:00
Anton Gorenko
25f7204de2 Add gemm_add_fastgelu instances and test 2025-06-02 15:30:21 +05:00
Anton Gorenko
22935c87e9 Prepare gemma_add tests for adding wmma 2025-06-02 15:09:31 +05:00
Anton Gorenko
db51d8a531 Make gemm_add_add_wmma to work with DeviceGemmMultipleD_Wmma_CShuffleV3 2025-05-30 13:58:57 +05:00
Anton Gorenko
f13b91324a Implement DeviceGemmMultipleD_Wmma_CShuffleV3 2025-05-29 17:11:44 +05:00
Anton Gorenko
51512066ce Clone example/65_gemm_multiply_multiply/gemm_add_add_xdl_fp16.cpp for wmma 2025-05-29 15:59:48 +05:00
Anton Gorenko
a038ba3877 Clone for device_gemm_wmma_cshuffle_v3.hpp for future Multiple D support 2025-05-29 14:37:19 +05:00
Anton Gorenko
99fc05e481 Use ThreadGroupTensorSliceTransfer_v7r3 2025-05-29 13:10:27 +05:00
Anton Gorenko
da5f962dfd Support multiple D in GridwiseGemm_wmma_cshuffle_v3 
DeviceGemm_Wmma_CShuffleV3 is changed for new template parameters.
 2025-05-29 13:09:08 +05:00
Anton Gorenko
a6ea6040ad Add FP16 KM_NK and KM_KN test suites for XDL 
These tests were added to common .inc for better testing of WMMA instances
 2025-05-23 19:05:20 +05:00
Anton Gorenko
2679c0aa2b Disable all tests and instances with f8 on gfx11 
Even though f8_f16 and f16_f8 don't require f8 WMMA instructions,
gfx11 still lacks hardware instructions for fast f8->f32 conversion.
 2025-05-23 15:17:36 +05:00
Anton Gorenko
668914cd58 Add mixed precision examples 2025-05-23 10:53:42 +05:00
Anton Gorenko
c143bf30a5 Fix cache flushing for pk_i4 2025-05-23 10:53:42 +05:00
Anton Gorenko
9e70603bb5 Use PermuteB with f16_i4 in most instances (as xdl) 
Some instances use PermuteB = false for checking correctness.
See also the previous commit.
 2025-05-23 10:53:42 +05:00
Anton Gorenko
83b1419bb7 Fix a bug that fp16_i4 validation passes only with PermuteB 
A permutation required by conversion from pk_i4_t to half_t does not
depend on PermuteB, they can be used independently.
 2025-05-23 10:53:42 +05:00
Anton Gorenko
05ad2146f9 Add test_gemm_universal_xdl/wmma_fp16 to REGRESSION_TESTS 
They take more than 30 seconds
 2025-05-23 10:53:42 +05:00
Anton Gorenko
eac7d359e8 Add missing Regular tests 2025-05-23 10:53:42 +05:00
Anton Gorenko
dd7ac95aff Add test case for bf16_i4 2025-05-23 10:53:42 +05:00
Anton Gorenko
02bf56a6de Add missing add_device_gemm_wmma_universal_f8_f8_bf16 declarations 2025-05-23 10:53:42 +05:00
Anton Gorenko
b1f50b5ba7 Remove instances that are too slow (mostly because of register spilling) 2025-05-23 10:53:42 +05:00
Anton Gorenko
8b5d340d09 Fix segfault in profiler when B is pk_i4_t 
b_device_buf's size in bytes is larger than b_k_n_permute so b_device_buf.ToDevice reads out-of-bounds.
 2025-05-23 10:53:41 +05:00
Anton Gorenko
c7d39a075e Enable profiling of mixed precision with f8 and int4 on WMMA 2025-05-23 10:53:41 +05:00
Anton Gorenko
5082a9cd69 Add instances for Interwave and Intrawave v1 2025-05-23 10:53:41 +05:00
Anton Gorenko
342bb570bf Add a workaround for fp8/bf8->f32 packed conversion issue 2025-05-23 10:53:41 +05:00
Anton Gorenko
17bc0fa4c5 Remove instances that are too slow (mostly because of register spilling) 2025-05-23 10:53:41 +05:00
Anton Gorenko
04d3fc7b07 Add instances with ABlockLdsExtraM and BBlockLdsExtraN = 0 2025-05-23 10:53:41 +05:00
Anton Gorenko
c94c3b4716 Add instances for Interwave and Intrawave v1 2025-05-23 10:53:41 +05:00
Anton Gorenko
c414097e31 Implement Intrawave and Interwave variants of pipeline v1 2025-05-23 10:53:41 +05:00
Anton Gorenko
ede7126dc6 Clone BlockwiseGemmXdlops_pipeline_v1 for WMMA implementation 2025-05-23 10:53:40 +05:00
Anton Gorenko
af9e9eda8a Load in M/NRepeat dims with thread copy's slice instead of a loop 2025-05-23 10:53:40 +05:00
Anton Gorenko
cc818b4bc7 Fix strides of a/b_thread_desc 
The descriptors are larger than needed (even though the compiler don't alloc registers for unused values).
 2025-05-23 10:53:40 +05:00
apoorva
501d957152 Fixed the build errors. 2025-05-22 13:16:48 +05:00
apoorva
9d8f1e4c21 Fixed the gemm specialization error 2025-05-22 13:16:48 +05:00
apoorva
c016164040 Modified some code as per gemm_universal_examples 2025-05-22 13:16:48 +05:00
apoorva
b0aa9339e9 Fixed compilation failure. 2025-05-22 13:16:48 +05:00
apoorva
9cc5702e89 Fixing build erros 2025-05-22 13:16:48 +05:00
apoorva
4b5a9acc95 Fixing clang format 2025-05-22 13:16:48 +05:00
apoorva
840b79dfa3 Fixed review comments and added executable to cmakeLists 2025-05-22 13:16:48 +05:00
apoorva
8f8e631355 Added the missing parameters 2025-05-22 13:16:47 +05:00
apoorva
b35a19599c Adding examples for gemm_wmma instances 2025-05-22 13:16:47 +05:00
apoorva
621012cb86 example for gemm_universal_bf16 2025-05-22 13:16:47 +05:00
Anca Hamuraru
15bfa00312 Moved f16xi4 test under the correct data layout group 2025-05-20 14:42:09 +00:00
Anca Hamuraru
b1a9a27bf1 Merge branch '33-wip' of projects.streamhpc.com:amd/ai/composable_kernel into 33-wip 2025-05-20 13:47:26 +00:00
Anca Hamuraru
185ea0fdc4 Fix typo 2025-05-20 13:45:48 +00:00