wafer-ai/composable_kernel

wafer-ai/composable_kernel

Description: Composable Kernel: Performance Portable Programming Model for Machine Learning Tensor Operators

License: NOASSERTION

Stars: 0

Forks: 0

Open issues: 0

Created: 2026-01-22T05:04:34Z

Pushed: 2026-01-22T05:04:59Z

Default branch: develop

Fork: yes

Parent repository: ROCm/composable_kernel

Archived: no

README:

Composable Kernel

> [!NOTE] > The published documentation is available at Composable Kernel in an organized, easy-to-read format, with search and a table of contents. The documentation source files reside in the docs folder of this repository. As with all ROCm projects, the documentation is open source. For more information on contributing to the documentation, see Contribute to ROCm documentation.

The Composable Kernel (CK) library provides a programming model for writing performance-critical kernels for machine learning workloads across multiple architectures (GPUs, CPUs, etc.). The CK library uses general purpose kernel languages, such as HIP C++.

CK uses two concepts to achieve performance portability and code maintainability:

• A tile-based programming model
• Algorithm complexity reduction for complex machine learning (ML) operators. This uses an innovative

technique called *Tensor Coordinate Transformation*.


The current CK library is structured into four layers:

• Templated Tile Operators
• Templated Kernel and Invoker
• Instantiated Kernel and Invoker
• Client API


General information

• [CK supported operations](include/ck/README.md)
• [CK Tile supported operations](include/ck_tile/README.md)
• [CK wrapper](client_example/25_wrapper/README.md)
• [CK codegen](codegen/README.md)
• [CK profiler](profiler/README.md)
• [Examples (Custom use of CK supported operations)](example/README.md)
• [Client examples (Use of CK supported operations with instance factory)](client_example/README.md)
• [Terminology](/TERMINOLOGY.md)
• [Contributors](/CONTRIBUTORS.md)

CK is released under the [MIT license](/LICENSE).

Building CK

We recommend building CK inside Docker containers, which include all necessary packages. Pre-built Docker images are available on DockerHub.

1. To build a new Docker image, use the Dockerfile provided with the source code:

DOCKER_BUILDKIT=1 docker build -t ck:latest -f Dockerfile .

2. Launch the Docker container:

docker run \
-it \
--privileged \
--group-add sudo \
-w /root/workspace \
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace \
ck:latest \
/bin/bash

3. Clone CK source code from the GitHub repository and start the build:

git clone https://github.com/ROCm/composable_kernel.git && \
cd composable_kernel && \
mkdir build && \
cd build

You must set the GPU_TARGETS macro to specify the GPU target architecture(s) you want to run CK on. You can specify single or multiple architectures. If you specify multiple architectures, use a semicolon between each; for example, gfx908;gfx90a;gfx942.

cmake \
-D CMAKE_PREFIX_PATH=/opt/rocm \
-D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc \
-D CMAKE_BUILD_TYPE=Release \
-D GPU_TARGETS="gfx908;gfx90a" \
..

If you don't set GPU_TARGETS on the cmake command line, CK is built for all GPU targets supported by the current compiler (this may take a long time). Tests and examples will only get built if the GPU_TARGETS is set by the user on the cmake command line.

NOTE: If you try setting GPU_TARGETS to a list of architectures, the build will only work if the architectures are similar, e.g., gfx908;gfx90a, or gfx1100;gfx1101;gfx11012. Otherwise, if you 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.

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:

../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:

../script/cmake-ck-dev.sh .. gfx1100

Or pass additional cmake arguments:

../script/cmake-ck-dev.sh .. gfx90a -DCMAKE_BUILD_TYPE=Release

5. Build the entire CK library:

make -j"$(nproc)"

6. Install CK:

make -j install

[See Note on -j](#notes)

Building for Windows

Install TheRock and run CMake configure as

cmake \
-D CMAKE_PREFIX_PATH="C:/dist/TheRock" \
-D CMAKE_CXX_COMPILER="C:/dist/TheRock/bin/hipcc.exe" \
-D CMAKE_BUILD_TYPE=Release \
-D GPU_TARGETS="gfx1151" \
-G Ninja \
..

Use Ninja to build either the whole library or individual targets.

Optional post-install steps

• Build examples and tests:

make -j examples tests

• Build and run all examples and tests:

make -j check

You can find instructions for running each individual example in [example](/example).

• Build and run smoke/regression examples and tests:

make -j smoke # tests and examples that run for = 30 seconds each

• Build ckProfiler:

make -j ckProfiler

You can find instructions for running ckProfiler in [profiler](/profiler).

• Build our documentation locally:

cd docs
pip3 install -r sphinx/requirements.txt
python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html

Notes

The -j option for building with multiple threads in parallel, which speeds up the build significantly. However, -j launches unlimited number of threads, which can cause the build to run out of memory and crash. On average, you should expect each thread to use ~2Gb of RAM. Depending on the number of CPU cores and the amount of RAM on your system, you may want to limit…