From 74b2592535ba2c0494bdf8533be7eb045c8ca6c7 Mon Sep 17 00:00:00 2001
From: carlushuang <carlus.huang@amd.com>
Date: Wed, 22 Jan 2025 17:34:27 +0800
Subject: [PATCH] add fp8 as dst (#1830)

[ROCm/composable_kernel commit: 052a72655cc8888199e69498214ad300e00f2b4c]
---
 .../moe_smoothquant_bf16_n1024_instance.cpp   |  13 ++-
 .../moe_smoothquant_bf16_n1536_instance.cpp   |  13 ++-
 .../moe_smoothquant_bf16_n2048_instance.cpp   |  13 ++-
 .../moe_smoothquant_bf16_n256_instance.cpp    |  10 +-
 .../moe_smoothquant_bf16_n3072_instance.cpp   |  12 ++-
 .../moe_smoothquant_bf16_n4096_instance.cpp   |  12 ++-
 ...moe_smoothquant_bf16_n4096_tp_instance.cpp |  12 ++-
 .../moe_smoothquant_bf16_n512_instance.cpp    |  13 ++-
 ...moe_smoothquant_bf16_n64_n128_instance.cpp |  10 +-
 .../moe_smoothquant_bf16_n768_instance.cpp    |  10 +-
 .../moe_smoothquant_fp16_n1024_instance.cpp   |  13 ++-
 .../moe_smoothquant_fp16_n1536_instance.cpp   |  13 ++-
 .../moe_smoothquant_fp16_n2048_instance.cpp   |  12 ++-
 .../moe_smoothquant_fp16_n256_instance.cpp    |  10 +-
 .../moe_smoothquant_fp16_n3072_instance.cpp   |  12 ++-
 .../moe_smoothquant_fp16_n4096_instance.cpp   |  12 ++-
 ...moe_smoothquant_fp16_n4096_tp_instance.cpp |  12 ++-
 .../moe_smoothquant_fp16_n512_instance.cpp    |  13 ++-
 ...moe_smoothquant_fp16_n64_n128_instance.cpp |  10 +-
 .../moe_smoothquant_fp16_n768_instance.cpp    |  10 +-
 .../instances/moe_smoothquant_fwd_api.cpp     | 100 ++++++++++--------
 .../moe_smoothquant_instance_common.hpp       |  19 ++--
 .../14_moe_smoothquant/moe_smoothquant.cpp    |  33 ++++--
 .../14_moe_smoothquant/moe_smoothquant.hpp    |  30 ++----
 .../14_moe_smoothquant/script/smoke_test.sh   |  52 ++++-----
 .../core/utility/unary_element_function.hpp   |  16 +--
 .../reference_rowwise_quantization2d.hpp      |   2 +-
 .../kernel/moe_smoothquant_kernel.hpp         |   3 +-
 .../smoothquant_pipeline_one_pass.hpp         |   2 +-
 .../smoothquant_pipeline_two_pass.hpp         |   2 +-
 30 files changed, 300 insertions(+), 194 deletions(-)

diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1024_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1024_instance.cpp
index f43626147f..39481e2c83 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1024_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1024_instance.cpp
@@ -15,8 +15,13 @@ template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 16,  4,  64, 1,  true
 template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  1, 256, 4,  true, false>>(const S&, A);
 #endif
 
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 1, 2,  128, 8,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 2,  128, 4,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 2,  128, 2,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 1, 2,  128, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 2,  128, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1,  256, 1,  true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 1, 2,  128, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 2,  128, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1,  256, 1,  true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1536_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1536_instance.cpp
index e380520fce..6feccbdaff 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1536_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n1536_instance.cpp
@@ -6,8 +6,13 @@
 
 // clang-format off
 //                                                  rm rn tm  tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 3, 4,  64, 8, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 3, 2, 128, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 3, 1, 256, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 6, 1, 256, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 4,  64, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 2, 128, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 1, 256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 6, 1, 256, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 4,  64, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 2, 128, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 1, 256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 6, 1, 256, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n2048_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n2048_instance.cpp
index 4d536cd61d..0e2c936633 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n2048_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n2048_instance.cpp
@@ -6,9 +6,14 @@
 
 // clang-format off
 //                                                   rm rn tm tn   vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 1, 1, 256, 8, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1, 256, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1, 256, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 8, 1, 256, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 1, 1, 256, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 1, 256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1, 256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 8, 1, 256, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 1, 1, 256, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 1, 256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1, 256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 8, 1, 256, 1, true, false>>(const S&, A);
 
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n256_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n256_instance.cpp
index b38a4733a4..373cb0352b 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n256_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n256_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                  rm  rn tm  tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n3072_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n3072_instance.cpp
index c5c170aef1..c0c778f36c 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n3072_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n3072_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn   vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1,  128, 8, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1,  256, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 6, 1,  256, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1, 1024, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 1,  128, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 1,  256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 6, 1,  256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 3, 1, 1024, 1, true, false>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 1,  128, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 1,  256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 6, 1,  256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 3, 1, 1024, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_instance.cpp
index 0e48a1b691..47cffd5fc2 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn   vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1,  256, 8,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 4,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1, 1024, 2,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1, 1024, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 1,  256, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1,  256, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 1, 1024, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1, 1024, 1,  true, false>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 1,  256, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1,  256, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 1, 1024, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1, 1024, 1,  true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_tp_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_tp_instance.cpp
index 4af42c6c80..726d6018a6 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_tp_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n4096_tp_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn   vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1,  256, 8,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 4,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1, 1024, 2,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1, 1024, 1,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 1,  256, 8,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 4, 1, 1024, 1,  true, true>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 1,  256, 8,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 1, 1024, 1,  true, true>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n512_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n512_instance.cpp
index ea611a1834..6026d509d0 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n512_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n512_instance.cpp
@@ -6,8 +6,13 @@
 
 // clang-format off
 //                                                  rm rn tm tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 1, 4, 64, 8, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 2, 4, 64, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 4, 4, 64, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 8, 4, 64, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t,  1, 1, 4, 64, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t,  1, 2, 4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t,  1, 4, 4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t,  1, 8, 4, 64, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 1, 4, 64, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 2, 4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 4, 4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 8, 4, 64, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n64_n128_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n64_n128_instance.cpp
index a6209820e6..3924662fe5 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n64_n128_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n64_n128_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                  rm rn  tm  tn  vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4, 64, 1,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4, 64, 2,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4, 64, 1,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  1,  4, 64, 1,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  1,  4, 64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  2,  4, 64, 1,  true , false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  1,  4, 64, 1,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  1,  4, 64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  2,  4, 64, 1,  true , false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n768_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n768_instance.cpp
index f569dedf35..00d5c980d7 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n768_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_bf16_n768_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                   rm rn tm tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  3, 4, 64, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1,  6, 4, 64, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::bf16_t, 1, 12, 4, 64, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  3, 4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1,  6, 4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::int8_t, 1, 12, 4, 64, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  3, 4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1,  6, 4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::bf16_t, ck_tile::fp8_t, 1, 12, 4, 64, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1024_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1024_instance.cpp
index 3793adb5c5..c908739efa 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1024_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1024_instance.cpp
@@ -15,8 +15,13 @@ template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 16,  4,  64, 1,  true
 template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  1, 256, 4,  true ,false>>(const S&, A);
 #endif
 
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 1, 2,  128, 8,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 2,  128, 4,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 2,  128, 2,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 1, 2,  128, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 2, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 4, 2,  128, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 4, 1,  256, 1,  true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 1, 2,  128, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 2, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 4, 2,  128, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 4, 1,  256, 1,  true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1536_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1536_instance.cpp
index 4bf9cb1a49..65e9470cde 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1536_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n1536_instance.cpp
@@ -6,8 +6,13 @@
 
 // clang-format off
 //                                                  rm rn tm   tn  vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 3, 4,   64, 8,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 3, 2,  128, 4,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 3, 1,  256, 2,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 6, 1,  256, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 3, 4,   64, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 3, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 3, 1,  256, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 6, 1,  256, 1,  true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 3, 4,   64, 8,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 3, 2,  128, 4,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 3, 1,  256, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 6, 1,  256, 1,  true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n2048_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n2048_instance.cpp
index eb0d0fe103..421352f45f 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n2048_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n2048_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                  rm rn tm  tn  vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 1, 1,  256, 8, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 2, 1,  256, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 4, 1,  256, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 8, 1,  256, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 1, 1,  256, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 2, 1,  256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 4, 1,  256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 8, 1,  256, 1, true, false>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 1, 1,  256, 8, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 2, 1,  256, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 4, 1,  256, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 8, 1,  256, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n256_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n256_instance.cpp
index 36bc0de150..f102cb6d60 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n256_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n256_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                  rm rn  tm  tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  1,  4, 64, 4, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  2,  4, 64, 2, true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  4,  4, 64, 1, true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n3072_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n3072_instance.cpp
index fa6f53b2d4..ad7b9e3d15 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n3072_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n3072_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn  vn  pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1,  128, 8,true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1,  256, 4,true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 6, 1,  256, 2,true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1, 1024, 1,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 3, 1,  128, 8,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 3, 1,  256, 4,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 6, 1,  256, 2,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t,  1, 3, 1, 1024, 1,true, false>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 3, 1,  128, 8,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 3, 1,  256, 4,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 6, 1,  256, 2,true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t,  1, 3, 1, 1024, 1,true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_instance.cpp
index 9b7462ab92..bb79ec7ab4 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn  vn    pd     2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1,  256, 8,  true,  false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 4,  true,  false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1, 1024, 2,  true,  false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1, 1024, 1,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 2, 1,  256, 8,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 4, 1,  256, 4,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 2, 1, 1024, 2,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 4, 1, 1024, 1,  true,  false>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 2, 1,  256, 8,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 4, 1,  256, 4,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 2, 1, 1024, 2,  true,  false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 4, 1, 1024, 1,  true,  false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_tp_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_tp_instance.cpp
index 8911bc2295..766c60689f 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_tp_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n4096_tp_instance.cpp
@@ -6,9 +6,13 @@
 
 // clang-format off
 //                                                   rm rn tm  tn   vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1,  256, 8,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 4,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1, 1024, 2,  true, true>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1, 1024, 1,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 2, 1,  256, 8,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 4, 1, 1024, 1,  true, true>>(const S&, A);
 
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 2, 1,  256, 8,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 4, 1, 1024, 1,  true, true>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n512_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n512_instance.cpp
index 07783ac168..6c24e1ebe0 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n512_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n512_instance.cpp
@@ -6,8 +6,13 @@
 
 // clang-format off
 //                                                  rm rn  tm  tn  vn   pd    2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4, 64, 8,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4, 64, 4,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  4,  4, 64, 2,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  8,  4, 64, 1,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  1,  4, 64, 8,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  2,  4, 64, 4,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  4,  4, 64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  8,  4, 64, 1,  true , false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  1,  4, 64, 8,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  2,  4, 64, 4,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  4,  4, 64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  8,  4, 64, 1,  true , false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n64_n128_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n64_n128_instance.cpp
index a5ab56a76c..df785eefef 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n64_n128_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n64_n128_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                  rm  rn  tm  tn  vn  pd      2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4,  64, 1,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4,  64, 2,  true, false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4,  64, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  1,  4,  64, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  1,  4,  64, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  2,  4,  64, 1,  true, false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  1,  4,  64, 1,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  1,  4,  64, 2,  true, false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  2,  4,  64, 1,  true, false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n768_instance.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n768_instance.cpp
index 4272cbafc6..d89f1c3bbf 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n768_instance.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fp16_n768_instance.cpp
@@ -6,7 +6,11 @@
 
 // clang-format off
 //                                                  rm  rn  tm  tn  vn  pd       2p
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  3,  4,  64, 4,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1,  6,  4,  64, 2,  true , false>>(const S&, A);
-template float moe_smoothquant_<trait_<ck_tile::fp16_t, 1, 12,  4,  64, 1,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  3,  4,  64, 4,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1,  6,  4,  64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::int8_t, 1, 12,  4,  64, 1,  true , false>>(const S&, A);
+
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  3,  4,  64, 4,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1,  6,  4,  64, 2,  true , false>>(const S&, A);
+template float moe_smoothquant_<trait_<ck_tile::fp16_t, ck_tile::fp8_t, 1, 12,  4,  64, 1,  true , false>>(const S&, A);
 // clang-format on
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fwd_api.cpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fwd_api.cpp
index a65d3fde66..9d86c54b1a 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fwd_api.cpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_fwd_api.cpp
@@ -4,7 +4,8 @@
 #include <ck_tile/core.hpp>
 #include "moe_smoothquant.hpp"
 
-template <typename DataType_,
+template <typename InType,
+          typename OutType,
           ck_tile::index_t Repeat_M_,         // each thread repeat along M
           ck_tile::index_t Repeat_N_,         // each thread repeat along N
           ck_tile::index_t ThreadPerBlock_M_, // num threads along M
@@ -12,7 +13,8 @@ template <typename DataType_,
           ck_tile::index_t Vector_N_,         // vector size along N
           bool kPadN_,
           bool kTwoPass_>
-using trait_ = moe_smoothquant_traits_<DataType_,
+using trait_ = moe_smoothquant_traits_<InType,
+                                       OutType,
                                        Repeat_M_,
                                        Repeat_N_,
                                        ThreadPerBlock_M_,
@@ -21,108 +23,108 @@ using trait_ = moe_smoothquant_traits_<DataType_,
                                        kPadN_,
                                        kTwoPass_>;
 
-template <typename data_type>
+template <typename in_type, typename out_type>
 float moe_smoothquant_dispatch(moe_smoothquant_traits /*t*/,
                                moe_smoothquant_args a,
                                const ck_tile::stream_config& s)
 {
     float r = -1;
     // clang-format off
-    //                                         rm  rn  tm  tn  vn   pd    2p
+    //                                                    rm  rn  tm  tn  vn   pd    2p
     if(a.hidden_size <= 64) {
-            r = moe_smoothquant_<trait_<data_type, 1,  1,  4,  64, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type, 1,  1,  4,  64, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 128) {
         if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type, 1,  1,  4,  64, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type, 1,  1,  4,  64, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type, 1,  2,  4,  64, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type, 1,  2,  4,  64, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 256) {
         if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 1,  4,  64, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 1,  4,  64, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2,  4,  64, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2,  4,  64, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 4,  4,  64, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4,  4,  64, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 512) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 1,  4,  64, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 1,  4,  64, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2,  4,  64, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2,  4,  64, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 4,  4,  64, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4,  4,  64, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 8,  4,  64, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 8,  4,  64, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 768) {
         if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3,  4,  64, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3,  4,  64, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 6,  4,  64, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 6,  4,  64, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1,12,  4,  64, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1,12,  4,  64, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 1024) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 1, 2,  128, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 1, 2,  128, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 2,  128, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 2,  128, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 2,  128, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 2,  128, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1,  256, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1,  256, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 1536) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 4,   64, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 4,   64, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 2,  128, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 2,  128, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 1,  256, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 1,  256, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 6, 1,  256, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 6, 1,  256, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 2048) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 1, 1,  256, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 1, 1,  256, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 1,  256, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 1,  256, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1,  256, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1,  256, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 8, 1,  256, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 8, 1,  256, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 3072) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 1,  128, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 1,  128, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 1,  256, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 1,  256, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 6, 1,  256, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 6, 1,  256, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 3, 1, 1024, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 3, 1, 1024, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size <= 4096) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 1,  256, 8,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 1,  256, 8,  true, false>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1,  256, 4,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1,  256, 4,  true, false>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 1, 1024, 2,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 1, 1024, 2,  true, false>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1, 1024, 1,  true, false>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1, 1024, 1,  true, false>>(s, a);
     }
     else if(a.hidden_size > 4096) {
         if (a.hidden_size % 8 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 1,  256, 8,  true, true>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 1,  256, 8,  true, true>>(s, a);
         else if (a.hidden_size % 4 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1,  256, 4,  true, true>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1,  256, 4,  true, true>>(s, a);
         else if (a.hidden_size % 2 == 0)
-            r = moe_smoothquant_<trait_<data_type,  1, 2, 1, 1024, 2,  true, true>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 2, 1, 1024, 2,  true, true>>(s, a);
         else
-            r = moe_smoothquant_<trait_<data_type,  1, 4, 1, 1024, 1,  true, true>>(s, a);
+            r = moe_smoothquant_<trait_<in_type, out_type,  1, 4, 1, 1024, 1,  true, true>>(s, a);
     }
     return r;
     // clang-format on
@@ -132,13 +134,21 @@ float moe_smoothquant(moe_smoothquant_traits t,
                       moe_smoothquant_args a,
                       const ck_tile::stream_config& s)
 {
-    if(t.data_type.compare("fp16") == 0)
+    if(t.in_type.compare("fp16") == 0 && t.out_type == "int8")
     {
-        return moe_smoothquant_dispatch<ck_tile::fp16_t>(t, a, s);
+        return moe_smoothquant_dispatch<ck_tile::fp16_t, ck_tile::int8_t>(t, a, s);
     }
-    else if(t.data_type.compare("bf16") == 0)
+    else if(t.in_type.compare("fp16") == 0 && t.out_type == "fp8")
     {
-        return moe_smoothquant_dispatch<ck_tile::bf16_t>(t, a, s);
+        return moe_smoothquant_dispatch<ck_tile::fp16_t, ck_tile::fp8_t>(t, a, s);
+    }
+    else if(t.in_type.compare("bf16") == 0 && t.out_type == "int8")
+    {
+        return moe_smoothquant_dispatch<ck_tile::bf16_t, ck_tile::int8_t>(t, a, s);
+    }
+    else if(t.in_type.compare("bf16") == 0 && t.out_type == "fp8")
+    {
+        return moe_smoothquant_dispatch<ck_tile::bf16_t, ck_tile::fp8_t>(t, a, s);
     }
     else
         throw std::runtime_error("Without supported instances!");
diff --git a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_instance_common.hpp b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_instance_common.hpp
index 19ee6670a1..885d9ff7bf 100644
--- a/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_instance_common.hpp
+++ b/example/ck_tile/14_moe_smoothquant/instances/moe_smoothquant_instance_common.hpp
@@ -11,7 +11,8 @@
 using S = ck_tile::stream_config;
 using A = moe_smoothquant_args;
 
-template <typename DataType_,
+template <typename InputType_,
+          typename OutputType_,
           ck_tile::index_t Repeat_M_,         // each thread repeat along M
           ck_tile::index_t Repeat_N_,         // each thread repeat along N
           ck_tile::index_t ThreadPerBlock_M_, // num threads along M
@@ -19,7 +20,8 @@ template <typename DataType_,
           ck_tile::index_t Vector_N_,         // vector size along N
           bool kPadN_,
           bool kTwoPass_>
-using trait_ = moe_smoothquant_traits_<DataType_,
+using trait_ = moe_smoothquant_traits_<InputType_,
+                                       OutputType_,
                                        Repeat_M_,
                                        Repeat_N_,
                                        ThreadPerBlock_M_,
@@ -31,14 +33,15 @@ using trait_ = moe_smoothquant_traits_<DataType_,
 template <typename Traits_>
 float moe_smoothquant_(const S& s, A a)
 {
-    using DataType = typename Traits_::DataType;
+    using InputType  = typename Traits_::InputType;
+    using OutputType = typename Traits_::OutputType;
 
     using PipelineProblem = ck_tile::SmoothquantPipelineProblem<
-        typename MoeSmoothquantTypeConfig<DataType>::XDataType,
-        typename MoeSmoothquantTypeConfig<DataType>::SmoothScaleDataType,
-        typename MoeSmoothquantTypeConfig<DataType>::ComputeDataType,
-        typename MoeSmoothquantTypeConfig<DataType>::YScaleDataType,
-        typename MoeSmoothquantTypeConfig<DataType>::QYDataType,
+        typename MoeSmoothquantTypeConfig<InputType, OutputType>::XDataType,
+        typename MoeSmoothquantTypeConfig<InputType, OutputType>::SmoothScaleDataType,
+        typename MoeSmoothquantTypeConfig<InputType, OutputType>::ComputeDataType,
+        typename MoeSmoothquantTypeConfig<InputType, OutputType>::YScaleDataType,
+        typename MoeSmoothquantTypeConfig<InputType, OutputType>::QYDataType,
         typename Traits_::Shape,
         Traits_::kPadN,
         Traits_::kTwoPass>;
diff --git a/example/ck_tile/14_moe_smoothquant/moe_smoothquant.cpp b/example/ck_tile/14_moe_smoothquant/moe_smoothquant.cpp
index 56daf4d523..dc5b397c85 100644
--- a/example/ck_tile/14_moe_smoothquant/moe_smoothquant.cpp
+++ b/example/ck_tile/14_moe_smoothquant/moe_smoothquant.cpp
@@ -63,7 +63,8 @@ auto create_args(int argc, char* argv[])
         .insert("stride", "-1", "stride per row, if -1 then equal to hidden_size")
         .insert("v", "1", "cpu validation or not")
         .insert("kname", "1", "print kernel name or not")
-        .insert("prec", "fp16", "precision")
+        .insert("prec_i", "fp16", "input precision, fp16/bf16")
+        .insert("prec_o", "int8", "precision, int8/fp8")
         .insert("warmup", "5", "cold iter")
         .insert("repeat", "20", "hot iter");
 
@@ -71,7 +72,7 @@ auto create_args(int argc, char* argv[])
     return std::make_tuple(result, arg_parser);
 }
 
-template <typename DataType>
+template <typename InputType, typename OutputType>
 bool run(const ck_tile::ArgParser& arg_parser)
 {
     ck_tile::index_t tokens      = arg_parser.get_int("t");
@@ -81,7 +82,8 @@ bool run(const ck_tile::ArgParser& arg_parser)
         stride = hidden_size;
     ck_tile::index_t experts = arg_parser.get_int("e");
     ck_tile::index_t topk    = arg_parser.get_int("k");
-    std::string data_type    = arg_parser.get_str("prec");
+    std::string prec_i       = arg_parser.get_str("prec_i");
+    std::string prec_o       = arg_parser.get_str("prec_o");
     int kname                = arg_parser.get_int("kname");
     int do_validation        = arg_parser.get_int("v");
     int warmup               = arg_parser.get_int("warmup");
@@ -89,7 +91,7 @@ bool run(const ck_tile::ArgParser& arg_parser)
 
     assert(stride >= hidden_size);
 
-    using TypeConfig = MoeSmoothquantTypeConfig<DataType>;
+    using TypeConfig = MoeSmoothquantTypeConfig<InputType, OutputType>;
 
     using XDataType           = typename TypeConfig::XDataType;
     using SmoothScaleDataType = typename TypeConfig::SmoothScaleDataType;
@@ -122,11 +124,11 @@ bool run(const ck_tile::ArgParser& arg_parser)
     smscale_buf.ToDevice(smscale_host.data());
     topk_ids_buf.ToDevice(topk_ids_host.data());
 
-    std::cout << "[" << data_type << "]"
+    std::cout << "[" << prec_i << "-" << prec_o << "]"
               << " tokens:" << tokens << ", hidden_size:" << hidden_size << ", stride:" << stride
               << ", experts:" << experts << ", topk:" << topk << std::flush;
 
-    moe_smoothquant_traits traits{data_type};
+    moe_smoothquant_traits traits{prec_i, prec_o};
 
     moe_smoothquant_args args{x_buf.GetDeviceBuffer(),
                               smscale_buf.GetDeviceBuffer(),
@@ -251,14 +253,23 @@ int main(int argc, char* argv[])
     if(!result)
         return -1;
 
-    const std::string data_type = arg_parser.get_str("prec");
-    if(data_type == "fp16")
+    const std::string prec_i = arg_parser.get_str("prec_i");
+    const std::string prec_o = arg_parser.get_str("prec_o");
+    if(prec_i == "fp16" && prec_o == "int8")
     {
-        return run<ck_tile::half_t>(arg_parser) ? 0 : -2;
+        return run<ck_tile::half_t, ck_tile::int8_t>(arg_parser) ? 0 : -2;
     }
-    else if(data_type == "bf16")
+    else if(prec_i == "fp16" && prec_o == "fp8")
     {
-        return run<ck_tile::bf16_t>(arg_parser) ? 0 : -2;
+        return run<ck_tile::half_t, ck_tile::fp8_t>(arg_parser) ? 0 : -2;
+    }
+    else if(prec_i == "bf16" && prec_o == "int8")
+    {
+        return run<ck_tile::bf16_t, ck_tile::int8_t>(arg_parser) ? 0 : -2;
+    }
+    else if(prec_i == "bf16" && prec_o == "fp8")
+    {
+        return run<ck_tile::bf16_t, ck_tile::fp8_t>(arg_parser) ? 0 : -2;
     }
 
     return -3;
diff --git a/example/ck_tile/14_moe_smoothquant/moe_smoothquant.hpp b/example/ck_tile/14_moe_smoothquant/moe_smoothquant.hpp
index b4c738ef36..c1b90b14b2 100644
--- a/example/ck_tile/14_moe_smoothquant/moe_smoothquant.hpp
+++ b/example/ck_tile/14_moe_smoothquant/moe_smoothquant.hpp
@@ -8,26 +8,13 @@
 #include "ck_tile/ops/smoothquant.hpp"
 #include <string>
 
-template <typename DataType>
-struct MoeSmoothquantTypeConfig;
-
-template <>
-struct MoeSmoothquantTypeConfig<ck_tile::half_t>
+template <typename InputType, typename OutputType>
+struct MoeSmoothquantTypeConfig
 {
-    using XDataType           = ck_tile::half_t;
+    using XDataType           = InputType;
     using SmoothScaleDataType = float;
     using YScaleDataType      = float;
-    using QYDataType          = ck_tile::int8_t;
-    using ComputeDataType     = float;
-};
-
-template <>
-struct MoeSmoothquantTypeConfig<ck_tile::bf16_t>
-{
-    using XDataType           = ck_tile::bf16_t;
-    using SmoothScaleDataType = float;
-    using YScaleDataType      = float;
-    using QYDataType          = ck_tile::int8_t;
+    using QYDataType          = OutputType;
     using ComputeDataType     = float;
 };
 
@@ -37,7 +24,8 @@ struct moe_smoothquant_args : public ck_tile::MoeSmoothquantHostArgs
 };
 
 // this is used to pattern-match internl kernel implementation, not to instantiate kernel
-template <typename DataType_,
+template <typename InputType_,
+          typename OutputType_,
           ck_tile::index_t Repeat_M_,         // each thread repeat along M
           ck_tile::index_t Repeat_N_,         // each thread repeat along N
           ck_tile::index_t ThreadPerBlock_M_, // num threads along M
@@ -47,7 +35,8 @@ template <typename DataType_,
           bool kTwoPass_>
 struct moe_smoothquant_traits_
 {
-    using DataType = ck_tile::remove_cvref_t<DataType_>;
+    using InputType  = ck_tile::remove_cvref_t<InputType_>;
+    using OutputType = ck_tile::remove_cvref_t<OutputType_>;
 
     static constexpr bool is_warp_per_row = ThreadPerBlock_N_ <= warpSize;
     static_assert((ThreadPerBlock_M_ * ThreadPerBlock_N_) % warpSize == 0);
@@ -108,7 +97,8 @@ float moe_smoothquant_(const ck_tile::stream_config& s, moe_smoothquant_args a);
 // This is the public API, will be generated by script
 struct moe_smoothquant_traits
 {
-    std::string data_type;
+    std::string in_type;  // input type
+    std::string out_type; // output type
 };
 
 float moe_smoothquant(moe_smoothquant_traits, moe_smoothquant_args, const ck_tile::stream_config&);
diff --git a/example/ck_tile/14_moe_smoothquant/script/smoke_test.sh b/example/ck_tile/14_moe_smoothquant/script/smoke_test.sh
index 3bb62d37b9..e01f3de10a 100755
--- a/example/ck_tile/14_moe_smoothquant/script/smoke_test.sh
+++ b/example/ck_tile/14_moe_smoothquant/script/smoke_test.sh
@@ -2,29 +2,31 @@
 EXE=build/bin/tile_example_moe_smoothquant
 
 for pr_i in "fp16" "bf16" ; do
-$EXE -prec=$pr_i -t=99  -h=13
-$EXE -prec=$pr_i -t=17  -h=16
-$EXE -prec=$pr_i -t=1   -h=100
-$EXE -prec=$pr_i -t=4   -h=128
-$EXE -prec=$pr_i -t=80  -h=127
-$EXE -prec=$pr_i -t=22  -h=255 -stride=256
-$EXE -prec=$pr_i -t=7   -h=599
-$EXE -prec=$pr_i -t=19  -h=512
-$EXE -prec=$pr_i -t=33  -h=313 -stride=1000
-$EXE -prec=$pr_i -t=11  -h=510
-$EXE -prec=$pr_i -t=171 -h=676 -stride=818
-$EXE -prec=$pr_i -t=91  -h=636
-$EXE -prec=$pr_i -t=12  -h=768 -stride=800
-$EXE -prec=$pr_i -t=100 -h=766 -stride=812
-$EXE -prec=$pr_i -t=31  -h=1024
-$EXE -prec=$pr_i -t=64  -h=1000 -stride=1004
-$EXE -prec=$pr_i -t=8   -h=1501
-$EXE -prec=$pr_i -t=3   -h=1826
-$EXE -prec=$pr_i -t=5   -h=2040
-$EXE -prec=$pr_i -t=7   -h=2734
-$EXE -prec=$pr_i -t=1   -h=3182
-$EXE -prec=$pr_i -t=9   -h=4096
-$EXE -prec=$pr_i -t=3   -h=8192
-$EXE -prec=$pr_i -t=1   -h=10547
-$EXE -prec=$pr_i -t=3   -h=17134
+for pr_o in "int8" "fp8" ; do
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=99  -h=13
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=17  -h=16
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=1   -h=100
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=4   -h=128
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=80  -h=127
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=22  -h=255 -stride=256
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=7   -h=599
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=19  -h=512
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=33  -h=313 -stride=1000
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=11  -h=510
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=171 -h=676 -stride=818
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=91  -h=636
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=12  -h=768 -stride=800
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=100 -h=766 -stride=812
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=31  -h=1024
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=64  -h=1000 -stride=1004
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=8   -h=1501
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=3   -h=1826
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=5   -h=2040
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=7   -h=2734
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=1   -h=3182
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=9   -h=4096
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=3   -h=8192
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=1   -h=10547
+$EXE -prec_i=$pr_i -prec_o=$pr_o -t=3   -h=17134
+done
 done
diff --git a/include/ck_tile/core/utility/unary_element_function.hpp b/include/ck_tile/core/utility/unary_element_function.hpp
index e24c0dc85c..ed3b464660 100644
--- a/include/ck_tile/core/utility/unary_element_function.hpp
+++ b/include/ck_tile/core/utility/unary_element_function.hpp
@@ -51,16 +51,18 @@ struct composes<F>
 template <typename... Ts>
 __host__ __device__ composes(Ts&&...)->composes<remove_cvref_t<Ts>...>;
 
-template <typename To>
+template <typename SaturateType>
 struct saturates
 {
-    template <typename From>
-    CK_TILE_HOST_DEVICE constexpr auto operator()(const From& from) const
-        -> std::enable_if_t<std::is_arithmetic_v<From>, From>
+    // NOTE: this function does not return SaturateType value
+    // it is user's responsiblity to do further cast or not
+    template <typename AccType>
+    CK_TILE_HOST_DEVICE constexpr auto operator()(const AccType& a_) const
+        -> std::enable_if_t<std::is_arithmetic_v<AccType>, AccType>
     {
-        return clamp(from,
-                     type_convert<From>(numeric<To>::lowest()),
-                     type_convert<From>(numeric<To>::max()));
+        return clamp(a_,
+                     type_convert<AccType>(numeric<SaturateType>::lowest()),
+                     type_convert<AccType>(numeric<SaturateType>::max()));
     }
 };
 
diff --git a/include/ck_tile/host/reference/reference_rowwise_quantization2d.hpp b/include/ck_tile/host/reference/reference_rowwise_quantization2d.hpp
index e9a398876f..aff5e78ff0 100644
--- a/include/ck_tile/host/reference/reference_rowwise_quantization2d.hpp
+++ b/include/ck_tile/host/reference/reference_rowwise_quantization2d.hpp
@@ -22,7 +22,7 @@ CK_TILE_HOST void reference_rowwise_quantization2d(const HostTensor<XDataType>&
             // scale = amax / 127 for int8
             auto v_scale = type_convert<XDataType>(scale_m(m));
             auto v_qx    = v_x / v_scale;
-            qx_m_n(m, n) = saturates<QXDataType>{}(v_qx);
+            qx_m_n(m, n) = type_convert<QXDataType>(saturates<QXDataType>{}(v_qx));
         }
     };
 
diff --git a/include/ck_tile/ops/smoothquant/kernel/moe_smoothquant_kernel.hpp b/include/ck_tile/ops/smoothquant/kernel/moe_smoothquant_kernel.hpp
index 62155ebf20..cb934c6c52 100644
--- a/include/ck_tile/ops/smoothquant/kernel/moe_smoothquant_kernel.hpp
+++ b/include/ck_tile/ops/smoothquant/kernel/moe_smoothquant_kernel.hpp
@@ -101,6 +101,7 @@ struct MoeSmoothquant
     template <> struct t2s<ck_tile::bf16_t> { static constexpr const char * name = "bf16"; };
     template <> struct t2s<ck_tile::fp8_t> { static constexpr const char * name = "fp8"; };
     template <> struct t2s<ck_tile::bf8_t> { static constexpr const char * name = "bf8"; };
+    template <> struct t2s<ck_tile::int8_t> { static constexpr const char * name = "i8"; };
     // clang-format on
 
     // in byte
@@ -118,7 +119,7 @@ struct MoeSmoothquant
 
         #define _SS_  std::string
         #define _TS_  std::to_string
-        return _SS_("moe_smoothquant_") + _SS_(t2s<XDataType>::name) + "_" +
+        return _SS_("moe_smoothquant_") + _SS_(t2s<XDataType>::name) + "_" +  _SS_(t2s<QYDataType>::name) + "_" +
              _TS_(S_::Block_M) + "x" + _TS_(S_::Block_N) + "_" + _TS_(S_::WarpPerBlock_M) + "x" + _TS_(S_::WarpPerBlock_N) + "_" +
              _TS_(S_::Warp_M) + "x" + _TS_(S_::Warp_N) + "_" + _TS_(S_::Vector_M) + "x" + _TS_(S_::Vector_N) + "_" +
              _SS_(Pipeline::name) + surfix;
diff --git a/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp b/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp
index 41111c0c12..f45afe3d2a 100644
--- a/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp
+++ b/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_one_pass.hpp
@@ -113,7 +113,7 @@ struct SmoothquantPipelineOnePass
         sweep_tile(qy, [&](auto idx) {
             constexpr auto i_idx = make_tuple(idx[number<0>{}]);
             auto qy_             = y[idx] / yscale[i_idx];
-            qy(idx)              = saturates<QYDataType>{}(qy_);
+            qy(idx)              = type_convert<QYDataType>(saturates<QYDataType>{}(qy_));
         });
         store_tile(qy_window, qy);
     }
diff --git a/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_two_pass.hpp b/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_two_pass.hpp
index 65b515d61f..ba9c6374f1 100644
--- a/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_two_pass.hpp
+++ b/include/ck_tile/ops/smoothquant/pipeline/smoothquant_pipeline_two_pass.hpp
@@ -136,7 +136,7 @@ struct SmoothquantPipelineTwoPass
             sweep_tile(qy, [&](auto idx) {
                 constexpr auto i_idx = make_tuple(idx[number<0>{}]);
                 auto qy_             = y[idx] / yscale[i_idx];
-                qy(idx)              = saturates<QYDataType>{}(qy_);
+                qy(idx)              = type_convert<QYDataType>(saturates<QYDataType>{}(qy_));
             });
             store_tile(qy_window, qy);