ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 11:34:54 +00:00
Code Issues Packages Projects Releases Wiki Activity
2,400 Commits 779 Branches 40 Tags
e3fccf0115ed1efdc99d17d898d96a2fb35604ec
Commit Graph

2400 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
kiefer
e3fccf0115 Fixup comments 2025-09-16 09:07:38 +00:00
kiefer
eea8476738 No need to time kernels for now. 2025-09-15 09:02:45 +00:00
kiefer
b9986de7ae Remove old years for newly created files. 2025-09-15 08:48:35 +00:00
kiefer
521970ce2f Add newer instances to DEVICE_INSTANCES so the main ckProfiler can build 2025-09-05 13:03:23 +00:00
kiefer
e7314a17f4 Remove wmma fp8 and bf8 instances when not targetting gfx12 2025-09-05 09:08:49 +00:00
kiefer
b8d4b01267 Implement rotating memory and flush cache. Requires ad-hoc buffer size calculations. 2025-09-04 14:22:26 +00:00
kiefer
52c42d5b47 Clean up device implementation: remove old todos, remove unnecessary comments and print statements, tweak description, wrap all prints in env check. 2025-09-02 09:26:24 +00:00
kiefer
bcf9279c3e Unify problem sizes across vanilla and clamp flavor tests 2025-09-02 08:07:36 +00:00
kiefer
9416c82bfa Add 2D and 3D clamp instances and tests 2025-09-01 15:00:59 +00:00
kiefer
0b8de9a0dc Add 3D bias clamp instances and tests 2025-09-01 13:45:22 +00:00
kiefer
4822517535 Add 2D bias clamp instances and tests 2025-09-01 12:22:26 +00:00
kiefer
bc2c2fd54e Add some more tests for vanilla grouped conv fwd 2025-09-01 12:19:47 +00:00
kiefer
812b48548c Add comp and mem instances for vanilla 3D grouped conv fwd. Skipped 2x and part2 instances, can be added later in the same instance lists. 2025-08-31 13:17:40 +00:00
kiefer
63f52e0618 Add comp and mem instances for all vanilla 2D grouped conv fwd types. Skipping "x2" and "part2" instance lists, can be added later without special names if necessary. 2025-08-29 13:36:31 +00:00
kiefer
382d6fea6b Add D layout tests to IsSupportedArgument() 2025-08-29 08:22:42 +00:00
kiefer
78635fd74e Add 1D instances 2025-08-28 15:57:30 +00:00
kiefer
68f9e73b5e Implement multi-AB support for grouped conv fwd and add example. 2025-08-28 13:12:40 +00:00
kiefer
a06e27650e Add merged groups instances for all 2D vanilla grouped conv fwd types and layouts. 2025-08-27 10:18:13 +00:00
kiefer
73521fe091 Implement merged groups in device impl and add instances for merged groups 3D vanilla conv fwd 2025-08-27 08:02:51 +00:00
kiefer
ca7b3121cd Add int8 instances for 2D vanilla grouped conv fwd all layouts. 2025-08-26 12:16:41 +00:00
kiefer
e325dab094 Add instances for all 8-bit 3D vanilla grouped conv fwd types, including mixed types but with the exception of deprecated f16 comp fp8. Adapt test so we can test 8-bit and mixed types. 2025-08-26 09:20:38 +00:00
kiefer
6ad73cd0cd Add CTranspose optimization for NCHW cases just like in xdl cshuffle non-v3 device implementation. 2025-08-24 12:44:01 +00:00
kiefer
b53c584eb9 Disable sharding for the new instances for now, has tendency to lead to linker errors on repeat builds. 2025-08-24 12:23:46 +00:00
kiefer
f906c706fb Reset output buffer after each run in profile_grouped_conv_fwd_impl(). 2025-08-24 11:57:08 +00:00
kiefer
9089f2cb99 Add instances for all vanilla 2D and 3D flavors for f16 and bf16, only one instance per instance list to save compile time for now. Also added incomplete set of comp instances and bias_clamp for f16 2D, just to make sure the multiple-D aspects of the device implementation are working. 2025-08-20 14:01:02 +00:00
kiefer
4354cefbca Make relevant profilers print the number of valid instances to aid testing. 2025-08-20 10:48:41 +00:00
kiefer
43f99d85bd Initial device implementation for grouped conv fwd multiABD wmma cshuffleV3. Functional but needs some fixups and extra features in the future. 2025-08-20 10:43:49 +00:00
kiefer
16920dee0f Add support for fwd conv in gridwise implementation. Identical to run function for bwd data. 2025-08-20 08:56:53 +00:00
Enrico Degregori
ccf696ad4e Fix unused template parameter 2025-08-06 14:04:17 +00:00
Enrico Degregori
9b7cf3ef81 Fix bug in device print function 2025-08-06 14:03:55 +00:00
Enrico Degregori
9ee5699e50 multi_abd wmma support: 
- Add multiple A and B support to multiple D implementation (gridwise level)
 - Add multi_abd GEMM (device level)
 - Add instances (xdl parity)
 - Add tests (both xdl and wmma)
 - Add examples
 - Add ckProfiler support (both xdl and wmma)
 2025-08-06 13:54:56 +00:00
Enrico Degregori
8f01112ae0 Fix compilation error 2025-08-06 13:28:36 +00:00
kiefer
c434378570 clang-format-18 2025-08-06 11:53:43 +00:00
Kiefer van Teutem
ec382804a9 Merge remote-tracking branch 'origin/develop' into 90-prepare-an-upstream-pr-for-multipled-based-gemms 2025-08-06 07:47:43 +00:00
Illia Silin
833ae1d051 Revert "Reduce build time tile engine (#2579)" (#2623) 
This reverts commit e5b79b26fa.
rocm-7.1.0 		2025-08-05 09:27:55 -07:00
Enrico Degregori
2203b0ddfe Add padding to 1x1Stride1Pad0 conv specialization (grouped conv bwd weight) (#2610) 
* Add padding 1x1Stride1Pad0 conv specialization

* Add gridwise checks for conv cshufflev3

* Merge padding with previous transforms

* Apply transform changes for padding to default specialization as well

---------

Co-authored-by: Bartłomiej Kocot <barkocot@amd.com>
 2025-08-05 15:23:19 +02:00
Thomas Ning
cbfecf8d7a Persistent grouped gemm CompV4 Enablement & Polish (#2605) 
* enable the persistent kernel for CompV4

* polish the example and clang format

* fix the non-persistent kernel error

---------

Co-authored-by: ThomasNing <thomasning@amd.com>
 2025-08-04 23:43:01 -07:00
Max Podkorytov
2a78da4708 fix build for test_ck_tile_fp8 on rhel8 (#2615) 2025-08-04 17:43:15 -07:00
Illia Silin
fb96b49666 fix test_mx_mfma errors (#2614) 2025-08-04 11:43:47 -07:00
rahjain-amd
59245df46d Fix Debug Build for ckProfiler (#2609) 
Problem
=======
relocation R_X86_64_32 out of range: 5405348154 is not in [0, 4294967295]

Solution
========
The problem was caused due the limitation comes from the 32 bit offsets
used in original DWARF standard.
We have the option to switch to 64bit offset for your libs which free
us from 4G size boundary.

add -gdwarf64 and -Og to avoid this limit.
 2025-08-04 11:28:09 -07:00
Jinchao Xu
15eb493152 Add -gsplit-dwarf flag to reduce debug section size and fix ckProfiler link errors (#2611) 
Resolves R_X86_64_32 relocation out of range errors in grouped conv2d instances
by splitting debug information into separate .dwo files.

Add explicit cast to avoid signed/unsigned comparison warning.
 2025-08-04 11:26:08 -07:00
Bartłomiej Kocot
8655ba989c Mark non-grouped convolutions instances as deprecated (#2595) 
* Mark non-grouped convolutions instances as deprecated

* Update CHANGELOG.md

Co-authored-by: John Afaganis <john.afaganis@amd.com>

* Update library/src/tensor_operation_instance/gpu/conv1d_bwd_data/device_conv1d_bwd_data_xdl_nwc_kxc_nwk_bf16_instance.cpp

Co-authored-by: John Afaganis <john.afaganis@amd.com>

---------

Co-authored-by: John Afaganis <john.afaganis@amd.com>
 2025-08-04 16:49:55 +02:00
Enrico Degregori
02cb1f2b6a Fix bug from merge 2025-08-04 08:37:40 +00:00
Max Podkorytov
0d9439760f remove std::format (#2604) 2025-08-01 19:22:07 -07:00
Illia Silin
b786d12e56 remove std=c++17 compiler flag (#2603) 2025-08-01 16:18:16 -07:00
Max Podkorytov
f36cb5b2aa [CK-tile] remove old ck-tile transpose test (#2591) 
* remove old ck-tile transpose test

* rename test exe for consistency

* replace batched transpose regression test
 2025-08-01 14:50:09 -07:00
Thomas Ning
e5b79b26fa Reduce build time tile engine (#2579) 
* Modify CMakeLists to allow for splitting.

* Modify CMakeLists for data and layout logic.

* Run tests and get build artifact.

* Test new Cmakelists for speedup.

* Further improvements for speedup.

* turn off the FMHA

* turn off the automatic tile engine gemm

* minor fix

* disable the transpose test first

* Address the comment

* Jenkinsfile

* change the make thread to 64

* change the compile thread to 32

* Try to use with less OS memory space

* Have the Unity build batch size to 2

* reduce the chunk size

---------

Co-authored-by: Vidyasagar Ananthan <vidyasagar.ananthan@amd.com>
 2025-08-01 14:42:33 -07:00
Illia Silin
788e8a878e update the switch condition for buffer built-ins (#2602) 2025-08-01 14:30:07 -07:00
Thomas Ning
7c44a763fa Fix the GFX 950 Universal GEMM (#2597) 
* solve the gfx950 error

* clang format

* fix a typo error

---------

Co-authored-by: ThomasNing <thomasning@amd.com>
 2025-08-01 09:32:24 -07:00
Illia Silin
e6104daecc Add a daily CI stage to test AITER with latest CK. (#2598) 
* add a CI stage for AITER testing
 2025-08-01 07:55:51 -07:00