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

..

amd_address_space.hpp

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

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

amd_arch.hpp

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

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

amd_buffer_addressing_builtins.hpp

Update AMD buffer coherency (#3403)

2025-12-18 10:16:22 +01:00

amd_buffer_addressing.hpp

Update AMD buffer coherency (#3403)

2025-12-18 10:16:22 +01:00

amd_buffer_coherence.hpp

Update AMD buffer coherency (#3403)

2025-12-18 10:16:22 +01:00

amd_ck_fp8.hpp

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

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

amd_gemm_dpp.hpp

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

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

amd_inline_asm.hpp

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

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

amd_lds.hpp

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

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

amd_smfmac.hpp

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

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

amd_transpose_load.hpp

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

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

amd_wave_read_first_lane.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

amd_wmma.hpp

Add support to gfx1153 and fix gfx115X WMMA config (#3496)

2026-01-05 10:03:30 -08:00

amd_xdlops.hpp

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

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

array_multi_index.hpp

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

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

array.hpp

WMMA grouped conv fwd large tensor extra flavors (#3582)

2026-01-23 12:19:51 +01:00

blkgemmpipe_scheduler.hpp

Simplify includes for CK builder reflection (#3357)

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

c_style_pointer_cast.hpp

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

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

common_header.hpp

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

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

container_element_picker.hpp

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

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

container_helper.hpp

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

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

data_type.hpp

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

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

debug.hpp

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

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

dtype_fp64.hpp

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

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

dtype_vector.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

dynamic_buffer.hpp

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

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

e8m0.hpp

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

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

enable_if.hpp

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

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

env.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

f8_utils.hpp

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

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

filter_tuple.hpp

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

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

flush_icache.hpp

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

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

functional2.hpp

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

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

functional3.hpp

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

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

functional4.hpp

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

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

functional.hpp

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

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

generic_memory_space_atomic.hpp

[CK] Allow tensors larger than 2GB in grouped conv bwd weight (#3169)

2026-01-08 08:02:02 +01:00

get_id.hpp

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

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

get_shift.hpp

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

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

ignore.hpp

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

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

inner_product_dpp8.hpp

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

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

inner_product.hpp

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

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

integral_constant.hpp

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

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

is_detected.hpp

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

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

is_known_at_compile_time.hpp

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

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

loop_scheduler.hpp

Simplify includes for CK builder reflection (#3357)

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

magic_division.hpp

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

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

math_v2.hpp

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

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

math.hpp

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

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

multi_index.hpp

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

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

mxf4_utils.hpp

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

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

mxf6_utils.hpp

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

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

mxf8_utils.hpp

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

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

mxfp_utils.hpp

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

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

number.hpp

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

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

numeric_limits.hpp

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

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

numeric_utils.hpp

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

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

pipeline_enum.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

random_gen.hpp

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

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

reduction_common.hpp

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

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

reduction_enums.hpp

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

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

reduction_functions_accumulate.hpp

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

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

reduction_operator.hpp

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

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

scaled_type_convert.hpp

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

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

scheduler_enum.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

sequence_helper.hpp

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

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

sequence.hpp

Optimize sequence metaprogramming utilities to reduce template instantiation depth (#3585)

2026-01-26 10:08:55 -08:00

span.hpp

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

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

static_buffer.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

statically_indexed_array_multi_index.hpp

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

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

statically_indexed_array.hpp

Optimize sequence metaprogramming utilities to reduce template instantiation depth (#3585)

2026-01-26 10:08:55 -08:00

synchronization.hpp

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

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

thread_group.hpp

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

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

transpose_vectors.hpp

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

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

tuple_helper.hpp

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

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

tuple.hpp

[Compiler] Addressing new compiler warnings (#3640)

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

type_convert.hpp

fix some syntax errors (#3658)

2026-01-27 09:59:39 -08:00

type.hpp

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

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

workgroup_barrier.hpp

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

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

workgroup_synchronization.hpp

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

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