Bf16*fp4 gemm

example/ck_tile/45_bf16_mxfp4_blockscale_gemm/CMakeLists.txt

@@ -0,0 +1,8 @@
+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_exampl_bf16_mxfp4_gemm PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})

example/ck_tile/45_bf16_mxfp4_blockscale_gemm/README.md

@@ -0,0 +1,36 @@
+# GEMM Matrix Multiplication
+
+This folder contains example for GEMM using ck_tile tile-programming implementation. Currently, it only supports the basic feature of the CK Tile GEMM, but creates the placeholders for the future support on different GEMM pipeline and different GEMM modules. In the near future, we will gradually migrate all the GEMM features from old CK to CK Tile.
+
+## build
+```
+# in the root of ck_tile
+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 pipeline method on the gemm calculation
+make tile_exampl_bf16_mxfp4_gemm -j
+
+```
+This will result in an executable `build/bin/tile_exampl_bf16_mxfp4_gemm`
+
+## example
+```
+args:
+          -b    batch size (default:1)
+          -m    m dimension (default:1024)
+          -n    n dimension (default:2048)
+          -k    k dimension (default:64)
+   -a_layout    Tensor A data layout (default: R)
+   -b_layout    Tensor B data layout (default: C)
+   -c_layout    Tensor C data layout (default: R)
+   -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
+          -e    Absolute error tolerance (default:1e-5)
+       -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/45_bf16_mxfp4_blockscale_gemm/benchmark_universal_fp4.sh

@@ -0,0 +1,25 @@
+#!/bin/bash
+
+EXE="$(find . -name tile_example_fp4_uint8_gemm -type f | head -n 1)"
+VALID=1
+
+N1=5888
+K1=3072
+N2=3072
+K2=2944
+
+
+#m_values=(1 16 32 64 128 256 512 1024 4096 16384)
+m_values=(1 16 64 256 512 1024 4096 16384)
+
+for m in "${m_values[@]}"; do
+    #echo "Running tests for m=$m"
+
+    # echo "Running test with m=$m, n=$N1, k=$K1"
+    $EXE -prec=pk_fp4_t -m=$m -n=$N1 -k=$K1 -v=1
+
+    # echo "Running test with m=$m, n=$N2, k=$K2"
+    $EXE -prec=pk_fp4_t -m=$m -n=$N2 -k=$K2 -v=1
+
+    # echo "Finished tests for m=$m"
+done