You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
{{ message }}
Repository navigation
[Issue]: [CK][CK BUILDER] CK Builder missing DeviceGroupedConvBwdWeight_Explicit support #4940
The builder cannot construct DeviceGroupedConvBwdWeight_Explicit instances. The CK instance library ships ~24 .cpp files instantiating these in its bwd weight conv factory, but the builder has no path to them.
DeviceGroupedConvBwdWeight_Explicit (include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_explicit.hpp) wraps a batched GEMM V3 op to perform 1x1/stride-1/pad-0 backward weight convolutions as explicit GEMM. It has an optional two-stage path with workspace when sizeof(WeiDataType) % 4 != 0. Covers FP16/BF16, 2D and 3D, NHWGC-family layouts only.
The instantiation chain is:
GEMM config tuples defined in .hpp headers (device_exp_gemm_xdl_universal_km_kn_mn_instance.hpp, device_exp_gemm_wmma_universal_km_kn_mn_instance.hpp)
.cpp files under library/src/tensor_operation_instance/gpu/grouped_convnd_bwd_weight/explicit_xdl/ and explicit_wmma/ call add_explicit_gemm_device_operation_instances() with those tuples to register concrete instances (15 XDL, 6 WMMA, 3 odd-dimension variants)
Operating System
All OSes
CPU
All CPUs
GPU
All GPUs
ROCm Version
ROCm 7.x
ROCm Component
composable_kernel
Steps to Reproduce
No response
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
Note: as part of my work porting the factories and underlying machinery to use the builder, I can convert the factory itself, the add_* functions, and the device_* tuples to use the builder, but would like assistance from CK to add support for these operations to the builder.
Activity
DarylHawkinsAMD commentedon Feb 26, 2026
Note: as part of my work porting the factories and underlying machinery to use the builder, I can convert the factory itself, the add_* functions, and the device_* tuples to use the builder, but would like assistance from CK to add support for these operations to the builder.