Add CK Tile Tutorials Folder with GEMM and COPY Kernel (#3038)

* feat: add tutorial folder with gemm tutorial

* chore: move copy kernel from examples folder to tutorial

* Update tutorial/ck_tile/01_naive_gemm/README.md

Co-authored-by: Copilot <[email protected]>

* Update tutorial/ck_tile/01_naive_gemm/README.md

Co-authored-by: Copilot <[email protected]>

* chore: remove handdrawn images

* docs: add write ups to explain the gemm kernel

* docs: add about block level pipeline and static distributed tensors

---------

Co-authored-by: Copilot <[email protected]>

Author: AviralGoelAMD

CMakeLists.txt

@@ -683,6 +683,12 @@ if(NOT GPU_ARCHS AND USER_GPU_TARGETS AND NOT MIOPEN_REQ_LIBS_ONLY)
         PACKAGE_NAME examples
    )
    add_subdirectory(example)
+
+   add_subdirectory(tutorial)
+   rocm_package_setup_component(tutorials
+        LIBRARY_NAME composablekernel
+        PACKAGE_NAME tutorials
+   )
    add_subdirectory(tile_engine)
    if(BUILD_TESTING)
        add_subdirectory(test)

example/ck_tile/CMakeLists.txt

@@ -25,7 +25,6 @@ add_subdirectory(22_gemm_multi_abd)
 add_subdirectory(35_batched_transpose)
 add_subdirectory(36_pooling)
 add_subdirectory(38_block_scale_gemm)
-add_subdirectory(39_copy)
 add_subdirectory(40_streamk_gemm)
 add_subdirectory(41_batched_contraction)
 

tutorial/CMakeLists.txt

@@ -0,0 +1,15 @@
+include_directories(BEFORE
+    ${PROJECT_SOURCE_DIR}/include
+    ${PROJECT_SOURCE_DIR}/library/include
+)
+
+message(STATUS "Building tutorials...")
+add_custom_target(tutorials)
+
+# add all tutorial subdir
+file(GLOB dir_list LIST_DIRECTORIES true *)
+FOREACH(subdir ${dir_list})
+    if(IS_DIRECTORY "${subdir}" AND EXISTS "${subdir}/CMakeLists.txt")
+        add_subdirectory(${subdir})
+    ENDIF()
+ENDFOREACH()

example/ck_tile/39_copy/CMakeLists.txt renamed to tutorial/ck_tile/00_copy_kernel/CMakeLists.txt

@@ -1,7 +1,9 @@
-add_executable(tile_example_copy EXCLUDE_FROM_ALL copy_basic.cpp)
+add_executable(tile_tutorial_copy_kernel EXCLUDE_FROM_ALL copy_basic.cpp)
 
 # Impact: This flag ensures that the compiler doesn't make 
 # assumptions about memory aliasing that could interfere with Composable Kernel's explicit memory access patterns.
-target_compile_options(tile_example_copy PRIVATE
+target_compile_options(tile_tutorial_copy_kernel PRIVATE
   -mllvm -enable-noalias-to-md-conversion=0
 )
+
+add_dependencies(tutorials tile_tutorial_copy_kernel)

example/ck_tile/39_copy/README.md renamed to tutorial/ck_tile/00_copy_kernel/README.md

@@ -54,10 +54,10 @@ bool run(const ck_tile::ArgParser& arg_parser)
     x_buf.ToDevice(x_host.data());
 
     // Define tile configuration
-    using ThreadTile = ck_tile::sequence<1, 4>;   // per-thread tile size along M and N
-    using WaveTile   = ck_tile::sequence<64, 4>;  // wave size along M and N dimension
-    using BlockWaves = ck_tile::sequence<4, 1>;   // number of waves along M dimension
-    using BlockTile  = ck_tile::sequence<512, 4>; // block size along M and N dimension
+    using ThreadTile = ck_tile::sequence<1, 4>;  // per-thread tile size along M and N
+    using WaveTile   = ck_tile::sequence<64, 4>; // per-wave tile size along M and N dimension
+    using BlockWaves = ck_tile::sequence<4, 1>; // number of waves per block along M and N dimension
+    using BlockTile  = ck_tile::sequence<512, 4>; // per-block tile size along M and N dimension
 
     // Calculate grid size
     ck_tile::index_t kGridSize =
@@ -68,14 +68,14 @@ bool run(const ck_tile::ArgParser& arg_parser)
     using Shape   = ck_tile::TileCopyShape<BlockWaves, BlockTile, WaveTile, ThreadTile>;
     using Problem = ck_tile::TileCopyProblem<XDataType, Shape>;
     using Policy  = ck_tile::TileCopyPolicy<Problem>;
-    using Kernel  = ck_tile::ElementWiseTileCopyKernel<Problem, Policy>;
-    // using Kernel  = ck_tile::TileCopyKernel<Problem, Policy>;
-    // using Kernel = ck_tile::TileCopyKernel_LDS<Problem, Policy>;
-
-    // question: Why do we not have a pipeline?
-    // answer: For basic copy operation, pipeline is not needed.
-    // we intentionally do not use pipeline for this example and let the kernel be composite of
-    // Problem and Policy
+    using Kernel  = ck_tile::ElementWiseTileCopyKernel<Problem, Policy>; // operates on element by
+                                                                         // element basis.
+
+    // We also implement two variations of the copy kernel:
+    // 1. TileCopyKernel: This is the basic copy kernel that operates on tile by tile basis.
+    // 2. TileCopyKernel_LDS: This is the copy kernel that operates on tile by tile basis and uses
+    // the LDS. using Kernel  = ck_tile::TileCopyKernel<Problem, Policy>; using Kernel =
+    // ck_tile::TileCopyKernel_LDS<Problem, Policy>;
 
     auto blockSize = Kernel::BlockSize();
 

example/ck_tile/39_copy/copy_basic.cpp renamed to tutorial/ck_tile/00_copy_kernel/copy_basic.cpp

@@ -4,7 +4,7 @@
 
 set -euo pipefail
 
-BIN="${BIN:-../../../build/bin/tile_example_copy}"
+BIN="${BIN:-../../../build/bin/tile_tutorial_copy_kernel}"
 WARMUP="${WARMUP:-20}"
 REPEAT="${REPEAT:-100}"
 VALIDATE="${VALIDATE:-1}"