Zoltán Lakatos 1a8bd3d34b 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

..

batchnorm_multiblock

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gemm_layernorm

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

normalization

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

block_to_ctile_map.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

epilogue_cshuffle_v3_reduce_wmma.hpp

Wmma support for gemm_bias_add_reduce (#3316)

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

epilogue_cshuffle_v3_welford_wmma.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

epilogue_cshuffle_v3_wmma_base.hpp

Add support for direct store in epilogue and padding support for wave transfer without transpose (#3465)

2026-01-14 11:02:19 +01:00

epilogue_cshuffle_v3_wmma.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

epilogue_direct_store.hpp

Add support for direct store in epilogue and padding support for wave transfer without transpose (#3465)

2026-01-14 11:02:19 +01:00

gridwise_2d_multiple_reduction_multiblock.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_2d_multiple_reduction_threadwise.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_2d_reduction_multiblock.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_2d_reduction_threadwise_multi_d.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_2d_reduction_threadwise.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_ab_transfer_thread_tiles_preshuffle.hpp

Multi AB support for wave transfer (#3578)

2026-01-29 10:29:40 -08:00

gridwise_ab_transfer_thread_tiles.hpp

Multi AB support for wave transfer (#3578)

2026-01-29 10:29:40 -08:00

gridwise_ab_transfer_wave_tiles_interleave.hpp

Multi AB support for wave transfer (#3578)

2026-01-29 10:29:40 -08:00

gridwise_ab_transfer_wave_tiles.hpp

Multi AB support for wave transfer (#3578)

2026-01-29 10:29:40 -08:00

gridwise_batched_gemm_gemm_wmma_cshuffle_v3.hpp

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

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

gridwise_batched_gemm_gemm_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_batched_gemm_multiple_d_softmax_gemm_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_batched_gemm_softmax_gemm_wmma_cshuffle.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_batched_gemm_softmax_gemm_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_batchnorm_backward_blockwise_welford.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_batchnorm_forward_blockwise_welford.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_elementwise_1d_scale.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_elementwise_2d.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_elementwise_layernorm_welford_variance.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_fpAintB_gemm_wmma.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_bias_add_reduce_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_dl_multiple_d.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_dl_v1r3.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_dpp.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_multiple_abd_xdl_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_multiple_d_wmma_cshuffle.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_multiple_d_xdl_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_multiple_d_xdl_splitk_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_pipeline_selector.hpp

Simplify includes for CK builder reflection (#3357)

2025-12-05 07:44:10 -08:00

gridwise_gemm_pipeline_v1.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_pipeline_v2.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_pipeline_v3.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_pipeline_v4_direct_load.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_reduce_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle_v2.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_waveletmodel.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_wmma_cshuffle_v3_ab_scale.hpp

Wmma support for grouped convolution bwd weight (#2947)

2025-12-17 15:58:58 -08:00

gridwise_gemm_wmma_cshuffle_v3_common.hpp

Multi AB support for wave transfer (#3578)

2026-01-29 10:29:40 -08:00

gridwise_gemm_wmma_cshuffle_v3.hpp

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

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

gridwise_gemm_wmma.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_gemm_xdl_cshuffle_common.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_conv_v3.hpp

Grouped Conv Bwd Weight Direct Load (#3648)

2026-01-28 15:31:54 -06:00

gridwise_gemm_xdl_cshuffle_streamk_v3.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v2.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_b_preshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_b_scale.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_b_preshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_blockscale_b_preshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_mx_bpreshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3_mx.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_cshuffle_v3.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_layernorm_cshuffle_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdl_waveletmodel_cshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_bwd_weight.hpp

[Compiler] Addressing new compiler warnings (#3640)

2026-02-02 09:39:48 -08:00

gridwise_gemm_xdlops_skip_b_lds_v1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_splitk_lds_direct_load.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_streamk.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v2r3.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v2r4.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v2r4r2.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v3r1.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v3r2.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_gemm_xdlops_v3r3.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_moe_gemm_blockscale.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_moe_gemm.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_moe_mx_gemm_bns.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_moe_mx_gemm_bpreshuffle.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_moe_mx_gemm.hpp

[ck] add gridwise base class for in all xdl kernel (#186) (#3544)

2026-01-27 12:49:47 -08:00

gridwise_permute.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_put_element_1d.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_set_buffer_value.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_set_multiple_buffer_value.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_softmax.hpp

fix undefined behaviour in softmax kernel (#3683)

2026-01-30 15:22:54 +08:00

gridwise_sparse_embeddings_forward_layernorm_builtins.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_sparse_embeddings_forward_layernorm.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00

gridwise_tensor_rearrange.hpp

chore(copyright): update copyright header for include directory (#3293)

2025-11-26 11:00:05 -07:00