"composable_kernel/include/utility/Sequence.hpp" did not exist on "0a2657312ec62a65e92a36cebd7d3b2a3c0712e1"
README.md 3.25 KB
Newer Older
Chao Liu's avatar
Chao Liu committed
1
2
3
4
5
6
7
8
9
# Composable Kernel

## Methodology
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 progarmming languages, like HIP C++.

CK utilizes two concepts to achieve performance portabilatity and code maintainbility:
* A tile-based programming model
* Algorithm complexity reduction for complex ML operators, using innovative technique we call "Tensor Coordinate Transformation".

Chao Liu's avatar
Chao Liu committed
10
![ALT](/doc/image/ck_component.png "CK Components")
Chao Liu's avatar
Chao Liu committed
11
12
13
14
15
16
17
18

## CK Structure
Current CK library are structure into 4 layers:
* "Templated Tile Operators"
* "Templated Kernel and Invoker" layer
* "Instantiated Kernel and Invoker" layer
* "Client API" layer

Chao Liu's avatar
Chao Liu committed
19
![ALT](/doc/image/ck_layer.png "CK Layers")
Chao Liu's avatar
Chao Liu committed
20

Chao Liu's avatar
update  
Chao Liu committed
21
## Contributors
Chao Liu's avatar
Chao Liu committed
22
The list of developers and contributors is here: [Contributors](/CONTRIBUTORS.md)
Chao Liu's avatar
Chao Liu committed
23
24

## Citation
Chao Liu's avatar
Chao Liu committed
25
CK paper will be freely available on arXiv soon: [Realizing Tensor Operators Using Coordinate Transformations and Tile Based Programming](???)
Chao Liu's avatar
Chao Liu committed
26
27

## License
Chao Liu's avatar
Chao Liu committed
28
CK is released under the MIT license. [License File](/LICENSE)
Chao Liu's avatar
Chao Liu committed
29
30
31
32


# Build CK

Chao Liu's avatar
Chao Liu committed
33
34
35
36
37
38
## Build docker image
```bash
DOCKER_BUILDKIT=1 docker build -t ck:latest -f Dockerfile .
```

## Launch docker
39
40
41
42
43
44
45
```bash
docker run                                     \
-it                                            \
--privileged                                   \
--group-add sudo                               \
-w /root/workspace                             \
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace  \
Chao Liu's avatar
Chao Liu committed
46
ck:latest                                      \
47
48
/bin/bash
```
Chao Liu's avatar
Chao Liu committed
49

Chao Liu's avatar
Chao Liu committed
50
## Build CK
51
52
53
```bash
mkdir build && cd build

Chao Liu's avatar
Chao Liu committed
54
55
56
57
58
59
60
# Need to specify target ID, example below is for gfx908 and gfx90a
cmake                                                                                             \
-D CMAKE_PREFIX_PATH=/opt/rocm                                                                    \
-D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                         \
-D CMAKE_CXX_FLAGS="-O3"                                                                          \
-D CMAKE_BUILD_TYPE=Release                                                                       \
-D GPU_TARGETS=gfx908;gfx90a                                                                      \
61
62
63
..
```

Chao Liu's avatar
Chao Liu committed
64
### Build examples and tests
65
```bash
66
 make -j examples tests
67
68
69
 make test
```

Chao Liu's avatar
Chao Liu committed
70
Instructions for running each individual examples are under [example](/example)
Chao Liu's avatar
Chao Liu committed
71
72


73
74
75
76
## Build ckProfiler
```bash
 make -j ckProfiler
```
Chao Liu's avatar
Chao Liu committed
77
Instructions for running ckProfiler are under [profiler](/profiler)
JD's avatar
JD committed
78

79
80
81
82
83
84
## Install CK
```bash
make install
```

## Using CK as pre-built kernel library
Chao Liu's avatar
Chao Liu committed
85
Instructions for using CK as a pre-built kernel library are under [client_example](/client_example)
JD's avatar
JD committed
86
87
88
89
90
91
92
93
94

## 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 verfication failure.
To work around it, do not 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.