Zoltán Lakatos 839a37780c Implement device grouped gemm fixed nk multi abd for rdna4 (#3619) ...

* device struct implementation

* added xdl grouped multi abd fixed nk testing

* wmma implementation fixed

* avoid unnecessary device mem allocation and code cleanups

* cleanup instances definitions

* wmma examples added

* code cleanups

* fix clang format

* typo and compilation fixes related to reference gemm

* fix compilation error due to std::remove_cvref_t

* added missing hip_check_error includes

* correction to example instances

* review commentes addressed

* removed split-k from testing

* code formatting

---------

Co-authored-by: Zoltán Lakatos <zoltan.lakatos@streamhpc.com>
Co-authored-by: illsilin_amdeng <Illia.Silin@amd.com>

[ROCm/composable_kernel commit: 301eb5cf08]

2026-02-02 13:58:11 -08:00

..

common.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

data_type_enum.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

grouped_convolution_forward_tile_algs.hpp

[CK TILE] Enable CK TILE Conv Fwd tests in CI and fix check_err (#3624)

2026-01-27 11:04:11 +02:00

grouped_convolution_signatures.hpp

[CK_BUILDER] conv bwd weight testing (#3618)

2026-01-26 23:50:15 +01:00

profile_avg_pool2d_bwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_avg_pool3d_bwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_contraction_multiple_d_impl.hpp

Implement batched gemm bias permute for RDNA4 (#3534)

2026-01-17 08:30:27 +01:00

profile_batched_gemm_add_relu_gemm_add_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_b_scale_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_bias_softmax_gemm_permute_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_gemm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_multiple_d_gemm_multiple_d_impl.hpp

Implement batched gemm add relu gemm add for rdna4 (#3391)

2026-01-20 13:06:59 -08:00

profile_batched_gemm_reduce_impl.hpp

WMMA support for batched_gemm_reduce (#3332)

2026-01-20 10:50:46 +01:00

profile_batched_gemm_softmax_gemm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batched_gemm_softmax_gemm_permute_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batchnorm_backward_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batchnorm_forward_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_batchnorm_infer_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_contraction_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_contraction_utils.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_conv_bwd_data_impl.hpp

[CK tests] Extend conv GPU reference (#3539)

2026-01-27 09:49:42 +01:00

profile_conv_fwd_bias_relu_add_impl.hpp

Profile resnet layout fixes (#3360)

2025-12-05 08:20:46 -08:00

profile_conv_fwd_bias_relu_impl.hpp

Profile resnet layout fixes (#3360)

2025-12-05 08:20:46 -08:00

profile_conv_fwd_impl.hpp

[CK tests] Extend conv GPU reference (#3539)

2026-01-27 09:49:42 +01:00

profile_conv_tensor_rearrange_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_elementwise_layernorm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_ab_scale_impl.hpp

Wmma support for gemm_ab_scale (#3314)

2025-12-11 09:06:20 +01:00

profile_gemm_add_add_fastgelu_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_fastgelu_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_multiply_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_relu_add_layernorm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_relu_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_add_silu_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_b_scale_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_bias_add_reduce_impl.hpp

Wmma support for gemm_bias_add_reduce (#3316)

2026-01-07 10:27:16 -08:00

profile_gemm_bilinear_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_blockscale_wp_impl.hpp

Wmma support for gemm_ab_scale (#3314)

2025-12-11 09:06:20 +01:00

profile_gemm_fastgelu_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_multi_abd_impl.hpp

Implement device grouped gemm fixed nk multi abd for rdna4 (#3619)

2026-02-02 13:58:11 -08:00

profile_gemm_multiply_add_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_multiply_multiply_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_multiply_multiply_wp_impl.hpp

Wmma support for gemm_multiply_multiply_wp (#3278)

2025-12-03 07:38:23 -08:00

profile_gemm_mx_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_quantization_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_reduce_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_splitk_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_streamk_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_universal_batched_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_universal_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_universal_preshuffle_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_universal_reduce_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_gemm_universal_streamk_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_grouped_conv_bwd_data_impl.hpp

[CK Profiler] Restore CPU tensor initialization when verification is not done on GPU (#3594)

2026-01-16 10:56:53 -08:00

profile_grouped_conv_bwd_weight_impl.hpp

[Conv] Enable bwd weight splitk autodeduction with cap (#3656)

2026-01-29 17:40:28 +00:00

profile_grouped_conv_fwd_bias_bnorm_clamp_impl.hpp

Grouped conv_fwd_bias_bnorm_clamp instances and tests (#3525)

2026-01-22 09:53:59 +01:00

profile_grouped_conv_fwd_bias_clamp_impl.hpp

[CK tests] Extend conv GPU reference (#3539)

2026-01-27 09:49:42 +01:00

profile_grouped_conv_fwd_bilinear_impl.hpp

[CK tests] Extend conv GPU reference (#3539)

2026-01-27 09:49:42 +01:00

profile_grouped_conv_fwd_convscale_add_impl.hpp

Adding remaining conv, dynamic_op, and scaleadd_scaleadd_relu flavors for grouped conv fwd (#3529)

2026-01-30 17:02:14 +01:00

profile_grouped_conv_fwd_impl.hpp

Adding remaining conv, dynamic_op, and scaleadd_scaleadd_relu flavors for grouped conv fwd (#3529)

2026-01-30 17:02:14 +01:00

profile_grouped_conv_fwd_outelementop_impl.hpp

Adding remaining conv, dynamic_op, and scaleadd_scaleadd_relu flavors for grouped conv fwd (#3529)

2026-01-30 17:02:14 +01:00

profile_grouped_conv_fwd_scaleadd_scaleadd_relu_impl.hpp

Adding remaining conv, dynamic_op, and scaleadd_scaleadd_relu flavors for grouped conv fwd (#3529)

2026-01-30 17:02:14 +01:00

profile_grouped_gemm_fastgelu_impl.hpp

Implement grouped gemm fastgelu for RDNA4 (#3303)

2026-01-07 10:20:44 -08:00

profile_grouped_gemm_fixed_nk_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_grouped_gemm_impl.hpp

Implement grouped gemm fastgelu for RDNA4 (#3303)

2026-01-07 10:20:44 -08:00

profile_grouped_gemm_multi_abd_fixed_nk_impl.hpp

Implement device grouped gemm fixed nk multi abd for rdna4 (#3619)

2026-02-02 13:58:11 -08:00

profile_grouped_gemm_multiply_tile_loop_impl.hpp

Implement grouped gemm tile loop for RDNA4 (#3304)

2026-01-13 07:14:23 +01:00

profile_grouped_gemm_tile_loop_generic_impl.hpp

Implement grouped gemm tile loop for RDNA4 (#3304)

2026-01-13 07:14:23 +01:00

profile_grouped_gemm_tile_loop_impl.hpp

Implement grouped gemm tile loop for RDNA4 (#3304)

2026-01-13 07:14:23 +01:00

profile_groupnorm_bwd_data_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_groupnorm_bwd_gamma_beta_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_groupnorm_fwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_layernorm_bwd_data_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_layernorm_bwd_gamma_beta_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_layernorm_fwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_max_pool2d_bwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_max_pool3d_bwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_permute_scale_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_pool2d_fwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_pool3d_fwd_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_reduce_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_softmax_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00

profile_transpose_impl.hpp

chore(copyright): update copyright header for profiler directory (#3205)

2025-11-14 11:19:25 -08:00