ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-16 16:51:26 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,705 Commits 779 Branches 40 Tags
aa15c49a6794de5436201b1516492e887899f652
Commit Graph

88 Commits

Author SHA1 Message Date
coderfeli
965c9f0c17 debug 16x16 load 2025-02-07 08:17:33 +00:00
coderfeli
f9abcf80e8 use offsets in transfer ok 2025-02-07 01:54:12 +00:00
coderfeli
f8d15f2af4 add others 2025-02-04 03:05:58 +00:00
coderfeli
00627feda4 results ok 2025-02-04 03:05:17 +00:00
coderfeli
6b51413b6e compile ok 2025-01-24 08:36:49 +00:00
aska-0096
487a05d612 refine blockgemm pipeline version as base struct. 2025-01-08 14:27:42 +00:00
aska-0096
22fe522d0c optimize software pipeline 2025-01-08 09:28:32 +00:00
aska-0096
0dbe537032 refine weight preshuffle format. 2025-01-02 13:59:58 +00:00
aska-0096
72c1ddacb9 Merge branch 'add_a8w8_preshuffle_ckprofiler' of https://github.com/ROCm/composable_kernel into update_cka8w8_uc 2024-12-31 07:23:50 +00:00
aska-0096
6f24c2d814 disable N, K Padding, splitk enabled 2024-12-31 06:31:06 +00:00
aska-0096
f60f9d5917 sanity pass, most tile size enabled. TODO: NWave!=4 2024-12-30 18:22:08 +00:00
coderfeli
db84352941 fix warnings and revert cmake and fix clang format 2024-12-30 08:24:50 +00:00
coderfeli
3f9dbcac63 use new pipeline for b preshuffle, run ok; revert olds to fix ckprofiler 2024-12-30 06:52:10 +00:00
coderfeli
54f44e6232 fix brepeat, kloop and lds two buffer; works ok now 2024-12-30 00:25:46 +00:00
coderfeli
2c056624af fix tail 2024-12-27 08:30:03 +00:00
coderfeli
174b46b04a add cpu shuffle 2024-12-27 07:31:14 +00:00
coderfeli
031ddf356d fix performance regression on blockgemm v3 pipe 2024-12-27 06:40:43 +00:00
aska-0096
7cec63a631 remove agpr usage when vgpr usage <256 2024-12-27 03:09:26 +00:00
coderfeli
e6f5a78b14 add double buffer scratch 2024-12-26 15:02:04 +00:00
coderfeli
3784329b68 can run 2024-12-26 13:01:07 +00:00
coderfeli
4a1ec81595 add bypass logic and build 2024-12-26 10:05:25 +00:00
aska-0096
1a089f6f88 sanity bug fix 2024-12-26 10:05:17 +00:00
aska-0096
c8c016ddad Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-12-13 09:18:50 +00:00
Adam Osewski
bfe983a151 Change block gemm pipeline local prefill loop order. (#1692) 
* Fix loop order.

* Fix loop order in pipeline v4
 2024-11-26 17:36:53 +01:00
aska-0096
ea90b01fc9 fix bug in enable f8 gemm inside ckProfiler 2024-11-20 09:33:39 +00:00
aska-0096
c99e3d595e Merge branch 'mem_gemm_opt' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-20 05:41:33 +00:00
aska-0096
2b840f5a85 reduce prefetch stage in blockwisepipev4 2024-11-18 07:32:30 +00:00
Illia Silin
922e42a039 fix compilation errors for gfx12 with clang20 (#1606) 2024-10-28 19:02:48 -07:00
macurtis-amd
aeb7c91f48 Fix compilation errors generated by forthcoming Clang changes (#1544) 
Without this change, the following diagnostic is generated:
  a template argument list is expected after a name prefixed by the template
  keyword [-Wmissing-template-arg-list-after-template-kw]

See C++17 spec [temp.names] p5.
 2024-10-02 13:56:22 -07:00
Illia Silin
42e6dceacc Fix compilation errors with Clang20.0. (#1533) 
* fix clang20 compilation errors for gfx90a

* fix clang20 compilation errors for gfx11 targets
 2024-09-25 13:45:38 -07:00
jakpiase
2a261afcdf Added structural sparsity blockwise gemm (#1435) 
* Implemented smfmac xdlops

* Added smfmac blockwise xdlops

* fixes

* add reviewers suggestions

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
 2024-09-11 15:19:42 +02:00
Haocong WANG
0b3a409d4f Merge branch 'develop' of https://github.com/ROCm/composable_kernel into mem_gemm_opt 2024-09-06 03:22:06 +00:00
Haocong WANG
5b10dae6a4 Add gemm universal bf16 instances (#1484) 
* revert ckprofiler change

* temp save

* Add test and test pass

* test pass

* Fix bug inside rotating buffer when tensor is not packed

* bug fix

* clang format

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
 2024-09-04 20:58:54 -07:00
aska-0096
dbfcb380cd temp save 2024-09-05 03:04:31 +00:00
aska-0096
41fcfbc64e clang format 2024-09-04 15:11:51 +00:00
aska-0096
6df91708a6 temp save 2024-09-04 14:32:02 +00:00
aska-0096
4885c38aa4 Merge branch 'transpose_opt' of https://github.com/ROCm/composable_kernel into rowwise_opt 2024-09-03 08:37:45 +00:00
aska-0096
5d9c964ece temp save 2024-09-02 09:55:07 +00:00
aska-0096
cbf14ee192 tempsave, epilogue optimization for universal gemm done. TODO: mulitpleD epilogue optimization 2024-09-01 08:10:43 +00:00
aska-0096
4f65f7b387 tempsave 2024-08-22 07:30:46 +00:00
Haocong WANG
8c90f25be3 [GEMM] F8 GEMM, performance optimized. (#1384) 
* add ab_scale init support

* enabled interwave

* add scale type; update isSupport

* adjust example

* clean

* enable f8 pure gemm rcr ckprofiler

* Add gemm_multiply_multiply instances

* clang format

* Optimize for ScaleBlockMNK=128

* enable abscale f8 gemm ck profiler

* Add pure f8 gemm test suite

* Reverting to the state of project at f60fd77

* update copyright

* clang format

* update copyright

---------

Co-authored-by: root <jizhan@amd.com>
 2024-07-19 22:06:52 +08:00
Illia Silin
941d1f7ce0 Merging the gfx12 code into public repo. (#1362) 2024-06-27 00:33:34 -07:00
Bartłomiej Kocot
510325a468 Fix cmake warnings (#1342) 
* Cmake add -Wno-nvcc-compt

* Remove template without initialization list

* dpp remove template without init list

* Fixes
 2024-06-21 09:47:58 +02:00
jakpiase
e2d139201b Switch to universal gemm in grouped gemm tile loop (#1335) 
* switch to universal gemm in grouped gemm tile loop

* minor fixes

* add reviewers comments

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
 2024-06-18 09:01:49 -05:00
zjing14
80db62f08d add f8 gemm multiD with both row/col wise scale (#1300) 
* add f8 gemm with multiD for both row/col wise

* change compute_type to fp8

* changed tuning parameters in the example

* add rcr example
 2024-05-28 12:04:22 -05:00
Illia Silin
c44137838e remove wrong use of nonexistent class members (#1290) 2024-05-15 08:08:17 -07:00
Haocong WANG
764164b488 [GEMM] UniversalGemm update (#1262) 
* Add bf16 instances

* Add bf16 gemm universal example

* tempsave

* Add guard to navi compilation

* workground on a specific mixed gemm instance ( bring back it when compiler fix upload)

* fix formatting condition statement issue

* solve conflict

---------

Co-authored-by: Jun Liu <Liu.Jun@amd.com>
 2024-04-26 12:56:07 -05:00
zjing14
0d0150db20 bf16A_Int8B with fastgelu/bias (#1264) 
* changed the copy function to v7r2

* adding multi_abd

* in-progress

* add post-load oob check

* debugging

* adjust instances

* add run_lds

* add elemntwise_op

* replace multi_abd_device with v3

* clean up

* clean

* clean

* Added LDSType

* profiling

* adjust oobcheck

* add missing file

* refactor

* clean

* add examples
 2024-04-26 07:26:30 -05:00
Haocong WANG
f83e9701e9 [GEMM] Gemm universal device operation (#1154) 
* Optimize GEMM on MI200/300:
1. Add new blockwise gemm pipeline
2. Add irregular splitk intances

* clang format + typo fix

* Fix a bug

* initial commit

* Add more instances to irregular splitk

* blkgemm pipeline v1~4 prototype

* Sanity Checked. Known issue:
1. Poor performance of splitk
2. Register spill on blkgemmpipeline v3

* Sanity and Performance fix:
1. fix a bug related to sanity in grouped b2c mapping
2. fix a bug related to sanity and performance in splitk offset

* Sanity and API update:
1. Remove prefetch stage
2. Fix valid check bug
3, Add first gemm_universal instance into ckProfiler

* Add NN instances for gemm universal

* 1. Add NT instances for gemm_universal
2. Fix a bug about Kpadding in gemm_universal

* Fix a bug regarding padding Odd K number

* remove kernel print

* Fix KPadding bug...

* Update safety check

* another try to fix kpadding..

* Sanity checked

* new instances..

* clang format+typo fix

* remove clang format script's change

* Add non-hotloop compile option

* 1. Add fp16xfp8 example
2. pull packed convert f8 from pr1150

* Some miscs.. opt and fix

* Add pipeline description docs

* Split universal gemm instance library to cut profiler compiling time

* uncomment cmakefile

* Fix a bug caused by blockwise_gemm_pipe_v2

* reduce default splitk to 1

* Add 224x256x64 tile size

* update, including:
1. Experiment pipeline 5~7
2. Optimization for pipeline 4
3. Organized instance library

* temp save

* temp save

* Permuted lds layout, sanity and function checked

* clang format

* Move OOB check from RunRead to RunWrite, for better software pipeline.
TODO: agpr spill when NN layout

* clangformat

* A/B splitpipe scheduler for v3

* Fix two bugs

* bug fix

* fix a bug in oob check

* Example for mixed fp16_fp8 gemm

* Clean experimental code blocks

* Add mixed precision gemm into profiler

* tempsave

* optimize m/n major lds layout

* Add RRR GEMM  mixed precision instances

* Optimize f8 matrix transpose

* Add test_gemm_universal

* A/B spilt schedule for blkpip v5

* Take ds_read2 into iglp scheduling scheme

* format

* fixed cmake

* Add llvm-option into CI cmake flag

---------

Co-authored-by: Jing Zhang <jizhan@amd.com>
 2024-04-13 21:03:18 -05:00
Bartłomiej Kocot
9c052804a7 Add elementwise with dynamic vector dim (#1198) 
* Add elementwise with dynamic vector dim

* Reduce number of instaces

* Fixes

* Fixes
 2024-03-22 10:40:43 +01:00