Notes on Composable Kernel library

Introduction

Many new developers contributing to the MIOpen are often confused by the fact that the same name, Composable Kernel (or simply CK), appears in different and logically unrelated parts of the codebase. This confusion is understandable: in reality, there isn't just one CK library, but actually three distinct libraries that share the same name.

Brief History

Initially, there was only a single CK library, and its source code was placed directly within the MIOpen's codebase. At that time, it was located in the composable_kernel directory. This library was introduced in PR#1821 (MLOpen).

Then the developers decided that a new kernel design was needed. The kernels from the first version of the library required passing tensor dimensions as constants, meaning each kernel had to be compiled specifically for each problem. This led to the creation of a new version of the library with dynamic kernels - they are generic and only need to be compiled once. A separated repository was created for the new library, but a full copy of that repository was added inside the MIOpen repository. The initial plan was to keep sources in sync. The last time this happened was in early 2022. After that, the sources inside the MIOpen repository were frozen, while the external repository continued to evolve. The new version of the library was added in PR#1059 (closed), PR#1071 (first attempt, reverted) and PR#1097 (final version). The directory with the old version of the library was renamed to static_composable_kernel.

At some point, the CK repository was completely reorganized. The new library was added as a dependency in PR#1692.

Current Status

We are currently working on standardizing the naming of all three CK libraries and separating their support code by placing related functions into distinct namespaces (PR#3749). The src/composable_kernel directory has been renamed to src/legacy_composable_kernel in PR#3843 to avoid confusion with CK. These libraries are now referred to as Static CK or static_ck (the earliest one), Legacy CK or legacy_ck (the later one), and CK (the modern version), respectively.

Usage of Libraries in MIOpen

Static CK is being used for convolutions (13 solvers). Legacy CK is being used in one convolution solver (disabled due to SWDEV-411729) and also for reduction. New CK is being used for convolutions (several solvers) and for layernorm (2 solvers).

All static convolution solvers based on Static CK are disabled by default for devices newer than mi200 and gfx10.