Files
derek.holloway 7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
Flatten subprojects into a single base monorepo
2026-09-02 16:34:47 -07:00
..

Client Example: Grouped N-Dimensional Convolution Forward

Theory

This client example demonstrates grouped N-dimensional convolution forward for 1D, 2D, and 3D inputs, supporting multiple data types (including BF8 and FP8). Grouped convolution is used in modern CNNs and vision transformers to reduce computation and enable channel-wise or expert-wise processing.

Mathematical Formulation: Given input X and weights W for G groups:

  • For each group g:
    
    Y^g[n, c_{out}, ...] = \sum_{c_{in}} \sum_{k_1} ... \sum_{k_n} X^g[n, c_{in}, ...] \cdot W^g[c_{out}, c_{in}, ...]
    
  • Each group operates on a subset of input/output channels.

Algorithmic Background:

Grouped Convolution Forward

Grouped convolution operation for 1D, 2D or 3D spatial dimensions. Convolution utilizes GEMM kernel after tensor coordinate transform. In CK Grouped Convolution Forward operation is called as DeviceGroupedConvFwdMultipleABD and requires following types as template parameters:

  • NumDimSpatial - number of spatial dimensions (1D, 2D, 3D).
  • InLayout - input layout (NHWGC, GNHWC, NGCHW).
  • WeiLayout - weight layout (GKYXC).
  • DsLayout - layouts for additional tensors for fused operations.
  • OutLayout - output layout (NHWGK, GNHWK, NGKHW).
  • ADataType - input data type. Pass tuple if there is fused operation with input.
  • BDataType - weight data type. Pass tuple if there is fused operation with weight.
  • DsDataType - data types for additional tensors for fused operations.
  • EDataType - Output data type.
  • AElementwiseOperation - fused operation on tensor A (input).
  • BElementwiseOperation - fused operation on tensor B (weight).
  • CDEElementwiseOperation - fused operation on tensor C (output).
  • AComputeType - compute data type of tensor A for mfma instruction (ADataType by default).
  • BComputeType - compute data type of tensor B for mfma instruction (AComputeType by default).

Grouped convolution forward support tensors larger than 2GB.

List of the device operations for grouped convolution forward in CK:

  • DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 - Device operation with XDL instructions. Optimized for AMD Instinct MI300 series.
  • DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle - Device operation with XDL instructions and support of fused operations to input, weight and output.
  • DeviceGroupedConvFwdMultipleD_Wmma_CShuffle - Device operation with WMMA instructions.
  • DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK - Device operation with DL instructions.

Table of supported cases by instance factory with XDL instruction:

NHWGC/GKYXC/NHWGK NGCHW/GKYXC/NGKHW NGCHW/GKCYX/NGKHW GNHWC/GKYXC/GNHWK
bf16 2D, 3D 2D 2D 1D, 2D, 3D
fp16 2D, 3D 2D 2D 1D, 2D, 3D
fp32 2D, 3D 2D 2D 1D, 2D, 3D
int8 2D, 3D 2D 2D 1D, 3D
fp8 3D ✗ ✗ ✗
bf8 3D ✗ ✗ ✗

Table of supported cases by instance factory with WMMA instruction:

NHWGC/GKYXC/NHWGK NGCHW/GKYXC/NGKHW GNHWC/GKYXC/GNHWK
fp16 2D, 3D ✗ 2D, 3D
int8 2D, 3D ✗ 2D, 3D

Table of supported cases by instance factory with DL instruction:

NHWGC/GKYXC/NHWGK NGCHW/GKYXC/NGKHW GNHWC/GKYXC/GNHWK
bf16 ✗ ✗ 2D
fp16 ✗ ✗ 2D
fp32 ✗ ✗ 2D
int8 ✗ ✗ 2D

Table of supported cases by instance factory with fused elementwise operation:

  • Dynamic elementwise operation - 2D/3D, NHWGC, bf16/fp16/fp32/int8
  • Bilinear - 3D, NHWGC, bf16/fp16/fp32/int8
  • ConvInvScale - 3D, NHWGC, fp8
  • ConvScale - 3D, NHWGC, fp8/bf8
  • ConvScale + Add - 3D, NHWGC, fp8
  • ConvScale + Relu - 3D, NHWGC, fp8
  • Scale - 3D, NHWGC, bf16/fp16/fp32/int8
  • Scale + Add (for A and B) - 3D, NHWGC, bf16/fp16/fp32/int8
  • Scale + Add + Scale + Add + Relu - 3D, NHWGC, bf16/fp16/fp32/int8

How to Run

Prerequisites

Please follow the instructions in the main Build Guide section as a prerequisite to building and running this example.

Build and run

cd composable_kernel/client_example/07_grouped_convnd_fwd
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j

# Example run (2D grouped convolution)
./grouped_conv2d_fwd

# Example run (3D grouped convolution, BF8)
./grouped_conv3d_fwd_bf8

# Example run (3D grouped convolution, FP8)
./grouped_conv3d_fwd_fp8

Source Code Structure

Directory Layout

client_example/07_grouped_convnd_fwd/
├── grouped_conv1d_fwd.cpp         # 1D grouped convolution
├── grouped_conv2d_fwd.cpp         # 2D grouped convolution (NCHW)
├── grouped_conv2d_fwd_ngchw.cpp   # 2D grouped convolution (NGCHW)
├── grouped_conv3d_fwd_bf8.cpp     # 3D grouped convolution (BF8)
├── grouped_conv3d_fwd_fp8.cpp     # 3D grouped convolution (FP8)
├── grouped_conv3d_fwd_bf8_fp8.cpp # 3D grouped convolution (BF8/FP8 mixed)
├── grouped_conv3d_fwd_fp8_bf8.cpp # 3D grouped convolution (FP8/BF8 mixed)
├── common.hpp                     # Common utilities for grouped convolution
├── CMakeLists.txt                 # Build configuration for the example

Key Functions

  • main() (in each .cpp):
    Sets up input tensors, configures grouped convolution parameters, launches the kernel, and verifies the result.
  • Grouped convolution kernel invocation:
    Uses the Composable Kernel device API to launch grouped convolution for different dimensions and data types.

This client example provides a comprehensive demonstration of grouped convolution for efficient CNN and vision transformer models.