diff --git a/clang_format_git.py b/clang_format_git.py
new file mode 100644
index 0000000000..bf4011febc
--- /dev/null
+++ b/clang_format_git.py
@@ -0,0 +1,140 @@
+#!/usr/bin/env python3
+
+import subprocess
+import sys
+import re
+import os
+from pathlib import Path
+from typing import List
+
+# 定义颜色代码
+class Color:
+    GREEN = '\033[0;32m'
+    YELLOW = '\033[0;33m'
+    RED = '\033[0;31m'
+    NC = '\033[0m'  # 无颜色
+
+def print_color(color: str, message: str):
+    """打印带颜色的消息"""
+    print(f"{color}{message}{Color.NC}")
+
+def check_command_exists(command: str) -> bool:
+    """检查命令是否存在"""
+    try:
+        subprocess.run(['which', command], stdout=subprocess.PIPE, stderr=subprocess.PIPE, check=True)
+        return True
+    except subprocess.CalledProcessError:
+        return False
+
+def is_git_repo() -> bool:
+    """检查当前目录是否是Git仓库"""
+    try:
+        subprocess.run(['git', 'rev-parse', '--is-inside-work-tree'], 
+                     stdout=subprocess.PIPE, stderr=subprocess.PIPE, check=True)
+        return True
+    except subprocess.CalledProcessError:
+        return False
+
+def get_git_modified_files() -> List[str]:
+    """获取Git中已修改但未提交的文件及未跟踪的新文件"""
+    try:
+        # 获取已修改的文件
+        modified_result = subprocess.run(['git', 'status', '--porcelain'], 
+                                      stdout=subprocess.PIPE, text=True, check=True)
+        
+        files = []
+        for line in modified_result.stdout.splitlines():
+            if not line.strip():
+                continue
+            
+            # 检查文件状态（M=修改, A=添加, ??=未跟踪）
+            if re.match(r'^\s*[AM]', line) or line.startswith('??'):
+                # 提取文件名
+                parts = line.strip().split(maxsplit=1)
+                if len(parts) > 1:
+                    files.append(parts[1])
+                else:
+                    # 处理未跟踪文件的情况
+                    files.append(parts[0][2:].strip())
+                    
+        return files
+    except subprocess.CalledProcessError as e:
+        print_color(Color.RED, f"获取Git修改文件失败: {e}")
+        return []
+
+def filter_cpp_files(files: List[str]) -> List[str]:
+    """筛选C++相关文件"""
+    cpp_extensions = ['.cpp', '.hpp', '.h', '.cc', '.c', '.cxx']
+    return [file for file in files if Path(file).suffix.lower() in cpp_extensions]
+
+def format_files(files: List[str]) -> tuple:
+    """使用clang-format格式化文件"""
+    success_count = 0
+    fail_count = 0
+    
+    for file in files:
+        print(f"格式化: {file} ... ", end="")
+        try:
+            # 先运行dos2unix确保文件使用Unix换行符
+            subprocess.run(['dos2unix', file], 
+                         stdout=subprocess.PIPE, stderr=subprocess.PIPE, check=True)
+            
+            # 使用clang-format格式化文件
+            subprocess.run(['clang-format-12', '-style=file', '-i', file], 
+                         stdout=subprocess.PIPE, stderr=subprocess.PIPE, check=True)
+            
+            print_color(Color.GREEN, "成功")
+            success_count += 1
+        except subprocess.CalledProcessError as e:
+            print_color(Color.RED, f"失败: {e}")
+            fail_count += 1
+    
+    return success_count, fail_count
+
+def main():
+    # 检查clang-format-12是否安装
+    if not check_command_exists('clang-format-12'):
+        print_color(Color.RED, "错误: clang-format-12 未安装或不在PATH中")
+        sys.exit(1)
+        
+    # 检查dos2unix是否安装
+    if not check_command_exists('dos2unix'):
+        print_color(Color.YELLOW, "警告: dos2unix 未安装，将跳过行尾符转换")
+        
+    # 检查是否在Git仓库中
+    if not is_git_repo():
+        print_color(Color.RED, "错误: 当前目录不是Git仓库")
+        sys.exit(1)
+    
+    print_color(Color.YELLOW, "获取Git修改的文件列表...")
+    
+    # 获取修改的文件
+    all_files = get_git_modified_files()
+    
+    # 筛选C++文件
+    cpp_files = filter_cpp_files(all_files)
+    
+    if not cpp_files:
+        print_color(Color.YELLOW, "没有找到需要格式化的C++文件")
+        sys.exit(0)
+    
+    print_color(Color.GREEN, "找到以下文件需要格式化:")
+    for file in cpp_files:
+        print(f"  - {file}")
+    
+    print_color(Color.YELLOW, "开始格式化文件...")
+    
+    # 格式化文件
+    success_count, fail_count = format_files(cpp_files)
+    
+    print()
+    print_color(Color.GREEN, "格式化完成!")
+    print_color(Color.GREEN, f"  - 成功: {success_count} 个文件")
+    if fail_count > 0:
+        print_color(Color.RED, f"  - 失败: {fail_count} 个文件")
+    
+    print()
+    print_color(Color.YELLOW, "提示: 您可以使用 'git diff' 查看格式化后的变更")
+
+if __name__ == "__main__":
+    main()
diff --git a/example/ck_tile/18_flatmm/flatmm_basic.cpp b/example/ck_tile/18_flatmm/flatmm_basic.cpp
index 7c55086c30..141f598737 100644
--- a/example/ck_tile/18_flatmm/flatmm_basic.cpp
+++ b/example/ck_tile/18_flatmm/flatmm_basic.cpp
@@ -13,7 +13,6 @@
 #include "flatmm_basic.hpp"
 #include <type_traits>
 
-
 template <typename T>
 constexpr const char* DataTypeToString()
 {
@@ -84,7 +83,6 @@ auto calculate_rtol_atol(const ck_tile::index_t K,
     return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k));
 }
 
-
 template <typename FlatmmConfig,
           typename ADataType,
           typename BDataType,
@@ -434,6 +432,14 @@ int run_flatmm_example(int argc, char* argv[])
                     argc, argv, Row{}, Col{}, Row{});
             }
         }
+        else if(data_type == "fp4")
+        {
+            run_flatmm_example_with_layouts<ck_tile::pk_fp4_t,
+                                            FlatmmConfig<ck_tile::pk_fp4_t>,
+                                            32,
+                                            32>(argc, argv, Row{}, Col{}, Row{});
+        }
+
         else
         {
             throw std::runtime_error("Unsupported data_type!");
diff --git a/example/ck_tile/18_flatmm/flatmm_basic.hpp b/example/ck_tile/18_flatmm/flatmm_basic.hpp
index 5e155358bf..422eb93cc5 100644
--- a/example/ck_tile/18_flatmm/flatmm_basic.hpp
+++ b/example/ck_tile/18_flatmm/flatmm_basic.hpp
@@ -129,6 +129,16 @@ struct GemmBasicTypeConfig<ck_tile::bf8_t>
     using CDataType   = ck_tile::half_t;
 };
 
+template <>
+struct GemmBasicTypeConfig<ck_tile::pk_fp4_t>
+{
+    using ADataType   = ck_tile::pk_fp4_t;
+    using BDataType   = ck_tile::pk_fp4_t;
+    using AccDataType = float;
+    using XDataType   = ck_tile::e8m0_bexp_t; // scale data type
+    using CDataType   = ck_tile::half_t;
+}
+
 template <typename T>
 struct DataTypeTraits;
 
diff --git a/example/ck_tile/18_flatmm/run_flatmm_example.inc b/example/ck_tile/18_flatmm/run_flatmm_example.inc
index 890e02b32c..ef94590a06 100644
--- a/example/ck_tile/18_flatmm/run_flatmm_example.inc
+++ b/example/ck_tile/18_flatmm/run_flatmm_example.inc
@@ -6,7 +6,7 @@ template <typename PrecType,
           typename FlatmmConfig,
           int ScaleGranularityM = -1,
           int ScaleGranularityN = -1,
-          typename ALayout,
+          int ScaleGranularityK = -1,
           typename BLayout,
           typename CLayout>
 int run_flatmm_example_with_layouts(int argc,
@@ -49,200 +49,252 @@ int run_flatmm_example_with_layouts(int argc,
     ck_tile::HostTensor<CDataType> c_rslt_host(
         ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{})));
 
-    ck_tile::HostTensor<AccDataType> per_token_scale(ck_tile::HostTensorDescriptor({M}, {1}));
-    ck_tile::HostTensor<AccDataType> per_channel_scale(ck_tile::HostTensorDescriptor({N}, {1}));
-
-    // TODO: add different init types
-    if(init_method == 0)
+    // mxfp4 gemm, scale data type is e8m0_bexp_t
+    if(std::is_same_v<ADataType, ck_tile::pk_fp4_t> || std::is_same_v<BDataType, ck_tile::pk_fp4_t>)
     {
-        ck_tile::FillUniformDistribution<ADataType>{-.5f, .5f}(a_host);
-        ck_tile::FillUniformDistribution<BDataType>{-.5f, .5f}(b_origin_host);
-        ck_tile::FillUniformDistribution<AccDataType>{-1.f, 1.f}(per_token_scale);
-        ck_tile::FillUniformDistribution<AccDataType>{-1.f, 1.f}(per_channel_scale);
-    }
-    else if(init_method == 1)
-    {
-        ck_tile::FillMonotonicSeq<ADataType>{}(a_host);
-        ck_tile::FillMonotonicSeq<BDataType>{}(b_origin_host);
-        ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_token_scale);
-        ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_channel_scale);
-    }
-    else if(init_method == 2)
-    {
-        ck_tile::FillUniformDistribution<ADataType>{1.f, 1.f}(a_host);
-        ck_tile::FillUniformDistribution<BDataType>{1.f, 1.f}(b_origin_host);
-        ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_token_scale);
-        ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_channel_scale);
-    }
-    else
-    {
-        a_host.SetZero();
-        b_origin_host.SetZero();
-    }
+        if(ScaleGranularityM != 32 || ScaleGranularityN != 32 || ScaleGranularityK != 32)
+            throw std::runtime_error("wrong! mxfp4 gemm blockscale size should be 32.\n");
 
-    ck_tile::DeviceMem a_dev_buf(a_host.get_element_space_size_in_bytes());
-    ck_tile::DeviceMem c_dev_buf(c_rslt_host.get_element_space_size_in_bytes());
+        if(K % ScaleGranularityK != 0)
+            throw std::runtime_error("wrong! K must be multiple of scaleGranulartiyK.\n");
 
-    ck_tile::DeviceMem per_token_scale_dev_buf(per_token_scale.get_element_space_size_in_bytes());
-    ck_tile::DeviceMem per_channel_scale_dev_buf(
-        per_channel_scale.get_element_space_size_in_bytes());
+        if(K % ck_tile::numeric_traits<ck_tile::pk_fp4_t>::PackedSize != 0)
+            throw std::runtime_error("wrong! K must be multiple of PackedSize.\n");
 
-    a_dev_buf.ToDevice(a_host.data());
-    c_rslt_host.SetZero();
-    per_token_scale_dev_buf.ToDevice(per_token_scale.data());
-    per_channel_scale_dev_buf.ToDevice(per_channel_scale.data());
+        auto Scale_Padded_M = (M + ScaleGranularityM - 1) / ScaleGranularityM * ScaleGranularityM;
 
-    // do pre-shuffle
-    ck_tile::HostTensor<BDataType> b_shuffle_host = shuffle_b<FlatmmConfig>(b_origin_host);
-    ck_tile::DeviceMem b_shuffle_dev_buf(b_shuffle_host.get_element_space_size_in_bytes());
-    b_shuffle_dev_buf.ToDevice(b_shuffle_host.data());
+        auto Scale_stride_AM = ck_tile::get_default_stride(
+            scale_Padded_M, ScaleGranularityM, -1, is_row_major(a_layout));
+        auto Scale_stride_BN =
+            ck_tile::get_default_stride(K / ScaleGranularityK, N, -1, is_row_major(b_layout));
 
-    auto per_token_scale_dev_ptr = ck_tile::FlatmmScalePointer<ScaleGranularityM>{
-        static_cast<float*>(per_token_scale_dev_buf.GetDeviceBuffer())};
-    auto per_channel_scale_dev_ptr = ck_tile::FlatmmScalePointer<ScaleGranularityN>{
-        static_cast<float*>(per_channel_scale_dev_buf.GetDeviceBuffer())};
+        using AscaleDataType = typename GemmBasicTypeConfig<PrecType>::XDataType;
+        using BscaleDataType = typename GemmBasicTypeConfig<PrecType>::XDataType;
 
-    invoke_flatmm<FlatmmConfig,
-                  ADataType,
-                  BDataType,
-                  ck_tile::tuple<>,
-                  AccDataType,
-                  CDataType,
-                  ALayout,
-                  BLayout,
-                  ck_tile::tuple<>,
-                  CLayout,
-                  decltype(per_token_scale_dev_ptr),
-                  decltype(per_channel_scale_dev_ptr)>(a_dev_buf,
-                                                       b_shuffle_dev_buf,
-                                                       c_dev_buf,
-                                                       M,
-                                                       N,
-                                                       K,
-                                                       stride_A,
-                                                       stride_B,
-                                                       stride_C,
-                                                       kbatch,
-                                                       per_token_scale_dev_ptr,
-                                                       per_channel_scale_dev_ptr,
-                                                       n_warmup,
-                                                       n_repeat);
+        // A, B scale tensors
+        ck_tile::HostTensor<AscaleDataType> a_m_k_scale(ck_tile::host_tensor_descriptor(
+            Scale_Padded_M, K / ScaleGranularityK, Scale_stride_AM, is_row_major(a_layout)));
+        ck_tile::HostTensor<BscaleDataType> b_k_n_scale(ck_tile::host_tensor_descriptor(
+            K / ScaleGranularityK, N, Scale_stride_BN, is_row_major(b_layout)));
 
-    c_dev_buf.FromDevice(c_rslt_host.data());
-    bool pass = true;
+        // A, B scale shuffled tensors
+        ck_tile::HostTensor<AscaleDataType> a_m_k_scale_shuffle(ck_tile::host_tensor_descriptor(
+            Scale_Padded_M, K / ScaleGranularityK, Scale_stride_AM, is_row_major(a_layout)));
+        ck_tile::HostTensor<BscaleDataType> b_k_n_scale_shuffle(ck_tile::host_tensor_descriptor(
+            K / ScaleGranularityK, N, Scale_stride_BN, is_row_major(b_layout)));
 
-    if(arg_parser.get_int("v") == 1)
-    {
-        if (ScaleGranularityM != -1 || ScaleGranularityN != -1)
-            throw std::runtime_error("ScaleAB is not supported for CPU verification!\n");
-        ck_tile::HostTensor<CDataType> c_ref_host(
-            ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{})));
-        c_ref_host.SetZero();
-
-        ck_tile::reference_gemm<ADataType, BDataType, AccDataType, CDataType>(
-            a_host, b_origin_host, c_ref_host);
-        const float max_accumulated_value =
-            *std::max_element(c_ref_host.mData.begin(), c_ref_host.mData.end());
-        const auto rtol_atol = calculate_rtol_atol<ADataType, BDataType, AccDataType, CDataType>(
-            K, kbatch, max_accumulated_value);
-        pass = ck_tile::check_err(c_rslt_host,
-                                  c_ref_host,
-                                  "Error: Incorrect results!",
-                                  rtol_atol.at(ck_tile::number<0>{}),
-                                  rtol_atol.at(ck_tile::number<1>{}));
-
-        std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
-                  << " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
-                  << std::endl;
-        std::cout << "The CPU veification result is:" << (pass ? "correct" : "fail") << std::endl;
-    }
-    else if(arg_parser.get_int("v") == 2)
-    {
-        ck_tile::DeviceMem b_origin_dev_buf(b_origin_host.get_element_space_size_in_bytes());
-        b_origin_dev_buf.ToDevice(b_origin_host.data());
-
-        ck_tile::HostTensor<CDataType> c_gpu_ref_host(
-            ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{})));
-        ck_tile::DeviceMem c_gpu_ref_dev_buf(c_gpu_ref_host.get_element_space_size_in_bytes());
-        c_gpu_ref_host.SetZero();
-        c_gpu_ref_dev_buf.SetZero();
-
-        ADataType* d_A;
-        BDataType* d_B;
-        CDataType* d_C;
-
-        ck_tile::hip_check_error(hipMalloc(&d_A, M * K * sizeof(ADataType)));
-        ck_tile::hip_check_error(hipMalloc(&d_B, N * K * sizeof(BDataType)));
-        ck_tile::hip_check_error(hipMalloc(&d_C, M * N * sizeof(CDataType)));
-
-        ck_tile::hip_check_error(hipMemcpy(
-            d_A, a_dev_buf.GetDeviceBuffer(), M * K * sizeof(ADataType), hipMemcpyHostToDevice));
-        ck_tile::hip_check_error(hipMemcpy(d_B,
-                                           b_origin_dev_buf.GetDeviceBuffer(),
-                                           N * K * sizeof(BDataType),
-                                           hipMemcpyHostToDevice));
-
-        if constexpr(ScaleGranularityM == -1 && ScaleGranularityN == -1)
+        if(init_method == 0)
         {
-            ck_tile::reference_gemm_gpu<ADataType,
-                                        BDataType,
-                                        AccDataType,
-                                        CDataType,
-                                        ALayout,
-                                        BLayout,
-                                        CLayout>(
-                d_A, d_B, d_C, M, N, K, stride_A, stride_B, stride_C);
+            ck_tile::FillUniformDistribution<ADataType>{-.5f, .5f}(a_host);
+            ck_tile::FillUniformDistribution<BDataType>{-.5f, .5f}(b_origin_host);
+            ck_tile::FillUniformDistribution<AscaleDataType>{-1.f, 1.f}(a_m_k_scale);
+            ck_tile::FillUniformDistribution<BscaleDataType>{-1.f, 1.f}(b_k_n_scale);
+        }
+    }
+    else // PTPC unversal flat gemm
+    {
+        ck_tile::HostTensor<AccDataType> per_token_scale(ck_tile::HostTensorDescriptor({M}, {1}));
+        ck_tile::HostTensor<AccDataType> per_channel_scale(ck_tile::HostTensorDescriptor({N}, {1}));
+
+        // TODO: add different init types
+        if(init_method == 0)
+        {
+            ck_tile::FillUniformDistribution<ADataType>{-.5f, .5f}(a_host);
+            ck_tile::FillUniformDistribution<BDataType>{-.5f, .5f}(b_origin_host);
+            ck_tile::FillUniformDistribution<AccDataType>{-1.f, 1.f}(per_token_scale);
+            ck_tile::FillUniformDistribution<AccDataType>{-1.f, 1.f}(per_channel_scale);
+        }
+        else if(init_method == 1)
+        {
+            ck_tile::FillMonotonicSeq<ADataType>{}(a_host);
+            ck_tile::FillMonotonicSeq<BDataType>{}(b_origin_host);
+            ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_token_scale);
+            ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_channel_scale);
+        }
+        else if(init_method == 2)
+        {
+            ck_tile::FillUniformDistribution<ADataType>{1.f, 1.f}(a_host);
+            ck_tile::FillUniformDistribution<BDataType>{1.f, 1.f}(b_origin_host);
+            ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_token_scale);
+            ck_tile::FillUniformDistribution<AccDataType>{1.f, 1.f}(per_channel_scale);
         }
         else
         {
-            ck_tile::reference_blockwise_gemm_gpu<ADataType,
-                                                  BDataType,
-                                                  AccDataType,
-                                                  CDataType,
-                                                  ALayout,
-                                                  BLayout,
-                                                  CLayout>(
-                d_A,
-                d_B,
-                d_C,
-                M,
-                N,
-                K,
-                stride_A,
-                stride_B,
-                stride_C,
-                ScaleGranularityM,
-                ScaleGranularityN,
-                K,
-                static_cast<float*>(per_token_scale_dev_buf.GetDeviceBuffer()),
-                static_cast<float*>(per_channel_scale_dev_buf.GetDeviceBuffer()));
+            a_host.SetZero();
+            b_origin_host.SetZero();
         }
 
-        ck_tile::hip_check_error(hipMemcpy(c_gpu_ref_dev_buf.GetDeviceBuffer(),
-                                           d_C,
-                                           M * N * sizeof(CDataType),
-                                           hipMemcpyDeviceToHost));
+        ck_tile::DeviceMem a_dev_buf(a_host.get_element_space_size_in_bytes());
+        ck_tile::DeviceMem c_dev_buf(c_rslt_host.get_element_space_size_in_bytes());
 
-        ck_tile::hip_check_error(hipFree(d_A));
-        ck_tile::hip_check_error(hipFree(d_B));
-        ck_tile::hip_check_error(hipFree(d_C));
+        ck_tile::DeviceMem per_token_scale_dev_buf(
+            per_token_scale.get_element_space_size_in_bytes());
+        ck_tile::DeviceMem per_channel_scale_dev_buf(
+            per_channel_scale.get_element_space_size_in_bytes());
 
-        c_gpu_ref_dev_buf.FromDevice(c_gpu_ref_host.data());
-        const float max_accumulated_value =
-            *std::max_element(c_gpu_ref_host.mData.begin(), c_gpu_ref_host.mData.end());
-        const auto rtol_atol = calculate_rtol_atol<ADataType, BDataType, AccDataType, CDataType>(
-            K, kbatch, max_accumulated_value);
-        pass = ck_tile::check_err(c_rslt_host,
-                                  c_gpu_ref_host,
-                                  "Error: Incorrect results!",
-                                  rtol_atol.at(ck_tile::number<0>{}),
-                                  rtol_atol.at(ck_tile::number<1>{}));
+        a_dev_buf.ToDevice(a_host.data());
+        c_rslt_host.SetZero();
+        per_token_scale_dev_buf.ToDevice(per_token_scale.data());
+        per_channel_scale_dev_buf.ToDevice(per_channel_scale.data());
 
-        std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
-                  << " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
-                  << std::endl;
-        std::cout << "The GPU veification result is: " << (pass ? "correct" : "fail") << std::endl;
+        // do pre-shuffle
+        ck_tile::HostTensor<BDataType> b_shuffle_host = shuffle_b<FlatmmConfig>(b_origin_host);
+        ck_tile::DeviceMem b_shuffle_dev_buf(b_shuffle_host.get_element_space_size_in_bytes());
+        b_shuffle_dev_buf.ToDevice(b_shuffle_host.data());
+
+        auto per_token_scale_dev_ptr = ck_tile::FlatmmScalePointer<ScaleGranularityM>{
+            static_cast<float*>(per_token_scale_dev_buf.GetDeviceBuffer())};
+        auto per_channel_scale_dev_ptr = ck_tile::FlatmmScalePointer<ScaleGranularityN>{
+            static_cast<float*>(per_channel_scale_dev_buf.GetDeviceBuffer())};
+
+        invoke_flatmm<FlatmmConfig,
+                      ADataType,
+                      BDataType,
+                      ck_tile::tuple<>,
+                      AccDataType,
+                      CDataType,
+                      ALayout,
+                      BLayout,
+                      ck_tile::tuple<>,
+                      CLayout,
+                      decltype(per_token_scale_dev_ptr),
+                      decltype(per_channel_scale_dev_ptr)>(a_dev_buf,
+                                                           b_shuffle_dev_buf,
+                                                           c_dev_buf,
+                                                           M,
+                                                           N,
+                                                           K,
+                                                           stride_A,
+                                                           stride_B,
+                                                           stride_C,
+                                                           kbatch,
+                                                           per_token_scale_dev_ptr,
+                                                           per_channel_scale_dev_ptr,
+                                                           n_warmup,
+                                                           n_repeat);
+
+        c_dev_buf.FromDevice(c_rslt_host.data());
+        bool pass = true;
+
+        if(arg_parser.get_int("v") == 1)
+        {
+            if(ScaleGranularityM != -1 || ScaleGranularityN != -1)
+                throw std::runtime_error("ScaleAB is not supported for CPU verification!\n");
+            ck_tile::HostTensor<CDataType> c_ref_host(
+                ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{})));
+            c_ref_host.SetZero();
+
+            ck_tile::reference_gemm<ADataType, BDataType, AccDataType, CDataType>(
+                a_host, b_origin_host, c_ref_host);
+            const float max_accumulated_value =
+                *std::max_element(c_ref_host.mData.begin(), c_ref_host.mData.end());
+            const auto rtol_atol =
+                calculate_rtol_atol<ADataType, BDataType, AccDataType, CDataType>(
+                    K, kbatch, max_accumulated_value);
+            pass = ck_tile::check_err(c_rslt_host,
+                                      c_ref_host,
+                                      "Error: Incorrect results!",
+                                      rtol_atol.at(ck_tile::number<0>{}),
+                                      rtol_atol.at(ck_tile::number<1>{}));
+
+            std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
+                      << " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
+                      << std::endl;
+            std::cout << "The CPU veification result is:" << (pass ? "correct" : "fail")
+                      << std::endl;
+        }
+        else if(arg_parser.get_int("v") == 2)
+        {
+            ck_tile::DeviceMem b_origin_dev_buf(b_origin_host.get_element_space_size_in_bytes());
+            b_origin_dev_buf.ToDevice(b_origin_host.data());
+
+            ck_tile::HostTensor<CDataType> c_gpu_ref_host(
+                ck_tile::host_tensor_descriptor(M, N, stride_C, is_row_major(CLayout{})));
+            ck_tile::DeviceMem c_gpu_ref_dev_buf(c_gpu_ref_host.get_element_space_size_in_bytes());
+            c_gpu_ref_host.SetZero();
+            c_gpu_ref_dev_buf.SetZero();
+
+            ADataType* d_A;
+            BDataType* d_B;
+            CDataType* d_C;
+
+            ck_tile::hip_check_error(hipMalloc(&d_A, M * K * sizeof(ADataType)));
+            ck_tile::hip_check_error(hipMalloc(&d_B, N * K * sizeof(BDataType)));
+            ck_tile::hip_check_error(hipMalloc(&d_C, M * N * sizeof(CDataType)));
+
+            ck_tile::hip_check_error(hipMemcpy(d_A,
+                                               a_dev_buf.GetDeviceBuffer(),
+                                               M * K * sizeof(ADataType),
+                                               hipMemcpyHostToDevice));
+            ck_tile::hip_check_error(hipMemcpy(d_B,
+                                               b_origin_dev_buf.GetDeviceBuffer(),
+                                               N * K * sizeof(BDataType),
+                                               hipMemcpyHostToDevice));
+
+            if constexpr(ScaleGranularityM == -1 && ScaleGranularityN == -1)
+            {
+                ck_tile::reference_gemm_gpu<ADataType,
+                                            BDataType,
+                                            AccDataType,
+                                            CDataType,
+                                            ALayout,
+                                            BLayout,
+                                            CLayout>(
+                    d_A, d_B, d_C, M, N, K, stride_A, stride_B, stride_C);
+            }
+            else
+            {
+                ck_tile::reference_blockwise_gemm_gpu<ADataType,
+                                                      BDataType,
+                                                      AccDataType,
+                                                      CDataType,
+                                                      ALayout,
+                                                      BLayout,
+                                                      CLayout>(
+                    d_A,
+                    d_B,
+                    d_C,
+                    M,
+                    N,
+                    K,
+                    stride_A,
+                    stride_B,
+                    stride_C,
+                    ScaleGranularityM,
+                    ScaleGranularityN,
+                    K,
+                    static_cast<float*>(per_token_scale_dev_buf.GetDeviceBuffer()),
+                    static_cast<float*>(per_channel_scale_dev_buf.GetDeviceBuffer()));
+            }
+
+            ck_tile::hip_check_error(hipMemcpy(c_gpu_ref_dev_buf.GetDeviceBuffer(),
+                                               d_C,
+                                               M * N * sizeof(CDataType),
+                                               hipMemcpyDeviceToHost));
+
+            ck_tile::hip_check_error(hipFree(d_A));
+            ck_tile::hip_check_error(hipFree(d_B));
+            ck_tile::hip_check_error(hipFree(d_C));
+
+            c_gpu_ref_dev_buf.FromDevice(c_gpu_ref_host.data());
+            const float max_accumulated_value =
+                *std::max_element(c_gpu_ref_host.mData.begin(), c_gpu_ref_host.mData.end());
+            const auto rtol_atol =
+                calculate_rtol_atol<ADataType, BDataType, AccDataType, CDataType>(
+                    K, kbatch, max_accumulated_value);
+            pass = ck_tile::check_err(c_rslt_host,
+                                      c_gpu_ref_host,
+                                      "Error: Incorrect results!",
+                                      rtol_atol.at(ck_tile::number<0>{}),
+                                      rtol_atol.at(ck_tile::number<1>{}));
+
+            std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
+                      << " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
+                      << std::endl;
+            std::cout << "The GPU veification result is: " << (pass ? "correct" : "fail")
+                      << std::endl;
+        }
+
+        return pass;
     }
-
-    return pass;
 }
diff --git a/include/ck_tile/core.hpp b/include/ck_tile/core.hpp
index 10dfdd7d28..82ae349706 100644
--- a/include/ck_tile/core.hpp
+++ b/include/ck_tile/core.hpp
@@ -38,6 +38,7 @@
 #include "ck_tile/core/numeric/numeric.hpp"
 #include "ck_tile/core/numeric/pk_fp4.hpp"
 #include "ck_tile/core/numeric/pk_int4.hpp"
+#include "ck_tile/core/numeric/e8m0.hpp"
 #include "ck_tile/core/numeric/type_convert.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
 #include "ck_tile/core/tensor/buffer_view.hpp"