Commit Graph

428 Commits

Author SHA1 Message Date
kiefer
78635fd74e Add 1D instances 2025-08-28 15:57:30 +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
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
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
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
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
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
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
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
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
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
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
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
John Shumway
47ae4b0955 Shard several of the most costly targets. (#2373)
* Shard several of the most costly targets.

Introduces a filter_tuple_by_modulo to break up tuples.

Drops build time of target from 21 minutes to under 14 minutes with 64
build processes, or 11 minutes with 128 build processes.

time ninja -j 64 device_grouped_conv3d_fwd_instance

* fix clang format

* Fix build errors in instantiation code.

I wasn't sure how to test the header-only instantiation code on my
initial commit. From Jenkins CI test results, I see that there is a
test target that depends on these headers:

ninja -j 128 test_grouped_convnd_fwd

This allowed me to test the build locally. I found three mistakes I
made, mostly related to early experiments on I tried on the code.
This was hard to find earlier because this PR is really too large.

I also discovered that there are five 2D convolution targets that now
dominate the compilation time. I will likely address those in a later
PR, rather than adding even more changes to this PR.

* Fix link errors from mismatched declarations.

Our pattern for instantiating MIOpen templates uses duplicate
declarations (instead of headers). This is fragile, and I didn't
notice that my last commit had a bunch of link errors. I fixed these
mistakes, and the bin/test_grouped_conv_fwd test target binary now links
correctly.

* Migrate the design to a code-generation approach.

Use a CMake function with template files to generate the source files for the
intantiating the kerenels and to generate the calling function.

* Shard the longest 2D convolution builds

Now that we have automated the shard instantiation, we can shard the 2D
convolution targets that take the longest to build. The target
test_grouped_conv2d_fwd now compiles in 15 minutes.

* Use PROJECT_SOURCE_DIR for submodule compatibility

I used CMAKE_SOURCE_DIR to refer to the top-level source directory in
the ShardInstantiation.cmake file, but this can cause issues with
git submodules.  Instead, we should use PROJECT_SOURCE_DIR to ensure
compatibility when this project is used as a submodule in another
project.

* Migrate the design to a code-generation approach.

Use a CMake function with template files to generate the source files for the
intantiating the kerenels and to generate the calling function.

* Migrate the design to a code-generation approach.

Use a CMake function with template files to generate the source files for the
intantiating the kerenels and to generate the calling function.

* Remove accidental copy of a file

* Remove accidental copies of template files.

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
2025-06-23 07:24:36 -07:00
Bartłomiej Kocot
663992e99b Grouped conv bias clamp fp32/fp16 support (#2366) 2025-06-20 11:41:04 +02:00
Zoltan Lakatos
5e454276e3 fp8 instances - not tested 2025-06-19 10:57:38 +00:00
apoorva
38d00277c2 Updated tests. 2025-06-19 10:56:18 +00:00
apoorva
ef781db305 Fixing build errors 2025-06-19 09:47:36 +00:00
apoorva
0cce81c3e7 Code update of template parameters modified. 2025-06-19 09:47:07 +00:00
Apoorva Kalyani
1fda4990ad added datatype and use clang-format-12
(cherry picked from commit ae4e853682ef1bb27784b2f965b4a66b3751ceec)

Co-authored-by: Cenxuan <cenxuan@streamhpc.com>
2025-06-19 09:47:06 +00:00
Apoorva Kalyani
455275de80 temp work saved, changed the BDataType to f16 or bf16 since wmma currently not support non-equal A and B datatype
(cherry picked from commit 22fbd68f1db458ab50780a394ee2544c7a1484d1)

Co-authored-by: Cenxuan <cenxuan@streamhpc.com>
2025-06-19 09:47:06 +00:00
Apoorva Kalyani
b42b6b67a5 adding gemm_add wmma_cshuffle and other support
(cherry picked from commit ec447e7f564095ea969eddc39ec77b843aa52976)

Co-authored-by: Cenxuan <cenxuan@streamhpc.com>
2025-06-19 09:45:34 +00:00
Zoltan Lakatos
40ce862cdf Merge remote-tracking branch 'origin/feature/multiple-d-gemms' into 64-implement-device_gemm_multiply_multiply_instance-for-rdna4 2025-06-18 14:46:38 +00:00
Illia Silin
cdfd7722bf Revert "Shard several of the most costly targets. (#2266)" (#2361)
This reverts commit 3a0cb27966.
2025-06-17 13:56:30 -07:00
Zoltan Lakatos
ac60286ed0 added wmma multiply_multiply instances 2025-06-17 19:46:22 +00:00
Bartłomiej Kocot
f6c2ff9dce Grouped convolution forward with clamp (#2334)
* Grouped convolution forward with clamp

* Optimize clamp

* unary fixes

* test gk bias

* Revert "test gk bias"

This reverts commit 8e42e29d7b.

* Revert "Revert "test gk bias""

This reverts commit e73c0550ce.

* workaround comment
2025-06-16 15:36:53 +02:00