Commit Graph

402 Commits

Author SHA1 Message Date
aska-0096
2e9901bfa6 fix cmake bug 2024-11-27 10:11:13 +00:00
aska-0096
0dcd489802 delete use less files 2024-11-26 07:30:13 +00:00
aska-0096
26d5174e15 update instance and lds layout strategy 2024-11-26 07:29:38 +00:00
aska-0096
ea90b01fc9 fix bug in enable f8 gemm inside ckProfiler 2024-11-20 09:33:39 +00:00
aska-0096
98fc138582 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-20 05:41:36 +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
Illia Silin
da0c21f661 add more fp32 dl gemm instances (#1675)
* add more fp32 dl gemm instances

* update the dates
2024-11-19 10:00:17 -08:00
aska-0096
4132643a55 update tile size for fp8 rowwise 2024-11-19 11:28:37 +00:00
aska-0096
ec6b000c77 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-19 08:50:37 +00:00
Illia Silin
8aba2724cc Add bf16 and int8 wmma gemms for Navi3x and Navi4x. (#1671)
* add bf16 gemms for gfx11/gfx12

* reduce the input values in test_gemm

* add int8 wmma gemm instances for gfx11/gfx12

* add example gemm_wmma_int8

* fix bug in gemm_wmma_int8 test

* increase bf16 gemm test tolerance

* update the dates and clean-up commented-out instances
2024-11-18 14:07:04 -08:00
Bartłomiej Kocot
754adc70e3 Batched GEMM Multiple D based on Universal GEMM (#1655)
* Batched GEMM Multiple D based on Universal GEMM

Co-authored-by: Jing Zhang <jizhan@fb.com>

* CI fixes

Co-authored-by: Jing Zhang <jizhan@fb.com>

---------

Co-authored-by: Jing Zhang <jizhan@fb.com>
2024-11-18 14:03:45 +01:00
aska-0096
f3bbfe3efe Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-18 07:32:39 +00:00
Bartłomiej Kocot
ea3640fdea Add generic instances for two stage conv bwd wei (#1643)
* Add generic instances for two stage conv bwd wei

* Update layout prefix
2024-11-08 10:04:33 +01:00
Illia Silin
75c5bfa364 enable compilation for generic navi targets (#1645) 2024-11-07 14:14:42 -08:00
aledudek
dcafb1de15 Generic threshold calculation after merge fixes (#1618)
* Generic threshold calculation add passing num of accums

* Generic threshold - after merge fixes

* Fix cmakelists

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
2024-11-06 10:44:58 +01:00
Illia Silin
b6e74be1aa Make sure cmake can handle the xnack+/xnack- targets. (#1633)
* make sure cmake can handle xnack targets

* dont build xdl instances for gfx906:xnack-

* dont build xdl tests for gfx906:xnack-
2024-11-05 08:53:10 -08:00
aska-0096
925c07198f Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-05 10:15:54 +00:00
Juan Manuel Martinez Caamaño
464abd235e [generate.py] Override blob list if it already exists (#1635)
Before, generate.py appended the list at the end of the output file.
When running the cmake configuration steps multiple times on the
examples, the blob list (such as fwd_blob_list.txt) would grow at every
configuration.
`library/src/tensor_operation_instance/gpu/mha/CMakeLists.txt` worked around
this issue by removing the output file if it exists.

Now, generate.py overrides the content of the output file.
There is no need for the workaround in the CMakeLists.txt;
and the issue is solved for the example projects too.
2024-11-05 10:09:52 +01:00
aska-0096
f20e48f1f4 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into update_cka8w8 2024-11-05 07:03:42 +00:00
aska-0096
b97c68764e update ck_a8w8 library, update flush cache timing api 2024-11-05 06:57:48 +00:00
Lin Sun
0c9012fb70 Linsun/convint8 fwd instances (#1626)
Add instances for int8 grouped conv2d fwd
---------

Co-authored-by: root <root@dell300x-pla-t28-03.pla.dcgpu>
Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
2024-11-04 16:33:20 -08:00
Bartłomiej Kocot
4f1fdbb6e3 Temporary disable part of dynamic op conv instances (#1630)
* Temporary disable part of dynamic op conv instances

* fix
2024-11-04 13:34:17 -08:00
Illia Silin
03c6448ba3 Reduce build time. (#1621)
* disable fp8 gemm_universal on gfx90a and gfx908 by default

* fix cmake syntax

* fix clang format

* add ifdefs in amd_xdlops

* disable fp8 gemm instances on gfx90a by default

* update readme
2024-11-01 13:52:23 +08:00
Bartłomiej Kocot
31bf253aeb Add dynamic elementwise op (#1426)
* Add dynamic elementwise op

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>

* CI issues fix

* Custom parameter value for dynamic functions - Comments addressed

---------

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>
Co-authored-by: ThruptiRajLakshmanaGowda <tlakshma@amd.com>
2024-10-26 15:22:37 +02:00
valarLip
37f7afed1e add int8 gemm multiply multiply a8w8 (#1591)
* add int8 gemm multiply multiply a8w8

* uncomment

* clang-format-12

* Add example_gemm_multiply_multiply_xdl_int8

* Remove shell scripts

* update preprocess number for mi308; bring back printout in ckprofiler

* format

---------

Co-authored-by: chenjun <junchen2@amd.com>
Co-authored-by: Haocong WANG <haocwang@amd.com>
Co-authored-by: carlushuang <carlus.huang@amd.com>
2024-10-26 16:39:34 +08:00
Rostyslav Geyyer
7d576f1748 Update GPU verification (#1596)
* Update inits

* Update static_cast to type_convert

* Add verification option selection
2024-10-25 08:13:46 -07:00
aledudek
9385caa306 Generic threshold calculation (#1546)
* Calculate generic relative threshold pool3dfwd

* Calculate absolute error threshold pool3d fwd

* Generic threshold calculation take max input for relative error pool3dfwd

* Remove max possible value for error calculation at runtime

* Remove debug print in pool3dfwd

* Pool3d fwd adjusted types in generic threshold calculation

* Generic threshold calculation take into account number of accumulations and accdatatype

* Generic threshold fix final error formula

* Generic threshold calculation - num of accs fix

* Generic threshold calculation - adjust absolute error

* Generic threshold calculation - OutDataType in absolute error
2024-10-25 12:46:24 +02:00
Bartłomiej Kocot
cedccd59c9 [POST MERGE PR] Enable grouped conv bwd wei bf16 NGCHW (#1594) 2024-10-23 12:02:33 +02:00
Haocong WANG
47294b4b22 Merge branch 'develop' into gemm_multiply_multiply_int8a8w8 2024-10-23 11:28:40 +08:00
Bartłomiej Kocot
82fc53835a Enable grouped conv bwd wei bf16 NGCHW (#1589)
* Enable grouped conv bwd wei bf16 NGCHW

* fixes

* fixes

* Fixes

* fixes

* fixes

* Fixes
2024-10-22 16:18:28 +02:00
chenjun
1670bba95f clang-format-12 2024-10-21 23:16:04 +08:00
Thomas Ning
560917b161 Ck profiler instance support (#1575)
* The draft on ckProfiler instance add

* support the ck profiler instance with same data types

* add a small feature on the M and N variable switch.

* Partially solve the incorrect result problem

* fix based on ci cd
2024-10-21 22:47:48 +08:00
chenjun
7fb0b3223c add int8 gemm multiply multiply a8w8 2024-10-21 21:57:41 +08:00
Haocong WANG
a285d6f9b5 disable bad instance detected on MI308CPX (#1584) 2024-10-18 08:46:11 -07:00
Rostyslav Geyyer
aa932445ea Add a gpu gemm reference kernel (#1528)
* Add a gpu gemm reference kernel

* Switch to gpu reference in gemm examples

* Remove redundant arguments

* Update all related examples

* Update more examples

* Try less threads per block

* Try even less threads per block

* Add support for all matrix layouts

* Increase block size

* Clean up

* Remove hardcoded strides

* Clean up

* Try a column-major case

* Revert back to row-major

* Run both CPU and GPU veriffication

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
2024-10-08 11:05:28 -05:00
Illia Silin
7d8ea5f08b Fix build logic using GRU_ARCHS. (#1536)
* update build logic with GPU_ARCHS

* fix the GPU_ARCHS build for codegen

* unset GPU_TARGETS when GPU_ARCHS are set
2024-10-07 08:18:23 -07:00
BrianHarrisonAMD
294cb82314 Add generating mha static library for gfx90a (#1540)
* Add generating mha static library for gfx90a

* Update comment to reflect changes
2024-10-02 09:26:11 -07:00
BrianHarrisonAMD
3528a523ff Add additional instances to device_mha_instance (#1522)
* Add additional instances to device_mha_instance

* Add comment to describe what receipt 3 option filters

---------

Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
2024-09-24 10:15:30 -06:00
Bartłomiej Kocot
4ba52b35dc Add support for NGCHW in grouped conv fwd (#1499)
* Support NGCHW in grouped conv fwd

* Remove not needed variable

* Fixes
2024-09-20 10:45:46 +02:00
aledudek
a793afc961 Extend pool3d fwd avg, max operations by f8_t, int8_t types (#1483)
* Extend pool3d fwd avg, max operations by f8_t, int8_t types

* Pack MaxPool3dFwd params together

* Fix MaxPool3dFwd AVG instances

* Decrease verification precision for bf16

* Adjust tests + review changes

* Adjust threshold for F8

* Adjusted compute types for MAX op instances

* Fix ComputeDataType mismatch in tests and profiler for AVG

* Fix naming from max_pool3d_fwd to pool3d_fwd

* Adjust CMakeLists

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
2024-09-17 15:57:10 +02:00
Mateusz Ozga
6834e5ee74 This commit contains implementation of max pool2d for f8 type (#1506)
Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
2024-09-16 10:15:06 +02:00
bibek
49e012dee1 Fix duplicate CMake tidy-target issue (#1513) 2024-09-13 21:15:04 -07:00
jakpiase
8f8a2ce396 Add pool2d int8 and fp8 instances (#1508)
* add pool2d fp8 and int8

* minor fixes

* add formatting

* add reviewer suggestions

* add reviewer suggestions
2024-09-13 10:18:21 -07:00
Jun Liu
81bc1496b2 Customize filesystem in CK for legacy systems (#1509)
* Legacy support: customized filesystem

* Update cmakefile for python alternative path

* fix build issues

* CK has no boost dependency

* More fixes to issues found on legay systems

* fix clang format issue

* Check if blob is correctly generated in cmake

* fix the python issues

* add a compiler flag for codegen when using alternative python

* use target_link_options instead of target_compile_options

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
2024-09-13 07:51:07 -07:00
Mateusz Ozga
448c0f56d8 Pool2d max/avg kernel in the BWD version (#1494)
* Add pool2d instance BWD AVG

* Add pool2d instance BWD MAX

* Fix: avg review

* Fix review: part2

* Fix - enable test when type is compiled

* Fix review part3
2024-09-12 11:47:52 +02:00
jakpiase
e8d2887cb2 Rewrite pool2d fwd (#1462)
* added pool2d fwd

* add tests

* add reviewers changes

* Revert "Merge remote-tracking branch 'origin/develop' into jakpiase/pool2d_fwd_new"

This reverts commit 6b2ba7ff89, reversing
changes made to 22c82bea0c.

* Revert "add reviewers changes"

This reverts commit 22c82bea0c.

* added reviewers comments

* revert some old files

* add reviewers requests

---------

Co-authored-by: Adam Osewski <19374865+aosewski@users.noreply.github.com>
2024-09-11 15:21:00 +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
81fb545fa7 tempsave 2024-09-05 12:46:55 +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