From 89bf0765fb17d7f9bd47e659ca5b495185fc8b67 Mon Sep 17 00:00:00 2001 From: ClementLinCF <162283536+ClementLinCF@users.noreply.github.com> Date: Fri, 25 Apr 2025 07:15:12 +0800 Subject: [PATCH] Update README.md --- example/ck_tile/99_toy_example/README.md | 57 ++++++++++++++++++++++++ 1 file changed, 57 insertions(+) diff --git a/example/ck_tile/99_toy_example/README.md b/example/ck_tile/99_toy_example/README.md index e49ab0d29b..dd22e01cf0 100644 --- a/example/ck_tile/99_toy_example/README.md +++ b/example/ck_tile/99_toy_example/README.md @@ -53,3 +53,60 @@ make -j basic_flash_attention_fwd ```sh ./bin/basic_flash_attention_fwd 1 1 ``` + +## Advanced part +#### **GEMM Example** +##### Follow these steps to build and run the different kernels: +```sh + +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 + +``` + +```sh +./bin/basic_gemm 1 +``` + +#### **Flash Attention Forward Example** +##### Follow these steps to build the kernels + +```sh +# 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" .. + +# 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 .. +``` +```sh +./bin/basic_flash_attention_fwd 1 1 +``` + + +##### Follow these steps to build the codegen instances + +```sh +mkdir build +cd build +../script/cmake-ck-release.sh .. gfx942 +make -j codegen_basic_flash_attention_fwd +``` + +```sh +./bin/codegen_basic_flash_attention_fwd 1 1 64 16384 16384 128 128 +```