Add smoothquant instance library

example/ck_tile/12_smoothquant/CMakeLists.txt

-set(EXAMPLE_SMOOTHQUANT "tile_example_smoothquant")
+set(TILE_SMOOTHQUANT "tile_smoothquant")
 # not using add_example_executable() to add this target, since we don't want this to have
 # to be included in "make all/install/check"
-message("adding example ${EXAMPLE_SMOOTHQUANT}")
+message("adding ${TILE_SMOOTHQUANT}")
+file(GLOB INSTANCE_SRCS instances/*.cpp)
+add_executable(${TILE_SMOOTHQUANT} EXCLUDE_FROM_ALL smoothquant.cpp)
+target_include_directories(${TILE_SMOOTHQUANT} PRIVATE ${CMAKE_CURRENT_LIST_DIR})
+target_sources(${TILE_SMOOTHQUANT} PRIVATE ${INSTANCE_SRCS})
 
-add_executable(${EXAMPLE_SMOOTHQUANT} EXCLUDE_FROM_ALL example_smoothquant.cpp)
-target_include_directories(${EXAMPLE_SMOOTHQUANT} PRIVATE ${CMAKE_CURRENT_LIST_DIR})
-set(EXAMPLE_SMOOTHQUANT_COMPILE_OPTIONS)
+set(TILE_SMOOTHQUANT_COMPILE_OPTIONS)
 
 # NOTE: we turn off undefined-func-template to let source compile without explicit declare function specializations
-list(APPEND EXAMPLE_SMOOTHQUANT_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal)
+list(APPEND TILE_SMOOTHQUANT_COMPILE_OPTIONS -Wno-undefined-func-template -Wno-float-equal)
+
+target_compile_options(${TILE_SMOOTHQUANT} PRIVATE ${TILE_SMOOTHQUANT_COMPILE_OPTIONS})
 
-target_compile_options(${EXAMPLE_SMOOTHQUANT} PRIVATE ${EXAMPLE_SMOOTHQUANT_COMPILE_OPTIONS})
+set(EXAMPLE_SMOOTHQUANT "tile_example_smoothquant")
+add_executable(${EXAMPLE_SMOOTHQUANT} EXCLUDE_FROM_ALL example_smoothquant.cpp)
+target_compile_options(${EXAMPLE_SMOOTHQUANT} PRIVATE ${TILE_SMOOTHQUANT_COMPILE_OPTIONS})
 
 # TODO: we have to turn off this global prop, otherwise the progress bar generated
 # by cmake will print too many files, execvp: /bin/sh: Argument list too long
 # however, this property may affect global
 # TODO: consider codegen a makefile by us
-set_property(GLOBAL PROPERTY RULE_MESSAGES OFF)
+set_property(GLOBAL PROPERTY RULE_MESSAGES OFF)
\ No newline at end of file

example/ck_tile/12_smoothquant/example_smoothquant.cpp

@@ -8,16 +8,16 @@
 template <typename DataType>
 auto get_elimit()
 {
-    double rtol = 1e-4;
-    double atol = 1e-4;
+    double rtol = 1e-5;
+    double atol = 1e-5;
     return ck_tile::make_tuple(rtol, atol);
 }
 
 template <>
 auto get_elimit<ck_tile::bf16_t>()
 {
-    double rtol = 1e-2;
-    double atol = 1e-2;
+    double rtol = 1e-5;
+    double atol = 1e-5;
     return ck_tile::make_tuple(rtol, atol);
 }
 
@@ -62,7 +62,7 @@ bool run(const ck_tile::ArgParser& arg_parser)
     assert(stride >= n);
 
     using XDataType       = DataType;
-    using XScaleDataType  = DataType;
+    using XScaleDataType  = float;
     using YScaleDataType  = DataType;
     using QYDataType      = ck_tile::int8_t;
     using ComputeDataType = float;
@@ -228,6 +228,10 @@ int main(int argc, char* argv[])
     {
         return run<ck_tile::half_t>(arg_parser) ? 0 : -2;
     }
+    else if(data_type == "bf16")
+    {
+        return run<ck_tile::bf16_t>(arg_parser) ? 0 : -2;
+    }
 
     return -3;
 }

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n1024_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+#if 0
+// template float smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4,  64, 8,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  4,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  8,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1, 16,  4,  64, 1,  true , false>>(const S&, A);
+
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  1, 256, 4,  true , false>>(const S&, A);
+#endif
+
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 1, 2,  128, 8,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 2,  128, 4,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 2,  128, 2,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 1,  true,  false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n1536_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 4,   64, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 2,  128, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1,  256, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 6, 1,  256, 1,  true, false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n2048_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 1, 1,  256, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1,  256, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 8, 1,  256, 1,  true, false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n256_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  4,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n3072_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1,  128, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1,  256, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 6, 1,  256, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 3, 1, 1024, 1,  true, false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n4096_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1,  256, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1, 1024, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1, 1024, 1,  true, false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n4096_tp_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1,  256, 8,  true, true>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t,  1, 4, 1, 1024, 1,  true, true>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n512_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4,  64, 8,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  4,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  8,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n64_n128_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4,  64, 1,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  1,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  2,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_bf16_n768_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  3,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1,  6,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::bf16_t, 1, 12,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n1024_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+#if 0
+// template float smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4,  64, 8,  true ,false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  4,  4,  64, 4,  true ,false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  8,  4,  64, 2,  true ,false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1, 16,  4,  64, 1,  true ,false>>(const S&, A);
+
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  1, 256, 4,  true ,false>>(const S&, A);
+#endif
+
+// template float smoothquant_<trait_<ck_tile::fp16_t,  1, 1, 2,  128, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 2,  128, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 2,  128, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 1,  true, false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n1536_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                       rm  rn  tm  tn  vn  pd   2p
+// template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 4,   64, 8,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 2,  128, 4,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1,  256, 2,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 6, 1,  256, 1,  true,  false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n2048_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm  rn  tm  tn  vn  pd     2p
+// template float smoothquant_<trait_<ck_tile::fp16_t,  1, 1, 1,  256, 8,  true,   false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1,  256, 4,  true,   false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 2,  true,   false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 8, 1,  256, 1,  true,   false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n256_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  4,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n3072_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn  vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1,  128, 8,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1,  256, 4,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 6, 1,  256, 2,  true, false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 3, 1, 1024, 1,  true, false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n4096_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm  rn  tm  tn  vn  pd     2p
+// template float smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1,  256, 8,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 4,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1, 1024, 2,  true,  false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1, 1024, 1,  true,  false>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n4096_tp_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1,  256, 4,  true, true>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 2, 1, 1024, 2,  true, true>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t,  1, 4, 1, 1024, 1,  true, true>>(const S&, A);
+
+// clang-format on

example/ck_tile/12_smoothquant/instances/smoothquant_fp16_n512_instance.cpp

+
+// SPDX-License-Identifier: MIT
+// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
+
+#include "smoothquant_instance_common.hpp"
+
+// clang-format off
+//                                                   rm rn tm  tn   vn   pd    2p
+// template float smoothquant_<trait_<ck_tile::fp16_t, 1,  1,  4,  64, 8,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  2,  4,  64, 4,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  4,  4,  64, 2,  true , false>>(const S&, A);
+template float smoothquant_<trait_<ck_tile::fp16_t, 1,  8,  4,  64, 1,  true , false>>(const S&, A);
+// clang-format on