Ck tile tutorial examples

[ClementLinCF]

Proposed changes

Examples of educational materials related to CK TILE.

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

• [ x] I have added tests relevant to the introduced functionality, and the unit tests are passing locally
• I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
• I have added inline documentation which enables the maintainers with understanding the motivation
• I have removed the stale documentation which is no longer relevant after this pull request
• (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
• [x ] I have run clang-format on all changed files
• [ x] Any dependent changes have been merged

[AviralGoelAMD]

Everything looks!

[spolifroni-amd]

Doesn't look like there's anything for docs to review

[ClementLinCF]

Doesn't look like there's anything for docs to review

We have a README file https://github.com/ROCm/composable_kernel/blob/ck_tile_tutorial/example/ck_tile/tutorial/README.md in the tutorial folder that contains the following content. It provides basic examples of ck_tile. In addition to the explanations found in the comments of the examples, there are currently some ongoing internal CK materials, and the demonstrations within them will refer to the examples located under this tutorial.

CK_TILE Tutorial Example

This folder demonstrates the examples implemented using ck_tile

Build Instructions

Follow these steps to build the examples:

cd composable_kernel
mkdir build
cd build

cmake -D CMAKE_PREFIX_PATH=/opt/rocm \
      -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc \
      -D CMAKE_BUILD_TYPE=Release \
      -D GPU_TARGETS="gfx942" \
      -Dkernel=N ..

Compile Examples

Basic Add Example

make -j add_basic

Elementwise Add Example

make -j add

GEMM Example

make -j basic_gemm

Flash Attention Forward Example

make -j basic_flash_attention_fwd

Running Examples

Basic Add Example

./bin/add_basic

Elementwise Add

./bin/add

GEMM Example

./bin/basic_gemm 1

Flash Attention Forward Example

./bin/basic_flash_attention_fwd 1 1

Advanced part

GEMM Example

Follow these steps to build and run the different kernels:

cd composable_kernel
mkdir build
cd build

# for naive kernel
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" -Dkernel=N .. && make -j basic_gemm

# for kernel A
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" -Dkernel=A .. && make -j basic_gemm

# for kernel B
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" -Dkernel=B .. && make -j basic_gemm

...

# for kernel H
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" -Dkernel=H .. && make -j basic_gemm

./bin/basic_gemm 1

Flash Attention Forward Example

Follow these steps to build the kernels

# for naive kernel
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" .. && make -j basic_flash_attention_fwd

# for optimized kernel
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Release -D GPU_TARGETS="gfx942" -DENABLE_TOY_FA_FWD_OPT=ON .. && make -j basic_flash_attention_fwd

./bin/basic_flash_attention_fwd 1 1

Follow these steps to build and run the codegen instances

mkdir build
cd build
../script/cmake-ck-release.sh .. gfx942
make -j codegen_basic_flash_attention_fwd

./bin/codegen_basic_flash_attention_fwd 1 1 64 16384 16384 128 128