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
Anca Hamuraru
ae21582a88
Updated the set of tests for FP16
2025-05-20 13:45:15 +00:00
Anca Hamuraru
b350bd212d
Updated the set of tests for FP16
2025-05-20 13:16:48 +00:00
Anca Hamuraru
0a5e6d48ec
Fixed typos
2025-05-20 13:16:12 +00:00
Anca Hamuraru
a08ca63b36
Padding not supported for when BDataType is pk_i4_t. Added fix for correct check and removed padding instances.
2025-05-20 13:14:07 +00:00
Anca Hamuraru
dc26ee3014
Added tests for fp16, KM_KN and KM_NK
2025-05-20 07:47:07 +00:00
Anca Hamuraru
7975c9cea6
Fixed typos for padded instances
2025-05-19 14:51:56 +00:00
Anca Hamuraru
a0a2bf27b1
Fixed typos for bf16xbf16xbf16 padding instances
2025-05-19 13:08:05 +00:00
Anca Hamuraru
bbda71f096
Added padding instances for f16xf8xf16, f16xi4xf16
2025-05-16 08:31:32 +00:00
Anca Hamuraru
e97a7f246d
Added padding instances for f8xf16xf16
2025-05-15 11:31:50 +00:00
Anca Hamuraru
537d19c459
Added packed instances for bf16xi4xbf16
2025-05-15 10:36:50 +00:00
Anca Hamuraru
a3d99c696d
Merge branch '32-add-the-remaining-combination-of-data-types-mixed-precision-for-gemm_universal-pipeline3-for' into '33-wip'
...
Added mixed precision instances for gemm_universal
See merge request amd/ai/composable_kernel!11
2025-05-15 06:53:55 +00:00
Anca Hamuraru
0482b83fe6
Added padding instances for bf16xbf16, f8xf8
2025-05-14 14:36:10 +00:00
Anca Hamuraru
e2855633ee
Fixes for padding instances for f16xf16xf16
2025-05-14 10:50:34 +00:00
Anca Hamuraru
ab28ac0b81
Merge branch '33-wip' of projects.streamhpc.com:amd/ai/composable_kernel into 33-wip
2025-05-14 09:43:36 +00:00
Anca Hamuraru
647024db6f
Ammending changes for adding support for padding instances for f16xf16xf16
2025-05-14 09:41:19 +00:00
Anca Hamuraru
7107bccda3
Fixed cmake build errors related to test_fp8
...
(cherry picked from commit 12176616b6anca@streamhpc.com >
2025-05-14 08:46:34 +00:00
Anca Hamuraru
527c844a06
Merge branch 'cherry-pick-5bdc993d' into '33-wip'
...
Fixed cmake errors related to gemm_bilinear. Previously, if the above flags...
See merge request amd/ai/composable_kernel!10
2025-05-14 08:46:15 +00:00
Anca Hamuraru
8b52033aae
Fixed cmake errors related to gemm_bilinear. Previously, if the above flags are set, cmake build fails: GPU_TARGETS="gfx1100;gfx1201" -D DTYPES="fp16;bf16;fp8"
...
(cherry picked from commit 5bdc993dbfanca@streamhpc.com >
2025-05-14 08:45:34 +00:00
Anca Hamuraru
dd47f39a0a
Added padding instances for f16xf16xf16
2025-05-14 07:54:46 +00:00
Thomas Ning
9d1e44e56a
Vectorized Transpose for Batched Transpose CK Tile Operator ( #2131 )
...
* Shared Memory for single data point
* CKTile Transpose vectorize CP1
* CKTile Transpose vectorize CP2
* CKTile Transpose vectorize CP2.1
* fixed the compile error of the transpose tile 2d
* Have the correct result for the current test sample
* Changes to printing tensor
* fp8 support added
* Debugging for transpose
* solving the corner issue
* Changed padding flag
* Intermideate Debugging
* Intermidiate Debugging
* Intermediate Debugging
* Finished debugging of the transpose op
* Code Cleanup
* Adding edge case smoke tests
* Adding Transpose test to CI/CD
* Adding Transpose test to CI/CD
* Adding Transpose test to CI/CD
* Addressing Review Comment
* Addressing Comments
* Addressing Comments
* Measuring Perf Tests
* Code Cleanup
* Changlog
* Added the running iterations
* clang format
* Fix the changelog
* Fix the compilation error
* change the printing factor
---------
Co-authored-by: ThruptiRajLakshmanaGowda <tlakshma@amd.com >
2025-05-12 00:41:45 -07:00
Khushbu Agarwal
d8faf1c6a1
Support for swizzle and transpose for MFMA_16x16x32_F16/BF16 ( #2172 )
...
* Changes for updating tile distribution for shuffle and transpose
* Fixed swizzle and transpose, removed comments
* clang formatted
* Adding support for bf16 type
* Addressing review comments
2025-05-10 22:40:05 -07:00
Bartłomiej Kocot
6fddb5708c
Add grouped conv fwd bias relu instances ( #2179 )
...
* Add grouped conv fwd bias relu instances
* fixes
* fix
2025-05-09 22:52:34 +02:00
jefyang1
6b1a339b6f
Fix grouped conv bwd data tests on gfx950 ( #2173 )
2025-05-09 09:01:06 -07:00
