README.md 7.74 KB
Newer Older
1
2
# Composable Kernel

3
4
5
> [!NOTE]
> The published documentation is available at [Composable Kernel](https://rocm.docs.amd.com/projects/composable_kernel/en/latest/) 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](https://rocm.docs.amd.com/en/latest/contribute/contributing.html).

6
7
8
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++.
Sam Wu's avatar
Sam Wu committed
9

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

* A tile-based programming model
13
14
* Algorithm complexity reduction for complex machine learning (ML) operators. This uses an innovative
   technique called *Tensor Coordinate Transformation*.
15

Sam Wu's avatar
Sam Wu committed
16
![ALT](/docs/data/ck_component.png "CK Components")
17

18
The current CK library is structured into four layers:
Sam Wu's avatar
Sam Wu committed
19

20
21
22
23
* Templated Tile Operators
* Templated Kernel and Invoker
* Instantiated Kernel and Invoker
* Client API
24

Sam Wu's avatar
Sam Wu committed
25
26
![ALT](/docs/data/ck_layer.png "CK Layers")

27
## General information
Sam Wu's avatar
Sam Wu committed
28

29
To build our documentation locally, use the following code:
Sam Wu's avatar
Sam Wu committed
30

31
``` bash
Sam Wu's avatar
Sam Wu committed
32
cd docs
Sam Wu's avatar
Sam Wu committed
33
pip3 install -r sphinx/requirements.txt
Sam Wu's avatar
Sam Wu committed
34
35
python3 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html
```
36

37
You can find a list of our developers and contributors on our [Contributors](/CONTRIBUTORS.md) page.
38

39
40
```note
If you use CK, cite us as follows:
Sam Wu's avatar
Sam Wu committed
41

42
43
* [Realizing Tensor Operators Using Coordinate Transformations and Tile Based Programming](???):
  This paper will be available on arXiv soon.
44
* [CITATION.cff](/CITATION.cff)
45
```
46

47
CK is released under the **[MIT license](/LICENSE)**.
Sam Wu's avatar
Sam Wu committed
48

49
## Building CK
50

51
52
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).
53

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

56
57
58
    ```bash
    DOCKER_BUILDKIT=1 docker build -t ck:latest -f Dockerfile .
    ```
Sam Wu's avatar
Sam Wu committed
59

60
2. Launch the Docker container:
61

62
63
64
65
66
67
68
69
70
71
    ```bash
    docker run                                     \
    -it                                            \
    --privileged                                   \
    --group-add sudo                               \
    -w /root/workspace                             \
    -v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace  \
    ck:latest                                      \
    /bin/bash
    ```
Sam Wu's avatar
Sam Wu committed
72

73
3. Clone CK source code from the GitHub repository and start the build:
Chao Liu's avatar
Chao Liu committed
74

75
    ```bash
76
    git clone https://github.com/ROCm/composable_kernel.git && \
77
78
79
80
    cd composable_kernel && \
    mkdir build && \
    cd build
    ```
Sam Wu's avatar
Sam Wu committed
81

82
83
84
    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`.
85

86
87
88
89
90
91
92
93
    ```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"                                                                    \
    ..
    ```
Sam Wu's avatar
Sam Wu committed
94

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

    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`.
103

104
4. Build the entire CK library:
105

106
107
108
    ```bash
    make -j
    ```
109

110
5. Install CK:
111

112
113
114
    ```bash
    make -j install
    ```
115

116
## Optional post-install steps
117

118
* Build examples and tests:
119

120
121
122
    ```bash
    make -j examples tests
    ```
Sam Wu's avatar
Sam Wu committed
123

124
125
126
127
128
* Build and run all examples and tests:

    ```bash
    make -j check
    ```
129

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

132
* Build ckProfiler:
133

134
135
136
137
138
139
    ```bash
    make -j ckProfiler
    ```

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

Illia Silin's avatar
Illia Silin committed
140
141
142
Note 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.
143
Depending on the number of CPU cores and the amount of RAM on your system, you may want to
Illia Silin's avatar
Illia Silin committed
144
limit the number of threads. For example, if you have a 128-core CPU and 128 Gb of RAM it's advisable to use `-j32`.
145
146
147
148
149
150
151
152
153
154
155

Additional cmake flags can be used to significantly speed-up the build:

* `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` (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.

Illia Silin's avatar
Illia Silin committed
156
157
158
159
160
* `CK_USE_FP8_ON_UNSUPPORTED_ARCH` (default is OFF) must be set to ON in order to build instances,
  such as `gemm_universal` and `gemm_multiply_multiply` for fp8 data type for GPU targets which do not
  have native support for fp8 data type, such as gfx908 or gfx90a. These instances are useful on
  architectures like the MI100/MI200 for the functional support only.

161
162
163
164
165
## Using sccache for building

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:
Sam Wu's avatar
Sam Wu committed
166

167
```bash
168
 sccache --start-server
169
```
JD's avatar
JD committed
170

171
then add the following flags to the cmake command line:
Sam Wu's avatar
Sam Wu committed
172

173
```bash
174
 -DCMAKE_CXX_COMPILER_LAUNCHER=sccache -DCMAKE_C_COMPILER_LAUNCHER=sccache
175
176
```

177
178
179
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.

180
## Using CK as pre-built kernel library
Sam Wu's avatar
Sam Wu committed
181

182
You can find instructions for using CK as a pre-built kernel library in [client_example](/client_example).
JD's avatar
JD committed
183

184
## Contributing to CK
185

186
187
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:
188
189
190
191
192

```bash
sudo script/install_precommit.sh
```

193
194
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.
195
196
197
198
199
200
201

If you need to uninstall hooks from the repository, you can do so by running the following command:

```bash
script/uninstall_precommit.sh
```

202
203
If you need to temporarily disable pre-commit hooks, you can add the `--no-verify` option to the
`git commit` command.