rename example folder

Author: eliotwang

example/ck_tile/40_fp4_uint8_gemm/CMakeLists.txt renamed to example/ck_tile/40_bf16_mxfp4_blockscale_gemm/CMakeLists.txt

@@ -1,8 +1,8 @@
-add_executable(tile_example_fp4_uint8_gemm EXCLUDE_FROM_ALL fp4_uint8_gemm.cpp)
+add_executable(tile_exampl_bf16_mxfp4_gemm  EXCLUDE_FROM_ALL bf16_mxfp4_blockscale_gemm.cpp)
 
 set(EXAMPLE_GEMM_COMPILE_OPTIONS)
 if(CK_USE_OCP_FP8)
   list(APPEND EXAMPLE_GEMM_COMPILE_OPTIONS -DCK_TILE_USE_OCP_FP8)
 endif()
 list(APPEND EXAMPLE_GEMM_COMPILE_OPTIONS -mllvm -enable-noalias-to-md-conversion=0)
-target_compile_options(tile_example_fp4_uint8_gemm PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
+target_compile_options(tile_exampl_bf16_mxfp4_gemm PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})

example/ck_tile/40_fp4_uint8_gemm/README.md renamed to example/ck_tile/40_bf16_mxfp4_blockscale_gemm/README.md

@@ -8,12 +8,11 @@ This folder contains example for GEMM using ck_tile tile-programming implementat
 mkdir build && cd build
 # you can replace <arch> with the appropriate architecture (for example gfx90a or gfx942) or leave it blank
 sh ../script/cmake-ck-dev.sh  ../ <arch>
-# The basic pipeline method on the gemm calculation
-make tile_example_gemm_basic -j
-# The memory bound pipeline on the gemm calculation
-make tile_example_gemm_universal -j
+# The pipeline method on the gemm calculation
+make tile_exampl_bf16_mxfp4_gemm -j
+
 ```
-This will result in an executable `build/bin/tile_example_gemm_basic` & `build/bin/tile_example_gemm_universal`
+This will result in an executable `build/bin/tile_exampl_bf16_mxfp4_gemm`
 
 ## example
 ```
@@ -28,9 +27,9 @@ args:
    -stride_a    Tensor A stride (default:0)
    -stride_b    Tensor B stride (default:0)
    -stride_c    Tensor C stride (default:0)
-          -v    0. No validation, 1. Validation on CPU, 2. Validation on GPU (default:2)
+          -v    0. No validation, 1. Validation on CPU
           -e    Absolute error tolerance (default:1e-5)
-       -prec    data type. fp16/bf16/fp8/bf8/int8 (default:fp16)
+       -prec    data type. fp16/bf16/fp8/bf8/int8/pk_fp4_t (default:fp16)
      -warmup    number of iterations before benchmark the kernel (default:10)
      -repeat    number of iterations to benchmark the kernel (default:100)
       -timer    gpu:gpu timer, cpu:cpu timer (default:gpu)

example/ck_tile/40_fp4_uint8_gemm/benchmark_universal_fp4.sh renamed to example/ck_tile/40_bf16_mxfp4_blockscale_gemm/benchmark_universal_fp4.sh

@@ -175,17 +175,16 @@ struct UniversalInvoker
 
                 auto size_a_buffer = a_m.get_element_space_size_in_bytes();
                 auto size_b_buffer = b_n.get_element_space_size_in_bytes();
-                               
+                                       
                 rotating_mem_ptr =
                     std::make_unique<ck_tile::RotatingMemWrapper<ADataType, BInDataType>>(
                         kargs.as_ptr[0],
                         kargs.bs_ptr[0],
                         s.rotating_count_,
                         size_a_buffer,
                         size_b_buffer);
-                rotating_mem_ptr->Print();
+                rotating_mem_ptr->Print();  
 
-
                 preprocess = [&]() {
                     ck_tile::flush_icache();
                     rotating_mem_ptr->Next();
@@ -297,53 +296,6 @@ int run_gemm_example(ck_tile::ArgParser& arg_parser)
     std::string b_layout  = arg_parser.get_str("b_layout");
 
     using Invoker = UniversalInvoker;
-    // if(data_type == "fp16")
-    // {
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t>, ck_tile::half_t>(
-    //         a_layout, b_layout, argc, argv);
-    // }
-    // else if(data_type == "bf16")
-    // {
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t>, ck_tile::bf16_t>(
-    //         a_layout, b_layout, argc, argv);
-    // }
-    // else if(data_type == "fp8")
-    // {
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::fp8_t>,
-    //                                       ck_tile::fp8_t,
-    //                                       ck_tile::fp8_t,
-    //                                       ck_tile::bf16_t>(a_layout, b_layout, argc, argv);
-    // }
-    // else if(data_type == "bf8")
-    // {
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::bf8_t>,
-    //                                       ck_tile::bf8_t,
-    //                                       ck_tile::bf8_t,
-    //                                       ck_tile::half_t>(a_layout, b_layout, argc, argv);
-    // }
-    // else if(data_type == "int8")
-    // {
-    //     return run_gemm_example_prec_type<GemmConfig<ck_tile::int8_t>,
-    //                                       ck_tile::int8_t,
-    //                                       ck_tile::int8_t,
-    //                                       ck_tile::int32_t>(a_layout, b_layout, argc, argv);
-    // }
-    // else if(data_type == "pk_int4_t")
-    // {
-    //     // TODO: Add support for bhalf_t ADataType
-    //     if constexpr(GemmConfig<ck_tile::half_t>::Pipeline == CK_TILE_PIPELINE_COMPUTE_V3)
-    //     {
-    //         return run_gemm_example_prec_type<GemmConfig<ck_tile::half_t>,
-    //                                           ck_tile::half_t,
-    //                                           ck_tile::pk_int4_t,
-    //                                           ck_tile::half_t>(a_layout, b_layout, argc, argv);
-    //     }
-    //     else
-    //     {
-    //         throw std::runtime_error("Unsupported pipeline for this operation !!!");
-    //     }
-    // }
-    // else if(data_type == "pk_fp4_t")
     if(data_type == "pk_fp4_t")
     {
         // TODO: Add support for bhalf_t ADataType