John Shumway 1f97dc1aee Add describe() method to device ops for runtime introspection (#3375) ...

Introduces a polymorphic describe() method to BaseOperator that enables runtime introspection of kernel configurations through a unified interface.

Key changes:

* Add virtual describe() method to BaseOperator returning Description objects
* Implement describe() in 6 device operation classes (conv fwd/bwd variants)
* Create conv_describe.hpp with factory function for ConvDescription
* Extract type definitions to conv_types.hpp to resolve circular dependencies
* Add InstanceStringDescription for kernels without full ConvDescription support

Other Improvements:

* Update tests to use describe() instead of GetInstanceString()
* Remove circular dependency include from conv_traits.hpp
* Add ODD_C to ConvFwdSpecialization enum and fix OddC mapping
* Replace silent fallback in conv_layout() with compile-time error

This provides a foundation for runtime kernel introspection and better tooling support for analyzing and debugging kernel configurations.

[ROCm/composable_kernel commit: 9ac51aa0f4]

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

..

codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

Fix template parameter macros (#3305)

2025-11-26 09:48:17 -08:00

device_avgpool2d_bwd_nhwc_nhwc.hpp

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

2025-11-18 10:17:18 -08:00

device_avgpool3d_bwd_ndhwc_ndhwc.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_contraction_multiple_d_wmma_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_contraction_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_e_permute_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_gemm_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_gemm_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_multi_d_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_multiple_d_dl.hpp

Add support for RDNA1 GPUs (#3220)

2025-11-20 10:45:57 -08:00

device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_multiple_d_xdl_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_reduce_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_wmma_cshuffle_v3_b_scale.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_batched_gemm_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_xdl_fpAintB_b_scale.hpp

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

2025-11-18 10:17:18 -08:00

device_batched_gemm_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_batchnorm_backward_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_batchnorm_forward_impl_obsolete.hpp

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

2025-11-18 10:17:18 -08:00

device_batchnorm_forward_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_cgemm_4gemm_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_column_to_image_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_contraction_multiple_abd_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_contraction_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_contraction_utils.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_backward_weight_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp

[CK, CK_TILE] Add GPU Reference Implementations for Grouped Convolution (#3216)

2025-12-03 21:14:21 +02:00

device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp

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

2025-11-18 10:17:18 -08:00

device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp

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

2025-11-18 10:17:18 -08:00

device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_elementwise_dynamic_vector_dims_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_elementwise_normalization_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_elementwise_scale_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_fpAintB_gemm_wmma.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_bias_add_reduce_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_dl.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_dpp.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_abd_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_abd_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_dl.hpp

Add support for RDNA1 GPUs (#3220)

2025-11-20 10:45:57 -08:00

device_gemm_multiple_d_layernorm_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_wmma_cshuffle_v3_ab_scale.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_multiple_d_wmma_cshuffle_v3_b_preshuffle.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_multiple_d_wmma_cshuffle_v3_blockscale_bpreshuffle.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_multiple_d_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_wmma_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_xdl_cshuffle_v3_ab_scale.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_multiple_d_xdl_cshuffle_v3_b_preshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_xdl_cshuffle_v3_blockscale_bpreshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_xdl_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_reduce_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_reduce_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_wmma_cshuffle_v3_b_scale.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_wmma_cshuffle_v3_common.hpp

Wmma support for gemm_ab_scale (#3314)

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

device_gemm_wmma_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_wmma_cshuffle_v3r1.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_wmma.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_lds_direct_load.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_streamk_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v2.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v3_b_preshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v3_b_scale.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v3_mx.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle_v3r1.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_layernorm_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_skip_b_lds.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_splitk_c_shuffle_lds_direct_load.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_splitk_c_shuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_streamk.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl_waveletmodel_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_gemm_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_contraction_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp

Fix template parameter macros (#3305)

2025-11-26 09:48:17 -08:00

device_grouped_conv_bwd_weight_dl.hpp

Add support for RDNA1 GPUs (#3220)

2025-11-20 10:45:57 -08:00

device_grouped_conv_bwd_weight_explicit_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_weight_wmma_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_weight_xdl_cshuffle_v3.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_bwd_weight_xdl_cshuffle.hpp

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

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

device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp

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

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

device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp

Add support for RDNA1 GPUs (#3220)

2025-11-20 10:45:57 -08:00

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp

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

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

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

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

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

device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_fwd_multiple_d_multiple_r.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp

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

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

device_grouped_conv_fwd_multiple_d_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp

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

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

device_grouped_conv_utils.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_multi_abd_xdl_fixed_nk.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_multiple_d_dl.hpp

Add support for RDNA1 GPUs (#3220)

2025-11-20 10:45:57 -08:00

device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_two_stage.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_wmma_splitk_cshuffle_v3.hpp

Add grouped gemm instances for RDNA4 (#3237)

2025-12-01 15:32:10 -08:00

device_grouped_gemm_xdl_fixed_nk.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_xdl_splitk_cshuffle.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_gemm_xdl.hpp

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

2025-11-18 10:17:18 -08:00

device_grouped_query_attention_forward_wmma.hpp

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

2025-11-18 10:17:18 -08:00

device_image_to_column_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_max_pool_bwd_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_moe_gemm_blockscale.hpp

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

2025-11-18 10:17:18 -08:00

device_moe_gemm.hpp

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

2025-11-18 10:17:18 -08:00

device_moe_mx_gemm_bns.hpp

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

2025-11-18 10:17:18 -08:00

device_moe_mx_gemm_bpreshuffle.hpp

fix static assert (#3178)

2025-11-20 17:27:05 -08:00

device_moe_mx_gemm.hpp

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

2025-11-18 10:17:18 -08:00

device_multi_query_attention_forward_wmma.hpp

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

2025-11-18 10:17:18 -08:00

device_multiple_reduce_multiblock.hpp

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

2025-11-18 10:17:18 -08:00

device_multiple_reduce_threadwise.hpp

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

2025-11-18 10:17:18 -08:00

device_normalization_bwd_data_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_normalization_bwd_gamma_beta_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_normalization_fwd_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_normalization_fwd_splitk_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_permute_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_pool2d_fwd_nhwc_nhwc.hpp

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

2025-11-18 10:17:18 -08:00

device_pool3d_fwd_ndhwc_ndhwc.hpp

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

2025-11-18 10:17:18 -08:00

device_put_element_impl.hpp

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

2025-11-18 10:17:18 -08:00

device_reduce_common.hpp

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

2025-11-18 10:17:18 -08:00

device_reduce_multiblock.hpp

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

2025-11-18 10:17:18 -08:00

device_reduce_threadwise_multi_d.hpp

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

2025-11-18 10:17:18 -08:00

device_reduce_threadwise.hpp

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

2025-11-18 10:17:18 -08:00

device_softmax_impl.hpp

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

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

device_sparse_embeddings_forward_layernorm.hpp

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

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

device_splitk_contraction_multiple_d_xdl_cshuffle.hpp

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

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

split_k_arg.hpp

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

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

split_k_utils.hpp

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

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