Enrico Degregori 12225ce645 Wmma support for multiple ABD GEMM (#2803) ...

* 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)

* Fix bug in device print function

* Fix unused template parameter

* Fix batched gemm for multiABD gridwise implementation

* Fix gemm_universal_reduce with multiABDs gridwise implementation

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 3d29bff2f0]

2025-09-22 18:49:06 -07:00

..

avg_pool2d_bwd

Pool2d max/avg kernel in the BWD version (#1494)

2024-09-12 11:47:52 +02:00

avg_pool3d_bwd

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

batched_gemm

upgrade from clang-format-12 to clang-format-18 (#2568)

2025-07-28 11:34:07 -07:00

batched_gemm_add_relu_gemm_add

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

batched_gemm_b_scale

Added Int4 mixed batch gemm support (#1839)

2025-02-10 11:17:02 +08:00

batched_gemm_bias_permute

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

batched_gemm_gemm

Implement batched gemm gemm for RDNA (3 and 4) (#2612)

2025-09-04 14:10:24 -07:00

batched_gemm_multi_d

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

batched_gemm_reduce

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

batched_gemm_softmax_gemm

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

batched_gemm_softmax_gemm_permute

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

batchnorm

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

column_to_image

Add support for groups in Img2Col/Col2Img (#1007)

2023-10-31 10:46:32 +01:00

contraction_bilinear

Extend support for contraction 6D (#1207)

2024-04-09 23:46:21 +02:00

contraction_scale

Extend support for contraction 6D (#1207)

2024-04-09 23:46:21 +02:00

conv1d_bwd_data

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

conv2d_bwd_data

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

conv2d_fwd

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

conv2d_fwd_bias_relu

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

conv2d_fwd_bias_relu_add

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

conv3d_bwd_data

Mark non-grouped convolutions instances as deprecated (#2595)

2025-08-04 16:49:55 +02:00

elementwise

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

elementwise_normalization

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

gemm

upgrade from clang-format-12 to clang-format-18 (#2568)

2025-07-28 11:34:07 -07:00

gemm_ab_scale

Add other layouts for FP8 block scaled gemm (#2665)

2025-08-18 01:46:10 -07:00

gemm_add

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_add_add_fastgelu

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_add_fastgelu

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_add_multiply

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_add_relu

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_add_relu_add_layernorm

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gemm_add_silu

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_b_scale

Support b_scale: (#2350)

2025-07-24 18:49:58 -07:00

gemm_bias_add_reduce

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

gemm_bilinear

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_blockscale_wp

delete all slp compilation flag in CK Tile (#2625)

2025-08-06 00:34:39 -07:00

gemm_fastgelu

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_multi_abd

Wmma support for multiple ABD GEMM (#2803)

2025-09-22 18:49:06 -07:00

gemm_multiply_add

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_multiply_multiply

Wmma support for multiple Ds based GEMMs (#2613)

2025-09-05 16:31:08 +02:00

gemm_multiply_multiply_wp

Convert CK (GeMM MulMul Weight Preshuffle) instances to use 16x16 xdl tile (#2229)

2025-06-10 09:37:14 -07:00

gemm_mx

[CK] Fixed MPerBlock=32 build issue for MXFP4 GEMM decode (#2512)

2025-07-18 14:35:54 +08:00

gemm_reduce

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

gemm_splitk

Optimized GEMMs for MX FP4/8 (#2294)

2025-06-05 13:54:15 -06:00

gemm_streamk

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

gemm_universal

Add gemm universal f8 f8 bf16 instances on gfx950 (#2662)

2025-08-14 13:25:24 -07:00

gemm_universal_batched

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

gemm_universal_preshuffle

[ckProfiler] Add infrastructure and instances to profile gemm_universal with B preshuffle (#2427)

2025-07-01 18:34:52 -07:00

gemm_universal_reduce

WMMA support for GEMM reduce (#2823)

2025-09-12 21:36:43 +02:00

gemm_universal_streamk

upgrade from clang-format-12 to clang-format-18 (#2568)

2025-07-28 11:34:07 -07:00

grouped_conv1d_bwd_weight

Enable grouped conv bwd wei bf16 NGCHW (#1589)

2024-10-22 16:18:28 +02:00

grouped_conv1d_fwd

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_conv2d_bwd_data

[CK][CONV] Support NCHW in class DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 (#2459)

2025-07-17 08:19:57 +08:00

grouped_conv2d_bwd_weight

Grouped conv bwd wei add for larger filter and Merge Groupes optimization (#2197)

2025-05-21 22:47:34 +02:00

grouped_conv2d_fwd

[CK][CONV] Support NCHW in class DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle (#2375)

2025-06-26 08:32:39 +08:00

grouped_conv2d_fwd_bias_bnorm_clamp

Grouped Convolution Forward Infer Bias Bnorm Activ (#2621)

2025-08-07 08:36:47 +02:00

grouped_conv2d_fwd_bias_clamp

Grouped conv bias clamp fp32/fp16 support (#2366)

2025-06-20 11:41:04 +02:00

grouped_conv2d_fwd_clamp

Grouped conv bias clamp fp32/fp16 support (#2366)

2025-06-20 11:41:04 +02:00

grouped_conv2d_fwd_dynamic_op

Temporary disable part of dynamic op conv instances (#1630)

2024-11-04 13:34:17 -08:00

grouped_conv3d_bwd_data

[CK][CONV] Support NCHW in class DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1 (#2459)

2025-07-17 08:19:57 +08:00

grouped_conv3d_bwd_data_bilinear

Integrate universal gemm with conv bwd data and add SplitK (#1315)

2025-04-28 23:54:49 +02:00

grouped_conv3d_bwd_data_scale

Integrate universal gemm with conv bwd data and add SplitK (#1315)

2025-04-28 23:54:49 +02:00

grouped_conv3d_bwd_weight

Grouped conv bwd wei add for larger filter and Merge Groupes optimization (#2197)

2025-05-21 22:47:34 +02:00

grouped_conv3d_bwd_weight_bilinear

Re-enable fp8 types for all architectures. (#1470)

2024-08-16 16:07:52 -06:00

grouped_conv3d_bwd_weight_scale

Re-enable fp8 types for all architectures. (#1470)

2024-08-16 16:07:52 -06:00

grouped_conv3d_fwd

TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)

2025-09-17 14:50:15 -07:00

grouped_conv3d_fwd_bias_bnorm_clamp

TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)

2025-09-17 14:50:15 -07:00

grouped_conv3d_fwd_bias_clamp

TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)

2025-09-17 14:50:15 -07:00

grouped_conv3d_fwd_bilinear

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_conv3d_fwd_clamp

TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)

2025-09-17 14:50:15 -07:00

grouped_conv3d_fwd_convinvscale

Add a convinvscale op, related instances and examples (#1307)

2024-06-10 14:48:49 -05:00

grouped_conv3d_fwd_convscale

Adding Instances and Examples for FP8-based Scaled Convolution and AMAX Reduction. (#1473)

2024-08-21 15:22:41 -07:00

grouped_conv3d_fwd_convscale_add

Adding more instances of grouped convolution 3d forward for FP8 with ConvScale+Bias element-wise operation. (#1412)

2024-07-24 15:49:55 -05:00

grouped_conv3d_fwd_convscale_relu

Adding Instances and Examples for FP8-based Scaled Convolution and AMAX Reduction. (#1473)

2024-08-21 15:22:41 -07:00

grouped_conv3d_fwd_dynamic_op

Temporary disable part of dynamic op conv instances (#1630)

2024-11-04 13:34:17 -08:00

grouped_conv3d_fwd_scale

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_conv3d_fwd_scaleadd_ab

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_conv3d_fwd_scaleadd_scaleadd_relu

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_convnd_bwd_weight

Grouped conv bwd wei explicit GEMM for odd C/K (#2306)

2025-06-10 11:17:12 +02:00

grouped_gemm

Remove CK_USE_AMD_MFMA_GFX950 (#1935)

2025-03-04 10:32:25 -08:00

grouped_gemm_bias

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_gemm_fastgelu

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

grouped_gemm_fixed_nk

Added bf16 instances grouped gemm fixed nk (#1825)

2025-01-20 09:13:09 -08:00

grouped_gemm_fixed_nk_multi_abd

bf16A_Int8B with fastgelu/bias (#1264)

2024-04-26 07:26:30 -05:00

grouped_gemm_tile_loop

upgrade from clang-format-12 to clang-format-18 (#2568)

2025-07-28 11:34:07 -07:00

image_to_column

Add support for groups in Img2Col/Col2Img (#1007)

2023-10-31 10:46:32 +01:00

max_pool_bwd

This commit contains implementation of max pool2d for f8 type (#1506)

2024-09-16 10:15:06 +02:00

mha

Label CMakeLists message() as DEBUG or STATUS for clean build output (#2301)

2025-06-10 10:46:47 -07:00

normalization_bwd_data

layernorm and groupnorm backward data (#1083)

2023-12-19 04:23:11 +08:00

normalization_bwd_gamma_beta

layernorm & groupnorm bwd gamma beta (#1133)

2024-01-25 19:53:15 +08:00

normalization_fwd

layernorm and groupnorm backward data (#1083)

2023-12-19 04:23:11 +08:00

permute_scale

Adding Instances and Examples for FP8-based Scaled Convolution with ReLU Activation and AMAX Reduction. (#1469)

2024-08-20 10:30:56 -05:00

pool2d_fwd

Add pool2d int8 and fp8 instances (#1508)

2024-09-13 10:18:21 -07:00

pool3d_fwd

OCP FP8 support for gfx12. (#1710)

2024-12-03 08:42:55 -08:00

quantization

Added wmma support for gemm quantization: (#2841)

2025-09-16 16:23:29 -07:00

reduce

Adding Instances and Examples for FP8-based Scaled Convolution with ReLU Activation and AMAX Reduction. (#1469)

2024-08-20 10:30:56 -05:00

softmax

enable compilation of INSTANCES_ONLY for Windows (#1082)

2023-12-20 14:34:53 -08:00

transpose

Hip tensor permute unit test (#1068)

2023-12-18 21:35:00 -06:00

CMakeLists.txt

TF32 POC in Conv3d on MI30x platform #2763 (second attempt) (#2852)

2025-09-17 14:50:15 -07:00