ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-18 01:28:27 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'develop' into tests_for_batched_grouped_gemm

Browse Source
This commit is contained in:
Aleksander Dudek
2025-11-28 10:28:57 +00:00
parent 3cdc751c32 f875ab0bbc
commit 9302ea4ace
5077 changed files with 129713 additions and 36007 deletions
Download Patch File Download Diff File Expand all files Collapse all files

27
.github/scripts/therock_configure_ci.py vendored
View File

@@ -6,6 +6,7 @@ import subprocess
import sys
from typing import Iterable, Optional, Mapping
def gha_set_output(vars: Mapping[str, str | Path]):
"""Sets values in a step's output parameters.
@@ -25,6 +26,7 @@ def gha_set_output(vars: Mapping[str, str | Path]):
with open(step_output_file, "a") as f:
f.writelines(f"{k}={str(v)}" + "\n" for k, v in vars.items())
def get_modified_paths(base_ref: str) -> Optional[Iterable[str]]:
"""Returns the paths of modified files relative to the base reference."""
try:
@@ -42,11 +44,13 @@ def get_modified_paths(base_ref: str) -> Optional[Iterable[str]]:
file=sys.stderr,
)
return None
GITHUB_WORKFLOWS_CI_PATTERNS = [
"therock*",
]
def is_path_workflow_file_related_to_ci(path: str) -> bool:
return any(
fnmatch.fnmatch(path, ".github/workflows/" + pattern)
@@ -56,11 +60,13 @@ def is_path_workflow_file_related_to_ci(path: str) -> bool:
for pattern in GITHUB_WORKFLOWS_CI_PATTERNS
)
def check_for_workflow_file_related_to_ci(paths: Optional[Iterable[str]]) -> bool:
if paths is None:
return False
return any(is_path_workflow_file_related_to_ci(p) for p in paths)
# Paths matching any of these patterns are considered to have no influence over
# build or test workflows so any related jobs can be skipped if all paths
# modified by a commit/PR match a pattern in this list.
@@ -70,23 +76,26 @@ SKIPPABLE_PATH_PATTERNS = [
"*.md",
"*.pre-commit-config.*",
"*LICENSE",
'Jenkinsfile',
'.github/ISSUE_TEMPLATE/*',
'.github/CODEOWNERS',
'.github/*.md',
'.github/dependabot.yml',
"Jenkinsfile",
".github/ISSUE_TEMPLATE/*",
".github/CODEOWNERS",
".github/*.md",
".github/dependabot.yml",
]
def is_path_skippable(path: str) -> bool:
"""Determines if a given relative path to a file matches any skippable patterns."""
return any(fnmatch.fnmatch(path, pattern) for pattern in SKIPPABLE_PATH_PATTERNS)
def check_for_non_skippable_path(paths: Optional[Iterable[str]]) -> bool:
"""Returns true if at least one path is not in the skippable set."""
if paths is None:
return False
return any(not is_path_skippable(p) for p in paths)
def should_ci_run_given_modified_paths(paths: Optional[Iterable[str]]) -> bool:
"""Returns true if CI workflows should run given a list of modified paths."""
@@ -118,16 +127,16 @@ def should_ci_run_given_modified_paths(paths: Optional[Iterable[str]]) -> bool:
)
return False
def main(args):
base_ref = args.get("base_ref")
modified_paths = get_modified_paths(base_ref)
print("modified_paths (max 200):", modified_paths[:200])
enable_jobs = should_ci_run_given_modified_paths(modified_paths)
output = {
'enable_therock_ci': json.dumps(enable_jobs)
}
output = {"enable_therock_ci": json.dumps(enable_jobs)}
gha_set_output(output)
if __name__ == "__main__":
args = {}
args["base_ref"] = os.environ.get("BASE_REF", "HEAD^1")

16
.github/workflows/pre-commit.yml vendored Normal file
View File

@@ -0,0 +1,16 @@
name: pre-commit
on:
pull_request:
push:
branches: [develop]
jobs:
pre-commit:
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@v3
- uses: actions/setup-python@v3
with:
python-version: '3.12'
- uses: pre-commit/action@v3.0.1

17
.github/workflows/therock-ci-linux.yml vendored
View File

@@ -20,7 +20,7 @@ jobs:
permissions:
id-token: write
container:
image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:044b113562629f4bd2ec5d2e64b32eee11562d48fb1a75d7493daec9dd8d8292
image: ghcr.io/rocm/therock_build_manylinux_x86_64@sha256:2f3ebd0beb04c449fdb36933e54bdc69483b914fb9005594d3fc9444c206b54b
options: -v /runner/config:/home/awsconfig/
env:
AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }}
@@ -35,6 +35,15 @@ jobs:
with:
repository: "ROCm/rocm-libraries"
- name: Pull DVC files for rocm-libraries # LOGNAME details here https://github.com/ROCm/rocm-libraries/pull/1617
run: |
if command -v dvc &> /dev/null; then
echo "dvc detected"
else
echo "Warning, dvc not detected!"
fi
LOGNAME=github-runner dvc pull -v
- name: Checkout composable_kernel repository
uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
with:
@@ -44,8 +53,8 @@ jobs:
uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
with:
repository: "ROCm/TheRock"
ref: dc05d637054ad197c84b00e24b6262af0ec797c6 # 10-03-2025 commit
path: "TheRock"
ref: f3f77a3161922df3eee006b888b439d75b2b4668 # 2025-10-29 commit
- name: Setup ccache
run: |
@@ -68,6 +77,8 @@ jobs:
- name: Patch rocm-libraries
run: |
git config --global --add safe.directory '*'
# Remove patches here if they cannot be applied cleanly, and they have not been deleted from TheRock repo
rm -f ./TheRock/patches/amd-mainline/rocm-libraries/0008-Revert-remove-options-no-enumerate-966.patch
git -c user.name="therockbot" -c "user.email=therockbot@amd.com" am --whitespace=nowarn ./TheRock/patches/amd-mainline/rocm-libraries/*.patch
- name: Install python deps
@@ -119,7 +130,7 @@ jobs:
run: |
python3 TheRock/build_tools/github_actions/post_build_upload.py \
--run-id ${{ github.run_id }} \
--amdgpu-family ${{ env.AMDGPU_FAMILIES }} \
--artifact-group ${{ env.AMDGPU_FAMILIES }} \
--build-dir TheRock/build \
--upload

5
.github/workflows/therock-test-component.yml vendored
View File

@@ -29,7 +29,7 @@ jobs:
--group-add video
--device /dev/kfd
--device /dev/dri
--group-add 992
--group-add 110
--env-file /etc/podinfo/gha-gpu-isolation-settings
strategy:
fail-fast: false
@@ -51,12 +51,13 @@ jobs:
uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
with:
repository: "ROCm/TheRock"
ref: f3f77a3161922df3eee006b888b439d75b2b4668 # 2025-10-29 commit
- name: Run setup test environment workflow
uses: './.github/actions/setup_test_environment'
with:
ARTIFACT_RUN_ID: ${{ env.ARTIFACT_RUN_ID }}
AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }}
ARTIFACT_GROUP: ${{ inputs.amdgpu_families }}
OUTPUT_ARTIFACTS_DIR: ${{ env.OUTPUT_ARTIFACTS_DIR }}
VENV_DIR: ${{ env.VENV_DIR }}
FETCH_ARTIFACT_ARGS: ${{ fromJSON(inputs.component).fetch_artifact_args }}

1
.github/workflows/therock-test-packages.yml vendored
View File

@@ -27,6 +27,7 @@ jobs:
uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
with:
repository: "ROCm/TheRock"
ref: f3f77a3161922df3eee006b888b439d75b2b4668 # 2025-10-29 commit
- name: "Configuring CI options"
env:

17
.gitignore vendored
View File

@@ -36,7 +36,10 @@ tags
# Editors
.vscode
# build-in-source directory
# Cline
.cline*
# build-in-source directory (see exceptions below)
build*
# emacs temporary/backup files
@@ -58,11 +61,17 @@ _doxygen/
docs/doxygen/html
docs/doxygen/xml
# JetBrains IDE
# JetBrains IDE (see build* exceptions below)
.idea/
cmake-build*/
build*/
# LSP configuration
.clangd
# User-defined CMake presets
CMakeUserPresets.json
# Python virtualenv
.venv/
@@ -71,3 +80,7 @@ __pycache__/
.cache/
# Exceptions to build* patterns above
# The experimental/builder directory should be tracked despite matching build*
!experimental/builder
!experimental/builder/**

45
.pre-commit-config.yaml
View File

@@ -1,11 +1,25 @@
repos:
- repo: local
- repo: https://github.com/pre-commit/mirrors-clang-format
rev: v18.1.3
hooks:
- id: clang-format
name: clang-format
entry: clang-format-18 -i --style=file
language: system
types_or: [c++, inc]
- repo: https://github.com/astral-sh/ruff-pre-commit
rev: v0.14.0
hooks:
- id: ruff-check
args: [ --fix ]
exclude: |
(?x)^(
docs/conf.py
)$
- id: ruff-format
exclude: |
(?x)^(
docs/conf.py
)$
- repo: local
hooks:
# - id: copyright-year-checker
# name: copyright-year-checker
# entry: script/check_copyright_year.sh
@@ -18,21 +32,12 @@ repos:
language: script
types_or: [c++, text]
verbose: true
- id: ruff-check
name: Ruff Linter
entry: ruff check --fix
- id: remod-ck-tile
name: Run ck_tile remod.py
entry: python script/remod_for_ck_tile.py
language: python
types: [python]
additional_dependencies: [ruff]
- id: ruff-format
name: Ruff Formatter
entry: ruff format
language: python
types: [python]
additional_dependencies: [ruff]
- id: run-remod-if-ck-tile-changed
name: Run remod.py if ck_tile files changed
entry: script/remod_for_ck_tile.sh
language: script
always_run: true
files: '^(include|example)/ck_tile/.*$'
additional_dependencies:
- dos2unix
- clang-format==18.1.3
pass_filenames: false

67
ACRONYMS.md Normal file
View File

@@ -0,0 +1,67 @@
# Acronyms in Composable Kernel
The following acronyms are used in the Composable Kernel codebase:
| Acronym | Expansion | Explanation |
|---------|-----------|-------------|
| BF16 | Brain Floating Point 16 | 1 Signed bit, 8 Exponent bits, 7 Significand bits |
| BF8 | 8-bit Brain Floating Point | 1 Signed bit, 3 Exponent bits, 4 Significand bits |
| DLA | Deep Learning Accelerator | Specialized hardware for deep learning workloads |
| DRAM | Dynamic Random-Access Memory | Main memory. Global memory on GPU |
| E2E | End-to-End | Complete pipeline or process from input to output |
| ELU | Exponential Linear Unit | Activation function: $x$ if $x>0$ else $\alpha(e^x-1)$ |
| FMHA | Fused Multi-Head Attention | Efficient transformer attention kernel, fusing softmax, masking, and matmul |
| FP16 | Half-Precision Floating Point | 16-bit IEEE floating point format |
| FP32 | Single-Precision Floating Point | 32-bit IEEE floating point format |
| FP64 | Double-Precision Floating Point | 64-bit IEEE floating point format |
| FP8 | 8-bit Floating Point | Experimental 8-bit floating point format for inference |
| GEMM | General Matrix Multiply | Matrix multiplication operation: $C = A \times B$ |
| GELU | Gaussian Error Linear Unit | Activation function: $x \cdot \Phi(x)$ |
| GQA | Grouped Query Attention | Variant of multi-head attention with grouped queries/keys/values |
| HBM | High Bandwidth Memory | Fast memory used in modern GPUs |
| HIP | Heterogeneous-Compute Interface for Portability | AMD's CUDA-like GPU programming API |
| INT8 | 8-bit Integer | Quantized integer format for inference |
| KVS | Key-Value Store | Data structure for storing key-value pairs (context: QKV in transformers) |
| L2/L1 | Level 2/Level 1 Cache | On-chip memory hierarchy in CPUs/GPUs |
| LDS | Local Data Share | Shared memory on AMD GPUs (equivalent to CUDA's shared memory) |
| LLM | Large Language Model | Transformer-based model for NLP tasks |
| LSE | Log-Sum-Exp | Numerically stable softmax computation: $\log(\sum \exp(x))$ |
| MHA | Multi-Head Attention | Attention mechanism with multiple heads in transformers |
| MFMA | Matrix Fused Multiply-Add | AMD GPU hardware instruction for matrix-matrix multiplication |
| MoE | Mixture of Experts | Neural network architecture with multiple expert subnetworks |
| MQA | Multi-Query Attention | Variant of multi-head attention with shared keys/values across heads |
| RCCL | ROCm Collective Communications Library | AMD Library for multi-GPU communication |
| NCHW | Batch, Channel, Height, Width | Tensor layout: batch-major, channels-first |
| NHWC | Batch, Height, Width, Channel | Tensor layout: batch-major, channels-last |
| OOM | Out Of Memory | Error when memory allocation fails |
| QAT | Quantization Aware Training | Training technique for quantized inference |
| QKV | Query, Key, Value | Components of transformer attention mechanism |
| RDMA | Remote Direct Memory Access | High-speed network memory access |
| RDQuant | Rowwise Dynamic Quantization | Quantization technique with per-row scaling for int8 inference |
| ReLU | Rectified Linear Unit | Activation function: $\max(0, x)$ |
| ROCm | Radeon Open Compute | AMD's open GPU computing stack |
| SGD | Stochastic Gradient Descent | Optimization algorithm for training neural networks |
| SM | Streaming Multiprocessor | GPU compute unit (NVIDIA terminology) |
| SWA | Sliding Window Attention | Attention mechanism with a limited window for each token |
| TLB | Translation Lookaside Buffer | Memory management unit cache for virtual-to-physical address translation |
| VGPR | Vector General Purpose Register | GPU register for vector operations |
| WARP | Group of Threads | Smallest scheduling unit on NVIDIA GPUs (32 threads) |
| WMMA | Warp Matrix Multiply-Accumulate | NVIDIA's matrix-multiply hardware primitive |
| XLA | Accelerated Linear Algebra | Compiler for optimizing ML computations (Google) |
### Common Variable Acronyms in Code
| Symbol | Meaning | Context |
|--------|---------|---------|
| M, N, K | Matrix dimensions | GEMM: $A[M,K] \times B[K,N] = C[M,N]$ |
| Q, K, V | Query, Key, Value | Transformer attention |
| S | Sequence length | NLP, transformers |
| D | Dimension | Hidden size, feature dim |
| B | Batch size | ML batch processing |
| H | Head count | Multi-head attention |
| C | Channel | CNNs, tensor layouts |
| T | Token | NLP, sequence models |
---
If you find an acronym not listed here, please submit a pull request or issue!

78
CHANGELOG.md
View File

@@ -2,15 +2,60 @@
Documentation for Composable Kernel available at [https://rocm.docs.amd.com/projects/composable_kernel/en/latest/](https://rocm.docs.amd.com/projects/composable_kernel/en/latest/).
## Composable Kernel 1.2.0 for ROCm 7.0.0
## Composable Kernel 1.2.0 for ROCm 7.2.0
### Added
* Added support for bf16 data type to grouped_gemm and grouped_gemm_preshuffle.
* Added support for mixed precision fp8 x bf8 universal GEMM and weight preshuffle GEMM
* Added a compute async pipeline in the CK TILE universal GEMM on gfx950
* Added support for B Tensor type pk_int4_t in the CK TILE weight preshuffle GEMM.
* Added the new api to load different memory sizes to SGPR.
* Added support for B Tensor Preshuffle in CK TILE Grouped GEMM.
* Added a basic copy kernel example and supporting documentation for new CK Tile developers.
* Added support for grouped_gemm kernels to perform multi_d elementwise operation.
* Added support for Multiple ABD GEMM
* Added benchmarking support for tile engine GEMM Multi D.
* Added block scaling support in CK_TILE GEMM, allowing flexible use of quantization matrices from either A or B operands.
* Added the row-wise column-wise quantization for CK_TILE GEMM & CK_TILE Grouped GEMM.
* Added support for f32 to FMHA (fwd/bwd).
* Added tensor-wise quantization for CK_TILE GEMM.
* Added support for batched contraction kernel.
* Added WMMA (gfx12) support for FMHA.
* Added pooling kernel in CK_TILE
* Added top-k sigmoid kernel in CK_TILE
* Added the blockscale 2D support for CK_TILE GEMM.
### Changed
* Removed `BlockSize` in `make_kernel` and `CShuffleEpilogueProblem` to support Wave32 in CK_TILE (#2594)
* Added an optional template parameter `Arch` (`gfx9_t`, `gfx12_t` etc.) to `make_kernel` to support linking multiple object files that have the same kernel compiled for different architectures.
* FMHA examples and tests can be built for multiple architectures (gfx9, gfx950, gfx12) at the same time.
### Upcoming changes
* Composable Kernel will be adopting C++20 features in an upcoming ROCm release, updating the minimum compiler requirement to C++20. Ensure that your development environment complies with this requirement to facilitate a seamless transition.
## Composable Kernel 1.1.0 for ROCm 7.1.1
### Upcoming changes
* Composable Kernel will be adopting C++20 features in an upcoming ROCm release, updating the minimum compiler requirement to C++20. Ensure that your development environment complies with this requirement to facilitate a seamless transition.
## Composable Kernel 1.1.0 for ROCm 7.1.0
### Added
* Added support for hdim as a multiple of 32 for FMHA (fwd/fwd_splitkv/bwd)
* Added support for elementwise kernel.
### Upcoming changes
* Non-grouped convolutions are deprecated. Their functionality is supported by grouped convolution.
## Composable Kernel 1.1.0 for ROCm 7.0.0
### Added
* Added support for bf16, f32, and f16 for 2D and 3D NGCHW grouped convolution backward data
* Added a fully asynchronous HOST (CPU) arguments copy flow for CK grouped GEMM kernels.
* Added support GKCYX layout for grouped convolution forward (NGCHW/GKCYX/NGKHW, number of instances in instance factory for NGCHW/GKYXC/NGKHW has been reduced).
@@ -19,37 +64,23 @@ Documentation for Composable Kernel available at [https://rocm.docs.amd.com/proj
* Added support for GKCYX layout for grouped convolution backward data (NGCHW/GKCYX/NGKHW).
* Added support for Stream-K version of mixed fp8/bf16 GEMM
* Added support for Multiple D GEMM
* Added support for Multiple ABD GEMM
* Added GEMM pipeline for microscaling (MX) FP8/FP6/FP4 data types
* Added support for FP16 2:4 structured sparsity to universal GEMM.
* Added support for Split K for grouped convolution backward data.
* Added logit soft-capping support for fMHA forward kernels.
* Added support for hdim as a multiple of 32 for FMHA (fwd/fwd_splitkv)
* Added support for hdim as a multiple of 32 for FMHA (fwd/fwd_splitkv/bwd)
* Added benchmarking support for tile engine GEMM.
* Added Ping-pong scheduler support for GEMM operation along the K dimension.
* Added rotating buffer feature for CK_Tile GEMM.
* Added int8 support for CK_TILE GEMM.
* Added support for elementwise kernel.
* Added benchmarking support for tile engine GEMM Multi D.
* Added block scaling support in CK_TILE GEMM, allowing flexible use of quantization matrices from either A or B operands.
* Added the row-wise column-wise quantization for CK_TILE GEMM & CK_TILE Grouped GEMM.
* Added support for f32 to FMHA (fwd/bwd).
* Added tensor-wise quantization for CK_TILE GEMM.
### Optimized
* Optimize the gemm multiply multiply preshuffle & lds bypass with Pack of KGroup and better instruction layout.
* Added Vectorize Transpose optimization for CK Tile
* Added the asynchronous copy for gfx950
* Optimize the gemm multiply multiply preshuffle & lds bypass with Pack of KGroup and better instruction layout. (#2166)
* Added Vectorize Transpose optimization for CK Tile (#2131)
* Added the asynchronous copy for gfx950 (#2425)
### Fixes
None
### Changes
### Changed
* Removed support for gfx940 and gfx941 targets (#1944)
* Replaced the raw buffer load/store intrinsics with Clang20 built-ins (#1876)
@@ -57,15 +88,6 @@ None
* Number of instances in instance factory for grouped convolution forward NGCHW/GKYXC/NGKHW has been reduced.
* Number of instances in instance factory for grouped convolution backward weight NGCHW/GKYXC/NGKHW has been reduced.
* Number of instances in instance factory for grouped convolution backward data NGCHW/GKYXC/NGKHW has been reduced.
* Removed `BlockSize` in `make_kernel` and `CShuffleEpilogueProblem` to support Wave32 in CK_TILE (#2594)
### Known issues
None
### Upcoming changes
* Non-grouped convolutions are deprecated. All of their functionality is supported by grouped convolution.
## Composable Kernel 1.1.0 for ROCm 6.1.0

25
CMakeLists.txt
View File

@@ -37,8 +37,14 @@ include(CTest)
option(ENABLE_CLANG_CPP_CHECKS "Enables clang tidy, cppcheck" ON)
option(MIOPEN_REQ_LIBS_ONLY "Build only the MIOpen required libraries" OFF)
option(CK_EXPERIMENTAL_BUILDER "Enable experimental builder" OFF)
option(BUILD_MHA_LIB "Build the static library for flash attention" OFF)
if(CK_EXPERIMENTAL_BUILDER)
add_definitions(-DCK_EXPERIMENTAL_BUILDER)
include_directories(${PROJECT_SOURCE_DIR}/experimental/builder/include)
endif()
# Usage: for customized Python location cmake -DCK_USE_ALTERNATIVE_PYTHON="/opt/Python-3.8.13/bin/python3.8"
# CK Codegen requires dataclass which is added in Python 3.7
# Python version 3.8 is required for general good practice as it is default for Ubuntu 20.04
@@ -116,7 +122,7 @@ add_compile_options(
# Recent change in compiler makes this warning ON by default, which led to compile errors.
add_compile_options(-Wno-nrvo)
if(NOT DISABLE_DL_KERNELS)
if(NOT DISABLE_DL_KERNELS AND GPU_TARGETS MATCHES "gfx101|gfx103|gfx10-1|gfx10-3")
add_definitions(-DDL_KERNELS)
set(DL_KERNELS "ON")
set(CK_ENABLE_DL_KERNELS "ON")
@@ -340,7 +346,6 @@ endif()
option(USE_BITINT_EXTENSION_INT4 "Whether to enable clang's BitInt extension to provide int4 data type." OFF)
option(USE_OPT_GFX11 "Whether to enable LDS cumode and Wavefront32 mode for GFX11 silicons." OFF)
option(ENABLE_ASM_DUMP "Whether to enable assembly dump for kernels." OFF)
option(ENABLE_JSON_DUMP "Whether to enable json dump for examples." OFF)
@@ -611,11 +616,11 @@ endif()
if(NOT MIOPEN_REQ_LIBS_ONLY)
# make check runs the entire set of examples and tests
add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR})
add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} USES_TERMINAL)
# make smoke runs the tests and examples that runs within 30 seconds on gfx90a
add_custom_target(smoke COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "SMOKE_TEST")
add_custom_target(smoke COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "SMOKE_TEST" USES_TERMINAL)
# make regression runs the tests and examples that runs for more 30 seconds on gfx90a
add_custom_target(regression COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "REGRESSION_TEST")
add_custom_target(regression COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "REGRESSION_TEST" USES_TERMINAL)
endif()
@@ -678,6 +683,12 @@ if(NOT GPU_ARCHS AND USER_GPU_TARGETS AND NOT MIOPEN_REQ_LIBS_ONLY)
PACKAGE_NAME examples
)
add_subdirectory(example)
add_subdirectory(tutorial)
rocm_package_setup_component(tutorials
LIBRARY_NAME composablekernel
PACKAGE_NAME tutorials
)
add_subdirectory(tile_engine)
if(BUILD_TESTING)
add_subdirectory(test)
@@ -692,6 +703,10 @@ if (NOT MIOPEN_REQ_LIBS_ONLY)
add_subdirectory(profiler)
endif()
if (CK_EXPERIMENTAL_BUILDER)
add_subdirectory(experimental/builder)
endif()
if(CK_USE_CODEGEN AND (SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR GPU_ARCHS))
add_subdirectory(codegen)
endif()

13
Dockerfile.aiter
View File

@@ -1,9 +1,8 @@
ARG BASE_DOCKER="rocm/composable_kernel-private:ck_aiter_base"
ARG BASE_DOCKER="rocm/pytorch:latest"
FROM $BASE_DOCKER
ARG AITER_BRANCH="main"
ARG CK_AITER_BRANCH="develop"
RUN groupadd irc && \
pip install pandas zmq einops && \
RUN pip install pandas zmq einops ninja && \
pip install numpy==1.26.2 && \
sudo mkdir /home/jenkins && \
sudo mkdir /home/jenkins/workspace && \
@@ -14,6 +13,10 @@ RUN groupadd irc && \
rm -rf 3rdparty/composable_kernel/ && \
git clone -b "$CK_AITER_BRANCH" https://github.com/ROCm/composable_kernel.git 3rdparty/composable_kernel/ && \
python3 setup.py develop && \
chown -R jenkins:jenkins /home/jenkins/workspace && \
chmod -R a+rwx /home/jenkins/workspace && \
groupadd -g 1001 jenkins && \
useradd -u 1001 -g 1001 -m -s /bin/bash jenkins && \
chown -R jenkins:jenkins /home/jenkins && \
chmod -R a+rwx /home/jenkins && \
chown -R jenkins:jenkins /tmp && \
chmod -R a+rwx /tmp && \
sudo usermod -aG irc jenkins

879
Jenkinsfile vendored
View File

File diff suppressed because it is too large Load Diff

35
README.md
View File

@@ -93,13 +93,44 @@ Docker images are available on [DockerHub](https://hub.docker.com/r/rocm/composa
want to build the library for a list of different architectures,
you should use the `GPU_ARCHS` build argument, for example `GPU_ARCHS=gfx908;gfx1030;gfx1100;gfx942`.
4. Build the entire CK library:
**Convenience script for development builds:**
Alternatively, you can use the provided convenience script `script/cmake-ck-dev.sh` which automatically
configures CK for development with sensible defaults. In the build directory:
```bash
../script/cmake-ck-dev.sh
```
This script:
* Cleans CMake cache files before configuring
* Sets `BUILD_DEV=ON` for development mode
* Defaults to GPU targets: `gfx908;gfx90a;gfx942`
* Enables verbose makefile output
* Sets additional compiler flags for better error messages
By default, it considers the parent directory to be the project source directory.
You can specify the source directory as the first argument.
You can specify custom GPU targets (semicolon-separated) as the second argument:
```bash
../script/cmake-ck-dev.sh .. gfx1100
```
Or pass additional cmake arguments:
```bash
../script/cmake-ck-dev.sh .. gfx90a -DCMAKE_BUILD_TYPE=Release
```
5. Build the entire CK library:
```bash
make -j"$(nproc)"
```
5. Install CK:
6. Install CK:
```bash
make -j install

56
client_example/01_gemm/README.md
View File

@@ -1,5 +1,22 @@
[Back to supported operations](../../../include/ck/README.md)
# Composable Kernel GEMM
# Client Example: Basic GEMM
## Theory
This client example demonstrates a basic **GEMM (General Matrix Multiplication)** operation using the Composable Kernel library. GEMM is a core operation in linear algebra and deep learning, computing the product of two matrices and optionally adding a bias or scaling.
**Mathematical Formulation:**
$$
C = \alpha (A \times B) + \beta D
$$
- $A$: [M, K] input matrix
- $B$: [K, N] weight matrix
- $D$: [M, N] optional bias or residual
- $C$: [M, N] output
- $\alpha, \beta$: scalars (often 1.0, 0.0)
**Algorithmic Background:**
- The operation is implemented using a tiled/blocking strategy for memory efficiency.
- GEMM is the computational backbone for transformer attention, MLPs, and CNNs (via im2col).
## GEMM
General matrix multiplications operation. In CK GEMM operation is called as `DeviceGemm` and requires following types as template parameters:
@@ -124,3 +141,38 @@ Table of supported cases by instance factory with XDL instruction for Row/Row/Ro
* **DeviceGemmReduce** - GEMM fused with reduction.
* **DeviceGemm_Streamk_V2** - GEMM stream K implementation. Implementation allows to use reduction instead of AtomicAdd.
* **DeviceGemmStreamK** - GEMM stream K implementation using AtomicAdd.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/01_gemm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run
./gemm
```
## Source Code Structure
### Directory Layout
```
client_example/01_gemm/
├── gemm.cpp # Main client example: sets up, runs, and verifies GEMM
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `gemm.cpp`):
Sets up input matrices, configures GEMM parameters, launches the GEMM kernel, and verifies the result.
- **GEMM kernel invocation**:
Uses the Composable Kernel device API to launch the GEMM operation.
This client example provides a minimal, end-to-end demonstration of using Composable Kernel for matrix multiplication in a user application.

2
client_example/01_gemm/gemm.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

65
client_example/02_gemm_add_add_fastgelu/README.md Normal file
View File

@@ -0,0 +1,65 @@
# Client Example: GEMM with Add, Add, and FastGELU Fusion
## Theory
This client example demonstrates **GEMM fused with two addition operations and FastGELU activation**. This pattern is common in transformer feed-forward networks and other neural architectures where a linear transformation is followed by bias addition, residual addition, and a non-linear activation.
**Mathematical Formulation:**
$$
E = \text{FastGELU}((A \times B) + D_0 + D_1)
$$
- $A$: [M, K] input matrix
- $B$: [K, N] weight matrix
- $D_0$: [N] bias vector (broadcasted)
- $D_1$: [M, N] residual tensor
- $E$: [M, N] output
FastGELU is an efficient approximation of GELU:
$$
\text{FastGELU}(x) = x \cdot \sigma(1.702 \cdot x)
$$
where $\sigma$ is the sigmoid function.
**Algorithmic Background:**
- The GEMM result is kept in registers, bias and residual are added, and FastGELU is applied before writing to global memory.
- No intermediate results are written to global memory.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/02_gemm_add_add_fastgelu
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run
./gemm_add_add_fastgelu
```
## Source Code Structure
### Directory Layout
```
client_example/02_gemm_add_add_fastgelu/
├── gemm_add_add_fastgelu.cpp # Main client example: GEMM+Add+Add+FastGELU
├── gemm_add_add_fastgelu_generic.cpp # Generic variant
├── gemm_add_fastgelu.cpp # GEMM+Add+FastGELU
├── gemm_add_fastgelu_generic.cpp # Generic variant
├── gemm_fastgelu.cpp # GEMM+FastGELU only
├── gemm_fastgelu_generic.cpp # Generic variant
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input matrices, configures GEMM and epilogue parameters, launches the fused kernel, and verifies the result.
- **Fused kernel invocation**:
Uses the Composable Kernel device API to launch the GEMM with fused addition and FastGELU.
This client example provides several variants to demonstrate different levels of fusion and genericity for transformer-style MLP layers.

2
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/02_gemm_add_add_fastgelu/gemm_add_add_fastgelu_generic.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/02_gemm_add_add_fastgelu/gemm_add_fastgelu_generic.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/02_gemm_add_add_fastgelu/gemm_fastgelu_generic.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

57
client_example/03_gemm_layernorm/README.md Normal file
View File

@@ -0,0 +1,57 @@
# Client Example: GEMM with LayerNorm Fusion
## Theory
This client example demonstrates **GEMM fused with layer normalization** and additional elementwise operations. This pattern is common in transformer feed-forward networks and other architectures where a linear transformation is followed by normalization and activation.
**Mathematical Formulation:**
- GEMM: $Y = A \times B$
- Additions: $Z = Y + D_0 + D_1$ (bias, residual, etc.)
- Activation: $A = \text{ReLU}(Z)$ (or other activation)
- LayerNorm: $\text{LayerNorm}(A) = \gamma \cdot \frac{A - \mu}{\sqrt{\sigma^2 + \epsilon}} + \beta$
$\mu$, $\sigma^2$ are mean and variance over the normalization axis; $\gamma$, $\beta$ are learnable scale and shift.
**Algorithmic Background:**
- The GEMM result is kept in registers, elementwise ops and layer normalization are fused in the epilogue.
- LayerNorm is typically applied over the last dimension (features).
- This fusion reduces memory traffic and is common in transformer MLP blocks.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/03_gemm_layernorm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (naive)
./gemm_add_add_layernorm_naive
# Example run (with ReLU and Welford)
./gemm_add_relu_add_layernorm_welford
```
## Source Code Structure
### Directory Layout
```
client_example/03_gemm_layernorm/
├── gemm_add_add_layernorm_naive.cpp # GEMM + Add + Add + LayerNorm (naive)
├── gemm_add_relu_add_layernorm_welford.cpp # GEMM + Add + ReLU + Add + LayerNorm (Welford)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input matrices, configures GEMM and epilogue parameters, launches the fused kernel, and verifies the result.
- **LayerNorm implementation**:
Demonstrates both naive and numerically stable (Welford) algorithms for mean/variance.
This client example provides variants to demonstrate different levels of fusion and normalization for transformer-style MLP layers.

2
client_example/03_gemm_layernorm/gemm_add_add_layernorm_naive.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/03_gemm_layernorm/gemm_add_relu_add_layernorm_welford.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

56
client_example/04_contraction/README.md Normal file
View File

@@ -0,0 +1,56 @@
# Client Example: General Tensor Contraction
## Theory
This client example demonstrates **general tensor contraction** operations, including bilinear and scaled contractions. Tensor contraction generalizes matrix multiplication to higher dimensions and is used in scientific computing, quantum chemistry, and advanced neural network layers.
**Mathematical Formulation:**
- General contraction: $C_{i,j} = \sum_k A_{i,k} \cdot B_{k,j}$
- Bilinear contraction: $C = \alpha (A \cdot B) + \beta D$
- Scale contraction: $C = \text{scale}(A, B)$ (elementwise or broadcasted scaling)
**Algorithmic Background:**
- Contraction can be performed over arbitrary axes and supports broadcasting.
- Bilinear and scale contractions are used for feature fusion, gating, and scientific workloads.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/04_contraction
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (bilinear FP32)
./contraction_bilinear_fp32
# Example run (scale FP64)
./contraction_scale_fp64
```
## Source Code Structure
### Directory Layout
```
client_example/04_contraction/
├── contraction_bilinear_fp32.cpp # Bilinear contraction (FP32)
├── contraction_bilinear_fp64.cpp # Bilinear contraction (FP64)
├── contraction_g1m2n3k1_add_xdl_fp16.cpp # Grouped contraction with addition (FP16)
├── contraction_scale_fp32.cpp # Scale contraction (FP32)
├── contraction_scale_fp64.cpp # Scale contraction (FP64)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures contraction parameters, launches the contraction kernel, and verifies the result.
- **Contraction kernel invocation**:
Uses the Composable Kernel device API to launch the contraction operation.
This client example provides several variants to demonstrate different contraction types and data types for scientific and ML workloads.

2
client_example/04_contraction/contraction_bilinear_fp32.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>

2
client_example/04_contraction/contraction_bilinear_fp64.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>

2
client_example/04_contraction/contraction_g1m2n3k1_add_xdl_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>

2
client_example/04_contraction/contraction_scale_fp32.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>

2
client_example/04_contraction/contraction_scale_fp64.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <numeric>

66
client_example/05_layernorm/README.md Normal file
View File

@@ -0,0 +1,66 @@
# Client Example: Layer Normalization (Forward and Backward)
## Theory
This client example demonstrates **layer normalization** in both forward and backward modes, for 2D and 4D tensors. Layer normalization is used in transformers and other neural networks to normalize activations across the feature dimension, improving training stability.
**Mathematical Formulation:**
Given input $X$:
- Mean: $\mu = \frac{1}{N} \sum_{i=1}^N X_i$
- Variance: $\sigma^2 = \frac{1}{N} \sum_{i=1}^N (X_i - \mu)^2$
- Normalized: $\hat{X}_i = \frac{X_i - \mu}{\sqrt{\sigma^2 + \epsilon}}$
- Output: $Y_i = \gamma \hat{X}_i + \beta$
$\gamma$, $\beta$ are learnable scale and shift parameters.
**Algorithmic Background:**
- Forward pass computes mean, variance, normalization, and affine transformation.
- Backward pass computes gradients with respect to input, gamma, and beta.
- Supports both 2D (batch, feature) and 4D (batch, channel, height, width) tensors.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/05_layernorm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (2D forward)
./layernorm2d_fwd
# Example run (4D forward)
./layernorm4d_fwd
# Example run (2D backward, data)
./layernorm2d_bwd_data
# Example run (2D backward, gamma/beta)
./layernorm2d_bwd_gamma_beta
```
## Source Code Structure
### Directory Layout
```
client_example/05_layernorm/
├── layernorm2d_fwd.cpp # 2D layernorm forward
├── layernorm4d_fwd.cpp # 4D layernorm forward
├── layernorm2d_bwd_data.cpp # 2D layernorm backward (data)
├── layernorm2d_bwd_gamma_beta.cpp # 2D layernorm backward (gamma/beta)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures normalization parameters, launches the forward or backward kernel, and verifies the result.
- **LayerNorm implementation**:
Demonstrates both forward and backward passes for different tensor shapes.
This client example provides a comprehensive demonstration of layer normalization for both inference and training in deep learning models.

2
client_example/05_layernorm/layernorm2d_bwd_data.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/05_layernorm/layernorm2d_bwd_gamma_beta.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/05_layernorm/layernorm2d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/05_layernorm/layernorm4d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

54
client_example/06_softmax/README.md Normal file
View File

@@ -0,0 +1,54 @@
# Client Example: 4D Softmax
## Theory
This client example demonstrates **Softmax computation over 4D tensors**. Softmax is a key operation in deep learning, especially in attention mechanisms and classification, converting logits into normalized probabilities.
**Mathematical Formulation:**
Given input $X$ and axis $a$:
$$
\text{softmax}(X)_i = \frac{\exp(X_i)}{\sum_j \exp(X_j)}
$$
**Algorithmic Background:**
- Softmax is implemented using a numerically stable algorithm:
1. Subtract the maximum value for numerical stability.
2. Exponentiate and sum.
3. Normalize by the sum.
- Efficient parallel Softmax requires careful reduction and memory access patterns.
- This example demonstrates Softmax over a 4D tensor, as used in attention and vision models.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/06_softmax
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run
./softmax4d
```
## Source Code Structure
### Directory Layout
```
client_example/06_softmax/
├── softmax4d.cpp # Main client example: sets up, runs, and verifies 4D softmax
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `softmax4d.cpp`):
Sets up input tensors, configures Softmax parameters, launches the Softmax kernel, and verifies the result.
- **Softmax kernel invocation**:
Uses the Composable Kernel device API to launch the Softmax operation.
This client example provides a demonstration of efficient, numerically stable Softmax for 4D tensors in deep learning models.

2
client_example/06_softmax/softmax4d.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>

66
client_example/07_grouped_convnd_fwd/README.md
View File

@@ -1,5 +1,18 @@
[Back to supported operations](../../../include/ck/README.md)
# Composable Kernel Grouped Convolution
# 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:
@@ -66,3 +79,52 @@ Table of supported cases by instance factory with fused elementwise operation:
* **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](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
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.

2
client_example/07_grouped_convnd_fwd/common.hpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/07_grouped_convnd_fwd/grouped_conv1d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/07_grouped_convnd_fwd/grouped_conv2d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/07_grouped_convnd_fwd/grouped_conv2d_fwd_ngchw.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <tuple>
#include <cstdlib>

2
client_example/07_grouped_convnd_fwd/grouped_conv3d_fwd_bf8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/07_grouped_convnd_fwd/grouped_conv3d_fwd_bf8_fp8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/07_grouped_convnd_fwd/grouped_conv3d_fwd_fp8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/07_grouped_convnd_fwd/grouped_conv3d_fwd_fp8_bf8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

89
client_example/08_fused_attention/README.md Normal file
View File

@@ -0,0 +1,89 @@
# Fused Attention Examples
This directory contains comprehensive examples demonstrating CK's high-performance fused attention implementations, which are critical for modern transformer architectures and large language models.
---
## Theory
**Fused Multi-Head Attention Operation:**
The fused attention mechanism performs the core transformer operation in a single, optimized kernel:
$$
\text{Attention}(Q, K, V) = \text{Softmax}(Q K^T / \sqrt{d_k}) V
$$
**Detailed Mathematical Steps:**
1. **Query-Key Attention Scores**: $S = Q K^T$
2. **Scale**: $S_{\text{scaled}} = S / \sqrt{d_k}$
3. **Softmax**: $A = \text{Softmax}(S_{\text{scaled}})$
4. **Weighted Value Sum**: $\text{Output} = A V$
- Multi-head extension: Each head computes attention independently, then results are concatenated and projected.
- Tensor shapes: Q, K, V, Output are typically [Batch, Seq_len, Num_heads, Head_dim].
**Algorithmic Background:**
- Fused attention combines two GEMMs and a softmax in a single kernel, minimizing memory traffic.
- Supports bias, masking, and permutation for transformer and LLM workloads.
---
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/08_fused_attention
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (basic fused attention)
./fused_attention
# Example run (fused attention with bias)
./fused_attention_bias
```
---
## Source Code Structure
### Directory Layout
```
client_example/08_fused_attention/
├── fused_attention.cpp # Main client example: fused attention (Q, K, V)
├── fused_attention_bias.cpp # Fused attention with bias
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up Q, K, V tensors, configures attention parameters, launches the fused kernel, and verifies the result.
- **Fused attention kernel invocation**:
Uses the Composable Kernel device API to launch the fused attention operation, optionally with bias.
---
## Additional Details
- Supports FP16, BF16, FP32, and mixed precision.
- Handles causal and generic masking for autoregressive and variable-length models.
- Optimized for memory efficiency (no intermediate attention matrix in global memory).
- Example parameters can be adjusted in the source for different transformer workloads.
---
## Related Examples
- [01_gemm](../01_gemm/README.md): GEMM for Q×K^T and Attn×V
- [06_softmax](../06_softmax/README.md): Softmax client API usage
- [03_gemm_layernorm](../03_gemm_layernorm/README.md): Fused GEMM + layer normalization
- [07_grouped_convnd_fwd](../07_grouped_convnd_fwd/README.md): Grouped convolution for vision transformers
---
[Back to Client Examples](../README.md)

2
client_example/08_fused_attention/fused_attention.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <vector>

2
client_example/08_fused_attention/fused_attention_bias.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iostream>
#include <vector>

85
client_example/09_quantization/README.md Normal file
View File

@@ -0,0 +1,85 @@
# Client Example: Quantization for GEMM and Conv2D
## Theory
This client example demonstrates **quantized GEMM and 2D convolution** operations, including per-layer and per-channel quantization, and fusion with bias and activation functions. Quantization reduces memory and computation by representing values with lower-precision integer types (e.g., int8), enabling efficient inference in deep learning.
**Mathematical Formulation:**
- Quantized GEMM: $C = \text{dequant}(A_q) \times \text{dequant}(B_q)$
- Quantized Conv2D: $Y = \text{dequant}(X_q) * \text{dequant}(W_q)$
- $\text{dequant}(x_q) = (x_q - z) \cdot s$ (scale $s$, zero-point $z$)
- Per-layer: one scale/zero-point per tensor
- Per-channel: scale/zero-point per output channel
**Algorithmic Background:**
- Quantized values are dequantized on-the-fly during computation.
- Accumulation is performed in higher precision for accuracy.
- Supports bias addition and activation fusion (ReLU, Tanh).
- Per-channel quantization improves accuracy for convolutional layers.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/09_quantization
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (GEMM quantization)
./gemm_quantization
# Example run (Conv2D per-layer quantization)
./conv2d_fwd_perlayer_quantization
# Example run (Conv2D per-channel quantization)
./conv2d_fwd_perchannel_quantization
# Example run (Conv2D + bias + ReLU + per-channel quantization)
./conv2d_fwd_bias_relu_perchannel_quantization
```
## Source Code Structure
### Directory Layout
```
client_example/09_quantization/
├── gemm_quantization.cpp # Quantized GEMM
├── conv2d_fwd_perlayer_quantization.cpp # Conv2D per-layer quantization
├── conv2d_fwd_perchannel_quantization.cpp # Conv2D per-channel quantization
├── conv2d_fwd_bias_relu_perlayer_quantization.cpp # Conv2D + bias + ReLU + per-layer quantization
├── conv2d_fwd_bias_relu_perchannel_quantization.cpp # Conv2D + bias + ReLU + per-channel quantization
├── conv2d_fwd_bias_tanh_perlayer_quantization.cpp # Conv2D + bias + Tanh + per-layer quantization
├── conv2d_fwd_bias_tanh_perchannel_quantization.cpp # Conv2D + bias + Tanh + per-channel quantization
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures quantization parameters, launches the quantized kernel, and verifies the result.
- **Quantization kernel invocation**:
Uses the Composable Kernel device API to launch quantized GEMM or Conv2D with optional bias and activation.
---
## Additional Details
- Supports int8 quantization, per-layer and per-channel scaling.
- Demonstrates fusion with bias and activation (ReLU, Tanh).
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [01_gemm](../01_gemm/README.md): GEMM for quantized matrix multiplication
- [14_gemm_quantization](../../example/14_gemm_quantization/README.md): GEMM quantization in the main example directory
- [40_conv2d_fwd_quantization](../../example/40_conv2d_fwd_quantization/README.md): Conv2D quantization in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/09_quantization/conv2d_fwd_bias_relu_perchannel_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/conv2d_fwd_bias_relu_perlayer_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/conv2d_fwd_bias_tanh_perchannel_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/conv2d_fwd_bias_tanh_perlayer_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/conv2d_fwd_perchannel_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/conv2d_fwd_perlayer_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

2
client_example/09_quantization/gemm_quantization.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

55
client_example/10_grouped_convnd_bwd_data/README.md
View File

@@ -1,4 +1,4 @@
[Back to supported operations](../../../include/ck/README.md)
[Back to supported operations](../../include/ck/README.md)
# Composable Kernel Grouped Convolution
## Grouped Convolution Backward Data
@@ -46,3 +46,56 @@ Table of supported cases by instance factory with fused elementwise operation:
* **Bilinear** - 3D, NHWGC, bf16/fp16/fp32
* **Scale** - 3D, NHWGC, bf16/fp16/fp32
---
## Theory
**Grouped convolution backward data** computes the gradient of the input tensor with respect to the loss, given the output gradient and the weights, for each group independently. This is essential for training CNNs and grouped/expert models.
**Mathematical Formulation:**
For each group $g$:
$$
\text{InputGrad}^g = \text{ConvBwdData}(\text{OutputGrad}^g, \text{Weights}^g)
$$
- Supports 1D, 2D, and 3D grouped convolutions.
- Utilizes implicit GEMM for efficient computation.
- Supports fused elementwise operations (e.g., bilinear, scale).
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/10_grouped_convnd_bwd_data
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (2D grouped convolution backward data)
./grouped_conv2d_bwd_data
```
## Source Code Structure
### Directory Layout
```
client_example/10_grouped_convnd_bwd_data/
├── grouped_conv1d_bwd_data.cpp # 1D grouped convolution backward data
├── grouped_conv2d_bwd_data.cpp # 2D grouped convolution backward data
├── grouped_conv3d_bwd_data.cpp # 3D grouped convolution backward data
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input/output tensors, configures grouped convolution parameters, launches the backward data kernel, and verifies the result.
- **Grouped convolution backward kernel invocation**:
Uses the Composable Kernel device API to launch grouped convolution backward data for different dimensions and data types.
This client example provides a comprehensive demonstration of grouped convolution backward data for efficient CNN and vision transformer training.

2
client_example/10_grouped_convnd_bwd_data/grouped_conv2d_bwd_data.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/10_grouped_convnd_bwd_data/grouped_conv2d_bwd_data_ngchw.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/10_grouped_convnd_bwd_data/grouped_conv3d_bwd_data.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/10_grouped_convnd_bwd_data/grouped_conv3d_bwd_data_input_fp16_comp_bf8f8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

62
client_example/11_grouped_conv_bwd_weight/README.md
View File

@@ -1,4 +1,4 @@
[Back to supported operations](../../../include/ck/README.md)
[Back to supported operations](../../include/ck/README.md)
# Composable Kernel Grouped Convolution
## Grouped Convolution Backward Weight
@@ -60,3 +60,63 @@ Table of supported cases by instance factory with fused elementwise operation:
* **Bilinear** - 3D, NHWGC, bf16(fp32 for weight)/fp16/fp32
* **Scale** - 3D, NHWGC, bf16(fp32 for weight)/fp16/fp32
---
## Theory
**Grouped convolution backward weight** computes the gradient of the weights with respect to the loss, given the input and output gradients, for each group independently. This is essential for training CNNs and grouped/expert models.
**Mathematical Formulation:**
For each group $g$:
$$
\text{WeightGrad}^g = \text{ConvBwdWeight}(\text{Input}^g, \text{OutputGrad}^g)
$$
- Supports 1D, 2D, and 3D grouped convolutions.
- Utilizes implicit GEMM for efficient computation.
- Supports fused elementwise operations (e.g., bilinear, scale).
- Uses splitK for large GEMM K dimensions.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/11_grouped_conv_bwd_weight
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (2D grouped convolution backward weight, FP16)
./grouped_conv2d_bwd_weight_fp16
# Example run (3D grouped convolution backward weight, FP32)
./grouped_conv3d_bwd_weight_fp32
```
## Source Code Structure
### Directory Layout
```
client_example/11_grouped_conv_bwd_weight/
├── grouped_conv1d_bwd_weight_fp16.cpp # 1D grouped convolution backward weight (FP16)
├── grouped_conv2d_bwd_weight_fp16.cpp # 2D grouped convolution backward weight (FP16)
├── grouped_conv3d_bwd_weight_fp16.cpp # 3D grouped convolution backward weight (FP16)
├── grouped_conv3d_bwd_weight_fp32.cpp # 3D grouped convolution backward weight (FP32)
├── grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8.cpp # 3D grouped convolution backward weight (FP16, BF8/FP8 mixed)
├── common.hpp # Common utilities for grouped convolution
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input/output tensors, configures grouped convolution parameters, launches the backward weight kernel, and verifies the result.
- **Grouped convolution backward weight kernel invocation**:
Uses the Composable Kernel device API to launch grouped convolution backward weight for different dimensions and data types.
This client example provides a comprehensive demonstration of grouped convolution backward weight for efficient CNN and vision transformer training.

2
client_example/11_grouped_conv_bwd_weight/common.hpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/11_grouped_conv_bwd_weight/grouped_conv1d_bwd_weight_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/11_grouped_conv_bwd_weight/grouped_conv2d_bwd_weight_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/11_grouped_conv_bwd_weight/grouped_conv3d_bwd_weight_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/11_grouped_conv_bwd_weight/grouped_conv3d_bwd_weight_fp16_comp_bf8_fp8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/11_grouped_conv_bwd_weight/grouped_conv3d_bwd_weight_fp32.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

69
client_example/12_elementwise_normalization/README.md Normal file
View File

@@ -0,0 +1,69 @@
# Client Example: Elementwise Layer Normalization
## Theory
This client example demonstrates **elementwise layer normalization** for 2D tensors. Layer normalization is used in transformers and other neural networks to normalize activations across the feature dimension, improving training stability. Elementwise normalization fuses normalization with other elementwise operations for efficiency.
**Mathematical Formulation:**
Given input $X$:
- Mean: $\mu = \frac{1}{N} \sum_{i=1}^N X_i$
- Variance: $\sigma^2 = \frac{1}{N} \sum_{i=1}^N (X_i - \mu)^2$
- Normalized: $\hat{X}_i = \frac{X_i - \mu}{\sqrt{\sigma^2 + \epsilon}}$
- Output: $Y_i = \gamma \hat{X}_i + \beta$
$\gamma$, $\beta$ are learnable scale and shift parameters.
**Algorithmic Background:**
- Computes mean and variance per row (sample).
- Applies normalization and affine transformation.
- Can be fused with other elementwise operations for efficiency.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/12_elementwise_normalization
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run
./elementwise_layernorm2d
```
## Source Code Structure
### Directory Layout
```
client_example/12_elementwise_normalization/
├── elementwise_layernorm2d.cpp # Main client example: elementwise layernorm for 2D tensors
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `elementwise_layernorm2d.cpp`):
Sets up input tensors, configures normalization parameters, launches the normalization kernel, and verifies the result.
- **Elementwise normalization kernel invocation**:
Uses the Composable Kernel device API to launch layer normalization, optionally fused with other elementwise ops.
---
## Additional Details
- Supports fusion with other elementwise operations for efficiency.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [05_layernorm](../05_layernorm/README.md): Layer normalization client API
- [27_layernorm2d_fwd](../../example/27_layernorm2d_fwd/README.md): Layer normalization in the main example directory
---
[Back to Client Examples](../README.md)

352
client_example/12_elementwise_normalization/elementwise_layernorm2d.cpp
View File

@@ -1,176 +1,176 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>
#include <iostream>
#include "ck/ck.hpp"
#include "ck/utility/reduction_enums.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_elementwise_normalization_impl.hpp"
#include "ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/library/tensor_operation_instance/gpu/elementwise_normalization.hpp"
using ADataType = ck::half_t; // Input 1
using BDataType = ck::half_t; // Input 2
using XDataType = ck::half_t;
using GammaDataType = ck::half_t;
using BetaDataType = ck::half_t;
using YDataType = ck::half_t;
using AccDataType = float;
using XElementwiseOperation = ck::tensor_operation::element_wise::Add;
using YElementwiseOperation = ck::tensor_operation::element_wise::PassThrough;
constexpr int Rank = 2;
constexpr int NumReduceDim = 1;
struct SimpleDeviceMem
{
SimpleDeviceMem() = delete;
SimpleDeviceMem(std::size_t mem_size) : p_mem_{}
{
(void)hipMalloc(static_cast<void**>(&p_mem_), mem_size);
}
void* GetDeviceBuffer() { return p_mem_; }
~SimpleDeviceMem() { (void)hipFree(p_mem_); }
void* p_mem_;
};
int main()
{
bool time_kernel = true;
ck::index_t M = 48 * 256;
ck::index_t N = 1024;
ck::index_t Stride = N;
auto mn_size = (M - 1) * Stride + N;
SimpleDeviceMem a_dev_buf(sizeof(ADataType) * mn_size);
SimpleDeviceMem b_dev_buf(sizeof(BDataType) * mn_size);
SimpleDeviceMem gamma_dev_buf(sizeof(GammaDataType) * N);
SimpleDeviceMem beta_dev_buf(sizeof(BetaDataType) * N);
SimpleDeviceMem y_dev_buf(sizeof(YDataType) * mn_size);
std::array<const void*, 2> ab_input = {a_dev_buf.GetDeviceBuffer(),
b_dev_buf.GetDeviceBuffer()};
std::vector<ck::index_t> abStride = {Stride, 1};
std::array<std::vector<ck::index_t>, 2> abStrides = {abStride, abStride};
using DeviceOp = ck::tensor_operation::device::DeviceElementwiseNormalization<
ck::Tuple<ADataType, BDataType>,
GammaDataType,
BetaDataType,
AccDataType,
YDataType,
XElementwiseOperation,
YElementwiseOperation,
Rank,
NumReduceDim>;
// get device op instances
const auto op_ptrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory<
DeviceOp>::GetInstances();
std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
std::string best_op_name;
bool found = false;
int best_op_id = -1;
float best_ave_time = std::numeric_limits<float>::max();
float best_gb_per_sec = 0;
// profile device operation instances
std::cout << "Run all instances and do timing" << std::endl;
for(int i = 0; i < op_ptrs.size(); ++i)
{
auto& op_ptr = op_ptrs[i];
auto argument_ptr = op_ptr->MakeArgumentPointer({M, N}, // lengths
abStrides,
{0, 1}, // gammaStrides
{0, 1}, // betaStrides
{Stride, 1}, // yStrides
{1}, // reduceDims
1e-4,
ab_input,
gamma_dev_buf.GetDeviceBuffer(),
beta_dev_buf.GetDeviceBuffer(),
y_dev_buf.GetDeviceBuffer(),
XElementwiseOperation{},
YElementwiseOperation{});
auto invoker_ptr = op_ptr->MakeInvokerPointer();
std::string op_name = op_ptr->GetTypeString();
if(op_ptr->IsSupportedArgument(argument_ptr.get()))
{
float ave_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, true});
std::size_t num_byte = sizeof(ADataType) * M * N + sizeof(BDataType) * M * N +
sizeof(GammaDataType) * N + sizeof(BetaDataType) * N +
sizeof(YDataType) * M * N;
float gb_per_sec = num_byte / 1.E6 / ave_time;
std::cout << "Perf: " << std::setw(10) << ave_time << " ms, " << gb_per_sec << " GB/s, "
<< op_name << std::endl;
if(ave_time < best_ave_time)
{
found = true;
best_op_id = i;
best_op_name = op_name;
best_ave_time = ave_time;
best_gb_per_sec = gb_per_sec;
}
}
else
{
std::cout << op_name << " does not support this problem" << std::endl;
}
}
std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
<< best_op_name << std::endl;
// run the best intance
if(found)
{
auto& op_ptr = op_ptrs[best_op_id];
std::cout << "Run the best instance without timing: " << op_ptr->GetTypeString()
<< std::endl;
auto argument_ptr = op_ptr->MakeArgumentPointer({M, N}, // lengths
abStrides,
{1}, // gammaStrides
{1}, // betaStrides
{Stride, 1}, // yStrides
{1}, // reduceDims
1e-4,
ab_input,
gamma_dev_buf.GetDeviceBuffer(),
beta_dev_buf.GetDeviceBuffer(),
y_dev_buf.GetDeviceBuffer(),
XElementwiseOperation{},
YElementwiseOperation{});
auto invoker_ptr = op_ptr->MakeInvokerPointer();
if(op_ptr->IsSupportedArgument(argument_ptr.get()))
{
invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false});
}
std::cout << "Done" << std::endl;
}
return 0;
}
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
#include <iomanip>
#include <vector>
#include <iostream>
#include "ck/ck.hpp"
#include "ck/utility/reduction_enums.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_elementwise_normalization_impl.hpp"
#include "ck/tensor_operation/gpu/device/reduction_operator_mapping.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
#include "ck/library/tensor_operation_instance/gpu/elementwise_normalization.hpp"
using ADataType = ck::half_t; // Input 1
using BDataType = ck::half_t; // Input 2
using XDataType = ck::half_t;
using GammaDataType = ck::half_t;
using BetaDataType = ck::half_t;
using YDataType = ck::half_t;
using AccDataType = float;
using XElementwiseOperation = ck::tensor_operation::element_wise::Add;
using YElementwiseOperation = ck::tensor_operation::element_wise::PassThrough;
constexpr int Rank = 2;
constexpr int NumReduceDim = 1;
struct SimpleDeviceMem
{
SimpleDeviceMem() = delete;
SimpleDeviceMem(std::size_t mem_size) : p_mem_{}
{
(void)hipMalloc(static_cast<void**>(&p_mem_), mem_size);
}
void* GetDeviceBuffer() { return p_mem_; }
~SimpleDeviceMem() { (void)hipFree(p_mem_); }
void* p_mem_;
};
int main()
{
bool time_kernel = true;
ck::index_t M = 48 * 256;
ck::index_t N = 1024;
ck::index_t Stride = N;
auto mn_size = (M - 1) * Stride + N;
SimpleDeviceMem a_dev_buf(sizeof(ADataType) * mn_size);
SimpleDeviceMem b_dev_buf(sizeof(BDataType) * mn_size);
SimpleDeviceMem gamma_dev_buf(sizeof(GammaDataType) * N);
SimpleDeviceMem beta_dev_buf(sizeof(BetaDataType) * N);
SimpleDeviceMem y_dev_buf(sizeof(YDataType) * mn_size);
std::array<const void*, 2> ab_input = {a_dev_buf.GetDeviceBuffer(),
b_dev_buf.GetDeviceBuffer()};
std::vector<ck::index_t> abStride = {Stride, 1};
std::array<std::vector<ck::index_t>, 2> abStrides = {abStride, abStride};
using DeviceOp = ck::tensor_operation::device::DeviceElementwiseNormalization<
ck::Tuple<ADataType, BDataType>,
GammaDataType,
BetaDataType,
AccDataType,
YDataType,
XElementwiseOperation,
YElementwiseOperation,
Rank,
NumReduceDim>;
// get device op instances
const auto op_ptrs = ck::tensor_operation::device::instance::DeviceOperationInstanceFactory<
DeviceOp>::GetInstances();
std::cout << "found " << op_ptrs.size() << " instances" << std::endl;
std::string best_op_name;
bool found = false;
int best_op_id = -1;
float best_ave_time = std::numeric_limits<float>::max();
float best_gb_per_sec = 0;
// profile device operation instances
std::cout << "Run all instances and do timing" << std::endl;
for(int i = 0; i < op_ptrs.size(); ++i)
{
auto& op_ptr = op_ptrs[i];
auto argument_ptr = op_ptr->MakeArgumentPointer({M, N}, // lengths
abStrides,
{0, 1}, // gammaStrides
{0, 1}, // betaStrides
{Stride, 1}, // yStrides
{1}, // reduceDims
1e-4,
ab_input,
gamma_dev_buf.GetDeviceBuffer(),
beta_dev_buf.GetDeviceBuffer(),
y_dev_buf.GetDeviceBuffer(),
XElementwiseOperation{},
YElementwiseOperation{});
auto invoker_ptr = op_ptr->MakeInvokerPointer();
std::string op_name = op_ptr->GetTypeString();
if(op_ptr->IsSupportedArgument(argument_ptr.get()))
{
float ave_time = invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, true});
std::size_t num_byte = sizeof(ADataType) * M * N + sizeof(BDataType) * M * N +
sizeof(GammaDataType) * N + sizeof(BetaDataType) * N +
sizeof(YDataType) * M * N;
float gb_per_sec = num_byte / 1.E6 / ave_time;
std::cout << "Perf: " << std::setw(10) << ave_time << " ms, " << gb_per_sec << " GB/s, "
<< op_name << std::endl;
if(ave_time < best_ave_time)
{
found = true;
best_op_id = i;
best_op_name = op_name;
best_ave_time = ave_time;
best_gb_per_sec = gb_per_sec;
}
}
else
{
std::cout << op_name << " does not support this problem" << std::endl;
}
}
std::cout << "Best Perf: " << best_ave_time << " ms, " << best_gb_per_sec << " GB/s, "
<< best_op_name << std::endl;
// run the best intance
if(found)
{
auto& op_ptr = op_ptrs[best_op_id];
std::cout << "Run the best instance without timing: " << op_ptr->GetTypeString()
<< std::endl;
auto argument_ptr = op_ptr->MakeArgumentPointer({M, N}, // lengths
abStrides,
{1}, // gammaStrides
{1}, // betaStrides
{Stride, 1}, // yStrides
{1}, // reduceDims
1e-4,
ab_input,
gamma_dev_buf.GetDeviceBuffer(),
beta_dev_buf.GetDeviceBuffer(),
y_dev_buf.GetDeviceBuffer(),
XElementwiseOperation{},
YElementwiseOperation{});
auto invoker_ptr = op_ptr->MakeInvokerPointer();
if(op_ptr->IsSupportedArgument(argument_ptr.get()))
{
invoker_ptr->Run(argument_ptr.get(), StreamConfig{nullptr, false});
}
std::cout << "Done" << std::endl;
}
return 0;
}

76
client_example/13_batchnorm/README.md Normal file
View File

@@ -0,0 +1,76 @@
# Client Example: Batch Normalization (Forward, Backward, Inference)
## Theory
This client example demonstrates **batch normalization** in forward, backward, and inference modes for NHWC tensors. Batch normalization is used in deep neural networks to normalize activations across the batch and spatial dimensions, improving training stability and convergence.
**Mathematical Formulation:**
Given input $X[N, H, W, C]$:
- Mean: $\mu_c = \frac{1}{NHW} \sum_{n,h,w} X_{n,h,w,c}$
- Variance: $\sigma^2_c = \frac{1}{NHW} \sum_{n,h,w} (X_{n,h,w,c} - \mu_c)^2$
- Normalized: $\hat{X}_{n,h,w,c} = \frac{X_{n,h,w,c} - \mu_c}{\sqrt{\sigma^2_c + \epsilon}}$
- Output: $Y_{n,h,w,c} = \gamma_c \hat{X}_{n,h,w,c} + \beta_c$
$\gamma_c$, $\beta_c$ are learnable scale and shift parameters per channel.
**Algorithmic Background:**
- Forward pass computes mean, variance, normalization, and affine transformation.
- Backward pass computes gradients with respect to input, gamma, and beta.
- Inference uses running mean and variance for normalization.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/13_batchnorm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (forward)
./batchnorm_fwd_nhwc
# Example run (backward)
./batchnorm_bwd_nhwc
# Example run (inference)
./batchnorm_infer_nhwc
```
## Source Code Structure
### Directory Layout
```
client_example/13_batchnorm/
├── batchnorm_fwd_nhwc.cpp # Batchnorm forward (NHWC)
├── batchnorm_bwd_nhwc.cpp # Batchnorm backward (NHWC)
├── batchnorm_infer_nhwc.cpp # Batchnorm inference (NHWC)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures batchnorm parameters, launches the forward, backward, or inference kernel, and verifies the result.
- **BatchNorm kernel invocation**:
Uses the Composable Kernel device API to launch batch normalization for different modes.
---
## Additional Details
- Supports NHWC layout for image and vision models.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [34_batchnorm](../../example/34_batchnorm/README.md): Batch normalization in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/13_batchnorm/batchnorm_bwd_nhwc.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>

2
client_example/13_batchnorm/batchnorm_fwd_nhwc.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>

2
client_example/13_batchnorm/batchnorm_infer_nhwc.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>

63
client_example/14_instance_id/README.md Normal file
View File

@@ -0,0 +1,63 @@
# Client Example: BatchNorm with Instance ID Selection
## Theory
This client example demonstrates **batch normalization** using explicit instance ID selection. In Composable Kernel, "instance ID" refers to a specific kernel configuration (tile sizes, vectorization, etc.) chosen for a given workload. This allows users to benchmark or select the best-performing kernel for their data shape.
**Mathematical Formulation:**
See [BatchNorm Theory](../13_batchnorm/README.md) for the mathematical details of batch normalization.
**Algorithmic Background:**
- The example shows how to enumerate and select a specific kernel instance by its ID.
- Useful for performance tuning, benchmarking, and debugging.
- BatchNorm is performed in NHWC layout.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/14_instance_id
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (selects a specific kernel instance)
./batchnorm_fwd_instance_id
```
## Source Code Structure
### Directory Layout
```
client_example/14_instance_id/
├── batchnorm_fwd_instance_id.cpp # Batchnorm forward with instance ID selection
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `batchnorm_fwd_instance_id.cpp`):
Sets up input tensors, enumerates available kernel instances, selects an instance by ID, launches the batchnorm kernel, and verifies the result.
- **Instance selection**:
Demonstrates how to use the Composable Kernel API to list and select kernel configurations.
---
## Additional Details
- Useful for kernel benchmarking and performance tuning.
- Example parameters and instance ID can be adjusted in the source.
---
## Related Examples
- [13_batchnorm](../13_batchnorm/README.md): Batch normalization client API
- [34_batchnorm](../../example/34_batchnorm/README.md): Batch normalization in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/14_instance_id/batchnorm_fwd_instance_id.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <functional>
#include <numeric>

73
client_example/15_convnd_bwd_data/README.md Normal file
View File

@@ -0,0 +1,73 @@
# Client Example: N-Dimensional Convolution Backward Data
## Theory
This client example demonstrates **N-dimensional convolution backward data** for 3D inputs, supporting multiple data types (FP16, FP32). The backward data operation computes the gradient of the input tensor with respect to the loss, given the output gradient and the weights. This is essential for training CNNs and 3D vision models.
**Mathematical Formulation:**
For input $X$, weights $W$, and output gradient $dY$:
$$
dX = \text{ConvBwdData}(dY, W)
$$
- Supports 3D convolution (ND can be extended).
- Utilizes implicit GEMM for efficient computation.
**Algorithmic Background:**
- The backward data operation is implemented as a convolution with transformed coordinates.
- Used in training pipelines for 3D CNNs, medical imaging, and volumetric data.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/15_convnd_bwd_data
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (3D backward data, FP16)
./conv3d_bwd_data_fp16
# Example run (3D backward data, FP32)
./conv3d_bwd_data_fp32
```
## Source Code Structure
### Directory Layout
```
client_example/15_convnd_bwd_data/
├── conv3d_bwd_data_fp16.cpp # 3D convolution backward data (FP16)
├── conv3d_bwd_data_fp32.cpp # 3D convolution backward data (FP32)
├── common.hpp # Common utilities for convolution
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input/output tensors, configures convolution parameters, launches the backward data kernel, and verifies the result.
- **Backward data kernel invocation**:
Uses the Composable Kernel device API to launch convolution backward data for different data types.
---
## Additional Details
- Supports FP16 and FP32 for 3D convolution.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [10_grouped_convnd_bwd_data](../10_grouped_convnd_bwd_data/README.md): Grouped convolution backward data
- [17_convnd_bwd_data](../../example/17_convnd_bwd_data/README.md): Convolution backward data in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/15_convnd_bwd_data/common.hpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/15_convnd_bwd_data/conv3d_bwd_data_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/15_convnd_bwd_data/conv3d_bwd_data_fp32.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

85
client_example/16_convnd_fwd/README.md Normal file
View File

@@ -0,0 +1,85 @@
# Client Example: N-Dimensional Convolution Forward
## Theory
This client example demonstrates **N-dimensional convolution forward** for 3D inputs, supporting multiple data types (FP16, FP32, FP8 composite). Convolution is a fundamental operation in deep learning, especially in convolutional neural networks (CNNs) for images, audio, and volumetric data.
**Mathematical Formulation:**
Given input $X$, weights $W$:
$$
Y = \text{Conv}(X, W)
$$
- Supports 3D convolution (ND can be extended).
- Utilizes implicit GEMM for efficient computation.
**Algorithmic Background:**
- The forward convolution operation is implemented as a convolution with transformed coordinates.
- Used in inference and training pipelines for 3D CNNs, medical imaging, and volumetric data.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/16_convnd_fwd
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (3D forward, FP16)
./conv3d_fwd_fp16
# Example run (3D forward, FP32)
./conv3d_fwd_fp32
# Example run (3D forward, FP16 compute with FP8)
./conv3d_fwd_fp16_comp_fp8
```
## Source Code Structure
### Directory Layout
```
client_example/16_convnd_fwd/
├── conv3d_fwd_fp16.cpp # 3D convolution forward (FP16)
├── conv3d_fwd_fp32.cpp # 3D convolution forward (FP32)
├── conv3d_fwd_fp16_comp_fp8.cpp # 3D convolution forward (FP16 compute, FP8)
├── common.hpp # Common utilities for convolution
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input/output tensors, configures convolution parameters, launches the forward kernel, and verifies the result.
- **Forward convolution kernel invocation**:
Uses the Composable Kernel device API to launch convolution forward for different data types.
---
## Additional Details
- Supports FP16, FP32, and FP8 composite for 3D convolution.
- Parameters can be adjusted in the source files for different workloads. The following parameters are configurable:
- `NumDimSpatial`: Number of spatial dimensions (default: 3 for 3D convolution)
- `G`: Number of groups (default: 1)
- `N`: Batch size (default: 64)
- `K`: Number of output channels (default: 128)
- `C`: Number of input channels (default: 64)
- `Z`, `Y`, `X`: Filter/kernel dimensions (default: 3x3x3)
- `Di`, `Hi`, `Wi`: Input dimensions - depth, height, width (default: 28x28x3)
- `Do`, `Ho`, `Wo`: Output dimensions - depth, height, width (default: 28x28x3)
---
## Related Examples
- [09_convnd_fwd](../../example/09_convnd_fwd/README.md): N-dimensional convolution in the main example directory
- [30_grouped_conv_fwd_multiple_d](../../example/30_grouped_conv_fwd_multiple_d/README.md): Grouped convolution forward with multiple D
---
[Back to Client Examples](../README.md)

2
client_example/16_convnd_fwd/common.hpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <cstdlib>
#include <iomanip>

2
client_example/16_convnd_fwd/conv3d_fwd_fp16.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/16_convnd_fwd/conv3d_fwd_fp16_comp_fp8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

2
client_example/16_convnd_fwd/conv3d_fwd_fp32.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"

71
client_example/17_grouped_gemm_fastgelu/README.md Normal file
View File

@@ -0,0 +1,71 @@
# Client Example: Grouped GEMM with FastGELU Activation
## Theory
This client example demonstrates **grouped GEMM fused with FastGELU activation**. Grouped GEMM performs multiple independent GEMM operations (with potentially different shapes) in a single kernel launch, and FastGELU is a fast approximation of the GELU activation used in transformers and MLPs.
**Mathematical Formulation:**
For $G$ groups, each with its own $A_g$, $B_g$:
- GEMM: $Y_g = A_g \times B_g$
- FastGELU: $E_g = \text{FastGELU}(Y_g)$
FastGELU is defined as:
$$
\text{FastGELU}(x) = x \cdot \sigma(1.702 \cdot x)
$$
where $\sigma$ is the sigmoid function.
**Algorithmic Background:**
- Each group can have different matrix sizes and strides.
- The kernel launches a grid covering all groups, with each block assigned to a group.
- FastGELU is applied in the epilogue for each group.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/17_grouped_gemm_fastgelu
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run
./grouped_gemm_fastgelu
```
## Source Code Structure
### Directory Layout
```
client_example/17_grouped_gemm_fastgelu/
├── grouped_gemm_fastgelu.cpp # Main client example: grouped GEMM + FastGELU
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `grouped_gemm_fastgelu.cpp`):
Sets up input matrices for each group, configures GEMM and epilogue parameters, launches the grouped kernel, and verifies the result.
- **Grouped GEMM kernel invocation**:
Uses the Composable Kernel device API to launch grouped GEMM with FastGELU activation.
---
## Additional Details
- Supports multiple groups with different matrix shapes.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [15_grouped_gemm](../../example/15_grouped_gemm/README.md): Grouped GEMM in the main example directory
- [04_gemm_add_add_fastgelu](../../example/04_gemm_add_add_fastgelu/README.md): GEMM with FastGELU fusion
---
[Back to Client Examples](../README.md)

2
client_example/17_grouped_gemm_fastgelu/grouped_gemm_fastgelu.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <iostream>

80
client_example/18_groupnorm/README.md Normal file
View File

@@ -0,0 +1,80 @@
# Client Example: Group Normalization (Forward and Backward)
## Theory
This client example demonstrates **group normalization** in both forward and backward modes, including fusion with Swish activation. Group normalization normalizes activations across groups of channels, improving training stability for small batch sizes or non-i.i.d. data.
**Mathematical Formulation:**
Given input $X[N, C, ...]$ divided into $G$ groups:
- For each group $g$:
- Mean: $\mu_g = \frac{1}{|g|} \sum_{i \in g} X_i$
- Variance: $\sigma^2_g = \frac{1}{|g|} \sum_{i \in g} (X_i - \mu_g)^2$
- Normalized: $\hat{X}_i = \frac{X_i - \mu_g}{\sqrt{\sigma^2_g + \epsilon}}$
- Output: $Y_i = \gamma \hat{X}_i + \beta$
$\gamma$, $\beta$ are learnable scale and shift parameters.
- Swish activation: $\text{Swish}(x) = x \cdot \sigma(x)$, where $\sigma$ is the sigmoid function.
**Algorithmic Background:**
- Forward pass computes mean, variance, normalization, and affine transformation per group.
- Backward pass computes gradients with respect to input, gamma, and beta.
- Swish activation can be fused with normalization for efficiency.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/18_groupnorm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (forward with Swish)
./groupnorm_swish_fwd
# Example run (backward, data)
./groupnorm_bwd_data
# Example run (backward, gamma/beta)
./groupnorm_bwd_gamma_beta
```
## Source Code Structure
### Directory Layout
```
client_example/18_groupnorm/
├── groupnorm_swish_fwd.cpp # Groupnorm forward with Swish activation
├── groupnorm_bwd_data.cpp # Groupnorm backward (data)
├── groupnorm_bwd_gamma_beta.cpp # Groupnorm backward (gamma/beta)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures groupnorm parameters, launches the forward or backward kernel, and verifies the result.
- **GroupNorm kernel invocation**:
Uses the Composable Kernel device API to launch group normalization for different modes.
---
## Additional Details
- Supports fusion with Swish activation for efficiency.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [42_groupnorm_fwd](../../example/42_groupnorm_fwd/README.md): Group normalization in the main example directory
- [54_groupnorm_bwd](../../example/54_groupnorm_bwd/README.md): Group normalization backward in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/18_groupnorm/groupnorm_bwd_data.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/18_groupnorm/groupnorm_bwd_gamma_beta.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/18_groupnorm/groupnorm_swish_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

80
client_example/19_pool/README.md Normal file
View File

@@ -0,0 +1,80 @@
# Client Example: Pooling Operations (2D Max, 3D Avg)
## Theory
This client example demonstrates **pooling operations** for 2D max pooling and 3D average pooling, including both forward and backward passes. Pooling is used in convolutional neural networks (CNNs) for spatial downsampling, translation invariance, and reducing computation.
**Mathematical Formulation:**
- **Max Pooling (2D):** $Y_{n,c,h,w} = \max_{i,j} X_{n,c,h \cdot s_H + i, w \cdot s_W + j}$
- **Average Pooling (3D):** $Y_{n,c,d,h,w} = \frac{1}{k_D k_H k_W} \sum_{i,j,k} X_{n,c,d \cdot s_D + i, h \cdot s_H + j, w \cdot s_W + k}$
Where $s_H, s_W, s_D$ are strides, $k_H, k_W, k_D$ are kernel sizes.
**Algorithmic Background:**
- Forward pass computes the pooled output.
- Backward pass computes the gradient with respect to the input.
- Handles padding and boundary conditions.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/19_pool
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (2D max pool forward)
./max_pool2d_fwd
# Example run (2D max pool backward)
./max_pool2d_bwd
# Example run (3D avg pool forward)
./avg_pool3d_fwd
# Example run (3D avg pool backward)
./avg_pool3d_bwd
```
## Source Code Structure
### Directory Layout
```
client_example/19_pool/
├── max_pool2d_fwd.cpp # 2D max pooling forward
├── max_pool2d_bwd.cpp # 2D max pooling backward
├── avg_pool3d_fwd.cpp # 3D average pooling forward
├── avg_pool3d_bwd.cpp # 3D average pooling backward
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in each `.cpp`):
Sets up input tensors, configures pooling parameters, launches the forward or backward kernel, and verifies the result.
- **Pooling kernel invocation**:
Uses the Composable Kernel device API to launch pooling operations for different modes.
---
## Additional Details
- Supports both max and average pooling, forward and backward.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [13_pool2d_fwd](../../example/13_pool2d_fwd/README.md): 2D pooling in the main example directory
- [48_pool3d_fwd](../../example/48_pool3d_fwd/README.md): 3D pooling in the main example directory
- [49_maxpool2d_bwd](../../example/49_maxpool2d_bwd/README.md): 2D max pool backward in the main example directory
- [51_avgpool3d_bwd](../../example/51_avgpool3d_bwd/README.md): 3D avg pool backward in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/19_pool/avg_pool3d_bwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/19_pool/avg_pool3d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/19_pool/max_pool2d_bwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

2
client_example/19_pool/max_pool2d_fwd.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

66
client_example/20_splitk_gemm/README.md Normal file
View File

@@ -0,0 +1,66 @@
# Client Example: Split-K GEMM
## Theory
This client example demonstrates **Split-K GEMM**, a technique for parallelizing matrix multiplication along the K dimension. Split-K is used to improve parallelism and memory bandwidth utilization for large GEMM operations, especially when K is large.
**Mathematical Formulation:**
- Standard GEMM: $C = A \times B$
- Split-K: Partition the K dimension into $K_s$ splits, compute partial results, then reduce:
$$
C = \sum_{s=1}^{K_s} (A_{[:, K_s]} \times B_{[K_s, :]})
$$
**Algorithmic Background:**
- Each split computes a partial GEMM over a chunk of K.
- Partial results are reduced (summed) to produce the final output.
- Useful for large K, limited workspace, or maximizing GPU occupancy.
## How to Run
### Prerequisites
Please follow the instructions in the main [Build Guide](../../README.md#building-ck) section as a prerequisite to building and running this example.
### Build and run
```bash
cd composable_kernel/client_example/20_splitk_gemm
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j
# Example run (FP16 compute, FP8 output)
./splitK_gemm_fp16_f8
```
## Source Code Structure
### Directory Layout
```
client_example/20_splitk_gemm/
├── splitK_gemm_fp16_f8.cpp # Main client example: Split-K GEMM (FP16 compute, FP8 output)
├── CMakeLists.txt # Build configuration for the example
```
### Key Functions
- **main()** (in `splitK_gemm_fp16_f8.cpp`):
Sets up input matrices, configures Split-K parameters, launches the Split-K GEMM kernel, and verifies the result.
- **Split-K kernel invocation**:
Uses the Composable Kernel device API to launch the Split-K GEMM operation.
---
## Additional Details
- Supports FP16 compute with FP8 output for memory efficiency.
- Example parameters can be adjusted in the source for different workloads.
---
## Related Examples
- [35_splitK_gemm](../../example/35_splitK_gemm/README.md): Split-K GEMM in the main example directory
---
[Back to Client Examples](../README.md)

2
client_example/20_splitk_gemm/splitK_gemm_fp16_f8.cpp
View File

@@ -1,5 +1,5 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <iomanip>
#include <vector>

Some files were not shown because too many files have changed in this diff Show More