Commit Graph

587 Commits

Author SHA1 Message Date
Enrico Degregori
ccf696ad4e Fix unused template parameter 2025-08-06 14:04:17 +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
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
Aviral Goel
546ef78d1d Disable fp8 instances on unsupported targets (#2592)
* Restrict building of gemm_universal_preshuffle_f8 instances to specific targets in CMakeLists.txt

* Add condition to skip gemm_xdl_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt

* Add conditions to skip unsupported targets for gemm_universal_preshuffle_f8 and gemm_xdl_universal_preshuffle_f8 instances in CMakeLists.txt

* Refine conditions to exclude gemm_universal_preshuffle_f8 instances for unsupported targets in CMakeLists.txt

---------

Co-authored-by: AviralGoelAMD <aviralgoel@amd.com>
2025-07-31 12:18:02 -07:00
Illia Silin
b80099cc5f Revert "Add gemm universal f8 f8 bf16 mk nk instances on gfx950 (#2558)" (#2584)
This reverts commit c64a0c65b9.
2025-07-29 13:04:51 -07:00
Illia Silin
504b101da3 upgrade from clang-format-12 to clang-format-18 (#2568)
* upgrade to clang-format-18

* update to clang-format-18 in pre-commit-config
2025-07-28 11:34:07 -07:00
Enrico Degregori
5dc21c5521 Merge branch 'develop' into feature/multiple-d-gemms 2025-07-28 17:18:18 +00:00
jefyang1
c64a0c65b9 Add gemm universal f8 f8 bf16 mk nk instances on gfx950 (#2558) 2025-07-28 09:03:54 -07:00
Enrico Degregori
b01a27ff22 Support b_scale: (#2350)
- extend pipeline v1 and v3
 - add instances
 - add tests
 - add example

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
2025-07-24 18:49:58 -07:00
Yi DING
4338346b10 Use filename but not path to filter compilation (#2556) 2025-07-24 17:38:14 +08:00
Mingtao Gu
0198257d79 [CK] Fixed MPerBlock=32 build issue for MXFP4 GEMM decode (#2512)
* added MPerBlock=32 for MXFP4 GEMM decode

* added two instance for M>128 scenario.

* added 1 instance

* format

---------

Co-authored-by: mtgu0705 <mtgu@amd.com>
Co-authored-by: felix <felix.li@amd.com>
2025-07-18 14:35:54 +08:00
linqunAMD
fbd9f32abe [CK][CONV] Support NCHW in class DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 (#2459)
1. Port NCHW support from ConvFwd (#2375) to conv bwd data
2. Add new instance device_grouped_conv_bwd_data_xdl_f16_nchw_instances for nchw

Co-authored-by: azhuang <anzhong.huang@amd.com>
2025-07-17 08:19:57 +08:00
Zoltan Lakatos
161fe6c9d2 fixed cmake comment 2025-07-14 12:04:11 +00:00
Zoltan Lakatos
e2a75d6653 Merge remote-tracking branch 'origin/feature/multiple-d-gemms' into 8-implement-device_gemm_add_multiply-for-rdna4 2025-07-14 11:59:26 +00:00
Apoorva Kalyani
27c0f95552 Merge branch '79-add-instances-and-examples-for-device_gemm_add_relu' into 'feature/multiple-d-gemms'
Resolve "Add instances and examples for device_gemm_add_relu"

See merge request amd/ai/composable_kernel!29
2025-07-14 11:46:24 +00:00
Apoorva Kalyani
9c1314de6d Merge branch '51-create-bf16-and-f16-instances-for-gemm_add-cshuffle_v3-for-rdna4' into 'feature/multiple-d-gemms'
Resolve "Create bf16 and f16 instances for gemm_add CShuffle_v3 for RDNA4"

See merge request amd/ai/composable_kernel!17
2025-07-14 11:45:57 +00:00
Andriy Roshchenko
25b359d630 MX GEMM - Add FP6 GEMM Test (#2488)
* Add F6 GEMM MX Test

* Add BF6 GEMM MX Test
2025-07-11 15:32:12 -06:00
Zoltan Lakatos
41d4500509 Merge remote-tracking branch 'origin/feature/multiple-d-gemms' into 8-implement-device_gemm_add_multiply-for-rdna4 2025-07-11 13:27:04 +00:00
Illia Silin
1b66f3f4a3 Add declarations for atomic add for fp16 and unsigned short. (#2483)
* add template for fp16 atomic add

* add template for unsigned short atomic add

* use atomicCAS in atomic add for fp16 and unsigned short

* revrt back to atomic add using casting
2025-07-10 07:18:56 -07:00
apoorva
e6ea4aaf6d Fixed typo. 2025-07-09 10:55:21 +00:00
apoorva
76f4bb0e60 Added missing wrapper instances 2025-07-09 09:22:02 +00:00
apoorva
32125077e7 Fixed the if condition MACROS. 2025-07-09 08:03:09 +00:00
apoorva
55299c924e Fixed the review comments 2025-07-09 07:33:14 +00:00
Illia Silin
93420ecf89 Revert "Add templates for fp16 and unsigned short atomic add to fix FBGEMM bu…" (#2474)
This reverts commit 112b47e885.
2025-07-08 19:01:26 -07:00
Illia Silin
112b47e885 Add templates for fp16 and unsigned short atomic add to fix FBGEMM builds. (#2471)
* add template for fp16 atomic add

* add template for unsigned short atomic add

* use atomicCAS in atomic add for fp16 and unsigned short
2025-07-08 18:09:30 -04:00
apoorva
516d1f5194 Updated wrapper for the v3 instances 2025-07-08 13:55:24 +00:00
apoorva
84b0b324cf Removed the old wmma instances. 2025-07-08 13:51:48 +00:00
apoorva
669befb25a Updated copyrights and added wrappers. 2025-07-08 12:06:38 +00:00
apoorva
9b64da2298 Added wrapper and renamed the wmma_v3 instances 2025-07-08 11:26:01 +00:00
apoorva
86ca6b827d Removed the old wmma instances. 2025-07-08 11:23:33 +00:00
Aviral Goel
36df1cbd0a [ckProfiler] Add infrastructure and instances to profile gemm_universal with B preshuffle (#2427)
* works on mi300

* fix(profiler): add error message for unsupported type/layout

* refactor(preshuffle.inc): add type aliases for code readability
2025-07-01 18:34:52 -07:00
apoorva
6ec0ad2758 Added test for gemm_add_relu wmma instance 2025-07-01 13:44:18 +00:00
Zoltan Lakatos
ff31873a19 addressed core review comments 2025-07-01 12:47:39 +00:00
apoorva
f5843dd22b Added v3 instances for gemm_add_relu 2025-07-01 12:37:46 +00:00
apoorva
6a116fa958 Modified the template parameters to make the instances work. 2025-07-01 11:23:44 +00:00
apoorva
cdaff7f210 Added instances to Cmake 2025-07-01 11:23:43 +00:00
apoorva
6f89183179 Added f16 wmma instance and corrected bf16 instance errors. 2025-07-01 11:23:17 +00:00
apoorva
35aab35d96 Added bf16 wmma instance for add_relu 2025-07-01 11:23:17 +00:00
Zoltan Lakatos
eaa0452b80 Merge remote-tracking branch 'origin/feature/multiple-d-gemms' into 64-implement-device_gemm_multiply_multiply_instance-for-rdna4 2025-06-30 11:06:09 +00:00
Zoltan Lakatos
6ba1dc66ac Merge remote-tracking branch 'origin/feature/multiple-d-gemms' into 8-implement-device_gemm_add_multiply-for-rdna4 2025-06-30 11:03:13 +00:00
huaiguxu
e1c5172fdb Huaiguxu/moe fp8 pertoken scale fix (#2391)
* fix pertoken_scale a_scale dimension

* clang-format

* Fix moe_gemm2_fp8 perTokenScale reference and example.
2025-06-27 10:24:34 +08:00
Zoltán Lakatos
686df332e2 Resolve "Implement device_gemm_bilinear for RDNA4" 2025-06-26 06:48:38 +00:00
linqunAMD
1749c0409e [CK][CONV] Support NCHW in class DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle (#2375)
1. When conv spec is 1x1 stride1 pad0, nchw is equal with matrix A + column major, we only need minor change in conv transformer to support it.
2. when out is NKHW, it is equal with matrix C with column major. we need swap A & B to get best performance.
3. Add new instance device_grouped_conv_fwd_xdl_f16_nchw_instances for nchw.
2025-06-26 08:32:39 +08:00
Zoltan Lakatos
3c3136be79 added fp8 instances 2025-06-24 16:05:42 +00:00
JonathanLichtnerAMD
87fdb368a7 Do not build "other" library for MIOpen (#2382)
MIOpen only needs the static CK library for convolutions.
2025-06-24 07:32:16 -07:00
Kiefer van Teutem
9e74ae7c89 Implement batched gemm wmma (RDNA batched gemm) based on wmma cshuffle v3 (#2319)
* Some prep work for adding batched_gemm_wmma_universal. Moved batched_gemm in general to gfx11 and gfx12 categories, and split existing batched_gemm test into xdl and wmma versions. Updated profiler and instance factory. For now only adding f16-row-row-row-GemmDefault. For now actual device instance list is empty.

* Add DeviceBatchedGemm_Wmma_CShuffleV3 based on DeviceGemm_Wmma_CShuffleV3 and make sure it's used in the instance factory and tests. Currently the new batched device level struct cannot actually handle batching, but it does pass tests with a trivial batch size of 1, meaning that the overall structure is good.

* Add custom kernel and Argument type to DeviceBatchedGemm_Wmma_CShuffleV3. Batching arguments not passed to kernel yet.

* Implement kernel-level batching logic for DeviceBatchedGemm_Wmma_CShuffleV3.  In principle the whole thing works now, just need to add other data types and perhaps do some cleanup.

* Add other layouts for batched gemm wmma chufflev3 f16 f16 f16. Now matching XDL (for f16).

* Add bf16 bf16 bf16 support for batched gemm wmma cshuffle v3 for all layouts.

* Fixup comments and TODOs

* Expand test cases for batched gemm wmma cshuffle v3 with more unusual shapes. Some of the original test cases for batched gemm do not work based on cshuffle v3 because the dimensions are too small.

* Fix argument order for calls to profile_batched_gemm_impl() ONLY in wmma tests.

* Take batching into account when using rotating memory or clearing the C tensor.

* Implement small refactors / comments etc. from review.

* Port recent gemm wmma updates to batched gemm wmma: V1 pipeline, non-main-k-block-loop, check compute type, packed buffer size calc. Ported new instance lists.

* Add MNKPadding instances to batched gemm wmma cshuffle v3, remove incompatible test problems.

* Put clearing the C matrix in a pre-process lambda for the non-flush case + small fixups.

* Once again switch order of strides and batch strides in calls to profile_batched_gemm_impl() from test_batched_gemm_wmma to match latest definition of that function.

---------

Co-authored-by: kiefer <kiefer.van.teutem@streamhpc.com>
2025-06-24 07:28:13 -07:00