Commit d9175d22 authored by Lisa Delaney's avatar Lisa Delaney
Browse files

readme updates

parent 4fcb9f20
# Composable Kernel
Composable Kernel (CK) library aims to provide a programming model for writing performance critical kernels for machine learning workloads across multiple architectures including GPUs, CPUs, etc, through general purpose kernel languages, like HIP C++.
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:
CK utilizes two concepts to achieve performance portability and code maintainability:
* A tile-based programming model
* Algorithm complexity reduction for complex ML operators, using innovative technique we call "Tensor Coordinate Transformation".
* Algorithm complexity reduction for complex machine learning (ML) operators. This uses an innovative
technique called *Tensor Coordinate Transformation*.
![ALT](/docs/data/ck_component.png "CK Components")
## Code Structure
The current CK library is structured into four layers:
Current CK library are structured into 4 layers:
* "Templated Tile Operators" layer
* "Templated Kernel and Invoker" layer
* "Instantiated Kernel and Invoker" layer
* "Client API" layer
* Templated Tile Operators
* Templated Kernel and Invoker
* Instantiated Kernel and Invoker
* Client API
![ALT](/docs/data/ck_layer.png "CK Layers")
## Documentation
## General information
Run the steps below to build documentation locally.
To build our documentation locally, use the following code:
```
``` bash
cd docs
pip3 install -r sphinx/requirements.txt
python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
```
## Contributors
You can find a list of our developers and contributors on our [Contributors](/CONTRIBUTORS.md) page.
page.
The list of developers and contributors is here: [Contributors](/CONTRIBUTORS.md)
```note
If you use CK, cite us as follows:
## Citation
If you use CK, please use following citations:
* CK paper will be freely available on arXiv soon: [Realizing Tensor Operators Using Coordinate Transformations and Tile Based Programming](???)
* [Realizing Tensor Operators Using Coordinate Transformations and Tile Based Programming](???):
This paper will be available on arXiv soon.
* [CITATION.cff](/CITATION.cff)
```
## License
CK is released under the **[MIT license](/LICENSE)**.
CK is released under the MIT license. [License File](/LICENSE)
## Building CK
We recommend building CK inside Docker containers, which include all necessary packages. Pre-built
Docker images are available on [DockerHub](https://hub.docker.com/r/rocm/composable_kernel/tags).
# Build Composable Kernel
1. To build a new Docker image, use the Dockerfile provided with the source code:
We recommend building Composable Kernel inside docker containers that include
all necessary packages. Pre-built docker images are available from this public repo:
```bash
DOCKER_BUILDKIT=1 docker build -t ck:latest -f Dockerfile .
```
https://hub.docker.com/r/rocm/composable_kernel/tags
2. Launch the Docker container:
In order to build a new docker image, you can use the Dockerfile provided with the source code as shown below:
```bash
docker run \
-it \
--privileged \
--group-add sudo \
-w /root/workspace \
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace \
ck:latest \
/bin/bash
```
```bash
DOCKER_BUILDKIT=1 docker build -t ck:latest -f Dockerfile .
```
3. Clone CK source code from the GitHub repository and start the build:
The docker container can then be launched, e.g., using the following command:
```bash
git clone https://github.com/ROCmSoftwarePlatform/composable_kernel.git && \
cd composable_kernel && \
mkdir build && \
cd build
```
```bash
docker run \
-it \
--privileged \
--group-add sudo \
-w /root/workspace \
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace \
ck:latest \
/bin/bash
```
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;gfx940`.
After launching the container you can clone Composable Kernel source code from the github repository and start the build:
```bash
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" \
..
```
```bash
git clone https://github.com/ROCmSoftwarePlatform/composable_kernel.git && \
cd composable_kernel && \
mkdir build && \
cd build
```
You will need to set the GPU_TARGETS macro to specify GPU target architecture(s) that you want
to execute CK on, e.g., gfx908, or gfx908;gfx90a;gfx940.
You can specify either single or multiple architectures (use semicolon to separate), e.g.:
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).
```bash
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 GPU_TARGETS is not set on the cmake command line, CK will be built for all GPU targets supported by the
current compiler (this may take quite a long time).
4. Build the entire CK library:
After that you can build the entire CK library with just
```bash
make -j
```
```bash
make -j
```
5. Install CK:
## Install CK
```bash
make -j install
```
```bash
make -j install
```
## Optional post-install steps
## Build examples and tests
* Build examples and tests:
```bash
make -j examples tests
```
```bash
make -j examples tests
```
## Build and run all examples and tests
* Build and run all examples and tests:
```bash
make -j check
```
```bash
make -j check
```
Instructions for running each individual examples are under [example](/example)
You can find instructions for running each individual example in [example](/example).
* Build ckProfiler:
## Build ckProfiler
```bash
make -j ckProfiler
```
```bash
make -j ckProfiler
```
Instructions for running ckProfiler are under [profiler](/profiler)
You can find instructions for running ckProfiler in [profiler](/profiler).
Please note the "-j" option for building with multiple threads in parallel. This speeds up the build significantly.
Depending on the number of CPU cores and the amount of RAM on your system, it may be advizable to limit the number of threads.
By default, "-j" will try to launch one thread per CPU core. This could potentially cause the build to run out of memory and crash,
for example if you have a 128-core CPU and 64Gb of RAM. In such cases, you can try to reduce the number of threads to 32 by using "-j32".
Note the `-j` option for building with multiple threads in parallel. This speeds up the build significantly.
Depending on the number of CPU cores and the amount of RAM on your system, you may want to
limit the number of threads. For example, if you have a 128-core CPU and 64 Gb of RAM.
By default, `-j` launches one thread per CPU core, which can cause the build to run out of memory and
crash. In such cases, you can reduce the number of threads to 32 by using `-j32`.
Additional cmake flags can be used to significantly speed-up the build:
INSTANCES_ONLY (by default is OFF) must be set to ON in order to build only the instances and library
while skipping all tests, examples, and profiler. This is useful in cases when you plan to use CK as a dependency and don't plan to
run any examples or tests.
* `INSTANCES_ONLY` (default is OFF) must be set to ON in order to build only the instances and library
while skipping all tests, examples, and profiler. This is useful in cases when you plan to use CK as a
dependency and don't plan to run any examples or tests.
DTYPES (by default not set) can be set to any subset of "fp64;fp32;fp16;fp8;bf16;int8" to build instances
of select data types only. The main default data types are fp32 and fp16, other data types can be safely skipped.
* `DTYPES` (default is not set) can be set to any subset of "fp64;fp32;fp16;fp8;bf16;int8" to build
instances of select data types only. The main default data types are fp32 and fp16; you can safely skip
other data types.
DL_KERNELS (by default is OFF) must be set to ON in order to build instances, such as, gemm_dl or batched_gemm_multi_d_dl.
Those instances are mostly useful on architectures such the NAVI2x, since of most other platforms miuch faster "xdl" or "wmma"
instances are available.
* `DL_KERNELS` (default is OFF) must be set to ON in order to build instances, such as `gemm_dl` or
`batched_gemm_multi_d_dl`. These instances are useful on architectures like the NAVI2x, as most
other platforms have faster instances, such as `xdl` or `wmma`, available.
## Using sccache for building
The default CK docker images come with pre-installed version of sccache which supports clang being used as hip-compiler
" -x hip". Using sccache can help reduce the time to re-build the code from hours to 1 - 2 minutes. In order to
invoke sccache, you need to run
The default CK Docker images come with a pre-installed version of sccache, which supports clang
being used as hip-compiler (" -x hip"). Using sccache can help reduce the time to re-build code from
hours to 1-2 minutes. In order to invoke sccache, you need to run:
```bash
sccache --start-server
```
and add the following flags to the cmake command line:
then add the following flags to the cmake command line:
```bash
-DCMAKE_CXX_COMPILER_LAUNCHER=sccache -DCMAKE_C_COMPILER_LAUNCHER=sccache
```
You may need to clean up the build folder and repeat the cmake and make steps in order to take advantage of the sccache
during the subsequent builds.
You may need to clean up the build folder and repeat the cmake and make steps in order to take
advantage of the sccache during subsequent builds.
## Using CK as pre-built kernel library
Instructions for using CK as a pre-built kernel library are under [client_example](/client_example)
You can find instructions for using CK as a pre-built kernel library in [client_example](/client_example).
## Contributing
## Contributing to CK
When you contribute to Composable Kernel, make sure to run `clang-format` on all the changed files. We highly recommend using git hooks that are managed by the `pre-commit` framework. To install hooks, run:
When you contribute to CK, make sure you run `clang-format` on all changed files. We highly
recommend using git hooks that are managed by the `pre-commit` framework. To install hooks, run:
```bash
sudo script/install_precommit.sh
```
This way, `pre-commit` will add the appropriate hooks to your local repository and automatically run `clang-format` (and possibly additional checks) before any commit is created.
With this approach, `pre-commit` adds the appropriate hooks to your local repository and
automatically runs `clang-format` (and possibly additional checks) before any commit is created.
If you need to uninstall hooks from the repository, you can do so by running the following command:
......@@ -181,14 +191,14 @@ If you need to uninstall hooks from the repository, you can do so by running the
script/uninstall_precommit.sh
```
If for any reason, you need to temporarily disable precommit hooks, you can add the `--no-verify` option to the `git commit` command.
If you need to temporarily disable pre-commit hooks, you can add the `--no-verify` option to the
`git commit` command.
## Caveat
### Kernel Timing and Verification
CK's own kernel timer will warn up kernel once, and then run it multiple times
to get average kernel time. For some kernels that use atomic add, this will cause
output buffer to be accumulated multiple times, causing verification failure.
To work around it, do not use CK's own timer and do verification at the same time.
**Kernel Timing and Verification**: CK's own kernel timer will warn up-kernel once, and then run it
multiple times to get average kernel time. For some kernels that use atomic add, this causes the
output buffer to be accumulated multiple times, causing verification failure. To work around this, don't
use CK's own timer and do verification at the same time.
CK's own timer and verification in each example and ckProfiler can be enabled or
disabled from command line.
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment