assistant-librarian[bot] bb15392230 [CK] Add fwd conv group merging to v3 conv instances (#4273) ...

## Proposed changes

Added conv group merging to the (universal) V3 fwd conv pipeline. The
new instance improves fwd conv performance when the number of
input/output channel per group is low.

On MI300 (`gfx942`) we get

| CK prof command | Baseline (TFLOPS) | V3 group merging (TFLOPS) |
|:-----|:------:|------:|
| grouped_conv_fwd 1 1 1 0 1 0 1 2 32 32 4 4 3 3 200 200 1 1 1 1 1 1 1 1
| 3.86035 | 8.36796 |
| grouped_conv_fwd 1 1 1 0 1 0 1 2 32 32 8 8 3 3 200 200 2 2 1 1 1 1 1 1
| 10.1867 | 13.4677 |
| grouped_conv_fwd 1 1 1 0 1 0 1 2 32 32 8 8 3 3 100 100 1 2 1 1 1 1 1 1
| 11.7875 | 16.3657 |



---
🔁 Imported from
[ROCm/composable_kernel#3675](https://github.com/ROCm/composable_kernel/pull/3675)
🧑‍💻 Originally authored by @vpietila-amd

---------

Co-authored-by: Ville Pietilä <>
Co-authored-by: Ville Pietilä <188998872+vpietila-amd@users.noreply.github.com>
Co-authored-by: systems-assistant[bot] <systems-assistant[bot]@users.noreply.github.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: Bartlomiej Kocot <barkocot@amd.com>

2026-02-08 12:34:59 +01:00

..

conv

[CK] Add fwd conv group merging to v3 conv instances (#4273)

2026-02-08 12:34:59 +01:00

impl

[CK] Add fwd conv group merging to v3 conv instances (#4273)

2026-02-08 12:34:59 +01:00

utils

[CK_BUILDER] Add grouped conv fwd ck tile profiler (#3518)

2026-01-19 22:29:01 -07:00

validation

[CK_BUILDER] Replace reference conv with old ck implementation (#3604)

2026-01-21 19:18:47 +01:00

CMakeLists.txt

[CK_BUILDER] conv bwd weight testing (#3618)

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

test_bwd_data_instance_traits.cpp

[CK-Tile] move out memory operation from cshuffle epilogue class (#3359)

2026-01-04 03:28:14 -08:00

test_bwd_weight_instance_traits.cpp

[CK-Tile] move out memory operation from cshuffle epilogue class (#3359)

2026-01-04 03:28:14 -08:00

test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ck_factory_grouped_convolution_forward_bias_clamp.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ck_factory_grouped_convolution_forward_bilinear.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ck_factory_grouped_convolution_forward_clamp.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ck_factory_grouped_convolution_forward_convscale.cpp

[CK_BUILDER] Improve CK Builder and CK Builder tests (#3382)

2025-12-11 09:50:00 +02:00

test_ck_factory_grouped_convolution_forward_dynamic_op.cpp

[CK_BUILDER] Improve CK Builder and CK Builder tests (#3382)

2025-12-11 09:50:00 +02:00

test_ck_factory_grouped_convolution_forward_scale.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp

[CK_BUILDER] Improve CK Builder and CK Builder tests (#3382)

2025-12-11 09:50:00 +02:00

test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp

[CK_BUILDER] Improve CK Builder and CK Builder tests (#3382)

2025-12-11 09:50:00 +02:00

test_ck_factory_grouped_convolution_forward.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_ckb_conv_builder.cpp

chore: update copyright header for misc files (#3402)

2025-12-11 08:25:29 -08:00

test_conv_builder.cpp

Fix copyright messages in experimental/builder. (#3253)

2025-11-20 17:40:55 -08:00

test_conv_description.cpp

[CK] Add fwd conv group merging to v3 conv instances (#4273)

2026-02-08 12:34:59 +01:00

test_fwd_instance_traits.cpp

[CK] Add fwd conv group merging to v3 conv instances (#4273)

2026-02-08 12:34:59 +01:00

test_inline_diff.cpp

[CK_BUILDER] validation (#3471)

2026-01-05 04:57:34 -08:00

test_instance_string_bwd_weight_grp_conv_dl.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_multiple_d_wmma_v3.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_multiple_d_xdl.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_two_stage_wmma_v3.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_two_stage_xdl.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_wmma_v3.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_wmma.cpp

[CK_BUILDER] Instance traits for conv bwd weight algorithms (#3498)

2025-12-31 15:41:15 -08:00

test_instance_string_bwd_weight_grp_conv_xdl_v3.cpp

Grouped Conv Bwd Weight Direct Load (#3648)

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

test_instance_string_bwd_weight_grp_conv_xdl.cpp

Add describe() method to device ops for runtime introspection (#3375)

2025-12-14 12:49:12 -08:00

test_instance_string_fwd_grp_conv_dl.cpp

Add describe() method to device ops for runtime introspection (#3375)

2025-12-14 12:49:12 -08:00

test_instance_string_fwd_grp_conv_large_tensor.cpp

Add describe() method to device ops for runtime introspection (#3375)

2025-12-14 12:49:12 -08:00

test_instance_string_fwd_grp_conv_v3.cpp

[CK] Add fwd conv group merging to v3 conv instances (#4273)

2026-02-08 12:34:59 +01:00

test_instance_string_fwd_grp_conv_wmma.cpp

Add describe() method to device ops for runtime introspection (#3375)

2025-12-14 12:49:12 -08:00

test_instance_string_fwd_grp_conv.cpp

Add describe() method to device ops for runtime introspection (#3375)

2025-12-14 12:49:12 -08:00

test_instance_traits_util.cpp

Simplify includes for CK builder reflection (#3357)

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

test_testing_utils.cpp

[CK_BUILDER] conv bwd weight testing (#3618)

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

testing_utils.cpp

[CK_BUILDER] conv bwd weight testing (#3618)

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

testing_utils.hpp

[CK_BUILDER] Integrate CKB validation with CK verification (#3649)

2026-01-28 17:41:02 +01:00

unit_conv_elementwise_op.cpp

[CK_BUILDER] Ck Tile Grouped convolution factory (#3352)

2025-12-08 10:32:56 +01:00

unit_conv_fwd_testing.cpp

[CK_BUILDER] Integrate CKB validation with CK verification (#3649)

2026-01-28 17:41:02 +01:00

unit_conv_tensor_layout.cpp

[CK_BUILDER] Replace reference conv with old ck implementation (#3604)

2026-01-21 19:18:47 +01:00

unit_conv_tensor_type.cpp

[CK_BUILDER] Debug utilities (#3528)

2026-01-08 10:14:13 +01:00

unit_conv_thread_block.cpp

[CK_BUILDER] Ck Tile Grouped convolution factory (#3352)

2025-12-08 10:32:56 +01:00

unit_conv_tuning_params.cpp

[CK_BUILDER] Add bwd weight factories (#3509)

2026-01-13 18:12:38 +02:00

unit_debug.cpp

[CK_BUILDER] Debug utilities (#3528)

2026-01-08 10:14:13 +01:00

unit_device_buffer.cpp

ck-builder: tensor input/output reflection (#3536)

2026-01-12 09:45:53 +01:00

unit_error.cpp

Fix build error from extra comma (#3516)

2026-01-06 11:08:54 -08:00

unit_tensor_descriptor.cpp

[CK_BUILDER] Add grouped conv fwd ck tile profiler (#3518)

2026-01-19 22:29:01 -07:00

unit_tensor_foreach.cpp

[CK_BUILDER] Integrate CKB validation with CK verification (#3649)

2026-01-28 17:41:02 +01:00

unit_validation.cpp

[CK_BUILDER] Integrate CKB validation with CK verification (#3649)

2026-01-28 17:41:02 +01:00