* chore(copyright) update library wide CMakeLists.txt files copyright header template * Fix build --------- Co-authored-by: Sami Remes <samremes@amd.com>
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.