Merge remote-tracking branch 'origin/ck-flash-attn' into gemm-perf

CMakeLists.txt

@@ -53,6 +53,12 @@ include(CTest)
 find_package(ROCM REQUIRED)
 find_package(Threads REQUIRED)
 
+if(WIN32)
+option(MIGRAPHX_ENABLE_PYTHON "Enable python bindings" OFF)
+else()
+option(MIGRAPHX_ENABLE_PYTHON "Enable python bindings" ON)
+endif()
+
 find_path(HALF_INCLUDE_DIR half.hpp PATH_SUFFIXES half)
 if (NOT HALF_INCLUDE_DIR)
     message(FATAL_ERROR "Could not find half.hpp - Please check that the install path of half.hpp has been added to CMAKE_PREFIX_PATH")
@@ -261,8 +267,6 @@ rocm_enable_cppcheck(
         MIGRAPHX_USE_CLANG_TIDY
 )
 
-enable_testing()
-
 include(ROCMCreatePackage)
 include(ROCMTest)
 

cmake/Embed.cmake

@@ -77,16 +77,17 @@ function(generate_embed_source EMBED_NAME)
         list(GET PARSE_FILES ${idx} FILE)
 
         set(START_SYMBOL "_binary_${SYMBOL}_start")
-        set(END_SYMBOL "_binary_${SYMBOL}_end")
+        set(LENGTH_SYMBOL "_binary_${SYMBOL}_length")
         if(EMBED_USE_LD)
             string(APPEND EXTERNS "
-                extern const char ${START_SYMBOL}[];
-                extern const char ${END_SYMBOL}[];
+extern const char ${START_SYMBOL}[];
+extern const size_t _binary_${SYMBOL}_size;
+const auto ${LENGTH_SYMBOL} = reinterpret_cast<size_t>(&_binary_${SYMBOL}_size);
             ")
         else()
             string(APPEND EXTERNS "
-                extern const char ${START_SYMBOL}[];
-                extern const char* ${END_SYMBOL};
+extern const char ${START_SYMBOL}[];
+extern const size_t ${LENGTH_SYMBOL};
             ")
         endif()
 
@@ -97,23 +98,22 @@ function(generate_embed_source EMBED_NAME)
         endif()
 
         string(APPEND INIT_KERNELS "
-            { \"${BASE_NAME}\", { ${START_SYMBOL}, ${END_SYMBOL}} },
-        ")
+            { \"${BASE_NAME}\", { ${START_SYMBOL}, ${LENGTH_SYMBOL}} },")
     endforeach()
 
     file(WRITE "${PARSE_HEADER}" "
+#include <string_view>
 #include <unordered_map>
-#include <string>
 #include <utility>
-const std::unordered_map<std::string, std::pair<const char*,const char*>>& ${EMBED_NAME}();
+std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}();
 ")
 
     file(WRITE "${PARSE_SRC}" "
 #include <${EMBED_NAME}.hpp>
 ${EXTERNS}
-const std::unordered_map<std::string, std::pair<const char*,const char*>>& ${EMBED_NAME}()
+std::unordered_map<std::string_view, std::string_view> ${EMBED_NAME}()
 {
-    static const std::unordered_map<std::string, std::pair<const char*,const char*>> result = {${INIT_KERNELS}};
+    static std::unordered_map<std::string_view, std::string_view> result = {${INIT_KERNELS}};
     return result;
 }
 ")
@@ -154,9 +154,10 @@ function(embed_file OUTPUT_FILE OUTPUT_SYMBOL FILE)
             # removes trailing comma
             string(REGEX REPLACE ", $" "" ARRAY_VALUES ${ARRAY_VALUES})
             file(WRITE "${OUT_FILE}" "
-                extern const char _binary_${SYMBOL}_start[] = { ${ARRAY_VALUES} };
-                extern const char* _binary_${SYMBOL}_end = _binary_${SYMBOL}_start + sizeof(_binary_${SYMBOL}_start);
-            \n")
+#include <cstddef>
+extern const char _binary_${SYMBOL}_start[] = { ${ARRAY_VALUES} };
+extern const size_t _binary_${SYMBOL}_length = sizeof(_binary_${SYMBOL}_start);
+")
         endif()
     endforeach()
 endfunction()

requirements.txt

@@ -28,5 +28,5 @@ ROCmSoftwarePlatform/half@rocm-5.6.0
 pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build
 msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off
 sqlite3@3.17 -DCMAKE_POSITION_INDEPENDENT_CODE=On
-ROCmSoftwarePlatform/composable_kernel@4b0b327b81668978249fd9b6eb1c35214e7d78ea -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
-ROCmSoftwarePlatform/rocMLIR@a48dfb1f163fb0b38369e73e580968b72e85b594 -DBUILD_FAT_LIBROCKCOMPILER=On
+ROCmSoftwarePlatform/composable_kernel@70eefcf4f263aa5c25f3c9ff0db8f6f199ef0fb9 -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
+ROCmSoftwarePlatform/rocMLIR@12748a3402c069f733ea7f2ba1f8d8a070b3622a -DBUILD_FAT_LIBROCKCOMPILER=On
\ No newline at end of file

src/CMakeLists.txt

@@ -282,7 +282,9 @@ add_subdirectory(driver)
 add_subdirectory(onnx)
 add_subdirectory(tf)
 
+if(MIGRAPHX_ENABLE_PYTHON)
 add_subdirectory(py)
+endif()
 add_subdirectory(targets/ref)
 target_link_libraries(migraphx_all_targets INTERFACE migraphx_ref)
 if(MIGRAPHX_ENABLE_CPU)

src/api/CMakeLists.txt

@@ -32,6 +32,10 @@ migraphx_generate_export_header(migraphx_c DIRECTORY migraphx/api)
 # bumped when binary compatibility is broken. 
 rocm_set_soversion(migraphx_c 3.0)   
 
+if(BUILD_TESTING)
+    target_compile_definitions(migraphx_c PRIVATE MIGRAPHX_BUILD_TESTING)
+endif()
+
 rocm_clang_tidy_check(migraphx_c)
 target_link_libraries(migraphx_c PRIVATE migraphx migraphx_tf migraphx_onnx)
 

src/api/api.cpp

@@ -38,26 +38,32 @@
 #include <migraphx/register_op.hpp>
 #include <migraphx/json.hpp>
 #include <migraphx/convert_to_json.hpp>
+#include <array>
 #include <algorithm>
 #include <cstdarg>
+
 namespace migraphx {
 
+#ifdef MIGRAPHX_BUILD_TESTING
 static thread_local bool disable_exception_catch = false; // NOLINT
 
 extern "C" MIGRAPHX_C_EXPORT void migraphx_test_private_disable_exception_catch(bool b)
 {
     disable_exception_catch = b;
 }
+#endif
 
 template <class F>
 migraphx_status try_(F f, bool output = true) // NOLINT
 {
+#ifdef MIGRAPHX_BUILD_TESTING
     if(disable_exception_catch)
     {
         f();
     }
     else
     {
+#endif
         try
         {
             f();
@@ -81,7 +87,9 @@ migraphx_status try_(F f, bool output = true) // NOLINT
         {
             return migraphx_status_unknown_error;
         }
+#ifdef MIGRAPHX_BUILD_TESTING
     }
+#endif
     return migraphx_status_success;
 }
 

src/api/include/migraphx/migraphx.h

@@ -26,6 +26,7 @@
 
 #include <stdlib.h>
 #include <stdbool.h>
+#include <stdint.h>
 
 #include <migraphx/api/export.h>
 

src/api/include/migraphx/migraphx.hpp

@@ -66,7 +66,7 @@ template <class PrivateMigraphTypeNameProbe>
 std::string compute_type_name()
 {
     std::string name;
-#ifdef _MSC_VER
+#if defined(_MSC_VER) && !defined(__clang__)
     name = typeid(PrivateMigraphTypeNameProbe).name();
     name = name.substr(7);
 #else

src/compile_src.cpp

@@ -46,7 +46,7 @@ std::vector<char> src_compiler::compile(const std::vector<src_file>& srcs) const
         fs::path full_path   = td.path / src.path;
         fs::path parent_path = full_path.parent_path();
         fs::create_directories(parent_path);
-        write_buffer(full_path.string(), src.content.first, src.len());
+        write_buffer(full_path.string(), src.content.data(), src.content.size());
         if(src.path.extension().string() == ".cpp")
         {
             params += " " + src.path.filename().string();

src/driver/CMakeLists.txt

@@ -48,7 +48,12 @@ rocm_clang_tidy_check(driver)
 file(STRINGS "${CMAKE_SOURCE_DIR}/test/onnx/.onnxrt-commit" String_output)
 target_compile_definitions(driver PUBLIC MIGRAPHX_ORT_SHA1="${String_output}")
 
-target_link_libraries(driver migraphx_all_targets migraphx_onnx migraphx_tf migraphx_py)
+target_link_libraries(driver migraphx_all_targets migraphx_onnx migraphx_tf)
+
+if(MIGRAPHX_ENABLE_PYTHON)
+    target_link_libraries(driver migraphx_py)
+    target_compile_definitions(driver PRIVATE MIGRAPHX_ENABLE_PYTHON)
+endif()
 
 rocm_install_targets(
   TARGETS driver

src/driver/main.cpp

@@ -32,7 +32,9 @@
 
 #include <migraphx/tf.hpp>
 #include <migraphx/onnx.hpp>
+#ifdef MIGRAPHX_ENABLE_PYTHON
 #include <migraphx/py.hpp>
+#endif
 #include <migraphx/stringutils.hpp>
 #include <migraphx/convert_to_json.hpp>
 #include <migraphx/load_save.hpp>
@@ -281,10 +283,12 @@ struct loader
                 options.format = "json";
                 p              = migraphx::load(file, options);
             }
+#ifdef MIGRAPHX_ENABLE_PYTHON
             else if(file_type == "py")
             {
                 p = migraphx::load_py(file);
             }
+#endif
             else if(file_type == "migraphx")
             {
                 p = migraphx::load(file);

src/include/migraphx/auto_register.hpp

@@ -62,10 +62,9 @@ const int auto_register<Action, T>::static_register = auto_register_action<Actio
 #define MIGRAPHX_AUTO_REGISTER_NAME_DETAIL(x) migraphx_auto_register_##x
 #define MIGRAPHX_AUTO_REGISTER_NAME(x) MIGRAPHX_AUTO_REGISTER_NAME_DETAIL(x)
 // NOLINTNEXTLINE
-#define MIGRAPHX_AUTO_REGISTER(...)                                                        \
-    void MIGRAPHX_AUTO_REGISTER_NAME(__LINE__)(migraphx::auto_register<__VA_ARGS__> x =    \
-                                                   migraphx::auto_register<__VA_ARGS__>{}) \
-        __attribute__((unused));
+#define MIGRAPHX_AUTO_REGISTER(...)                              \
+    [[maybe_unused]] void MIGRAPHX_AUTO_REGISTER_NAME(__LINE__)( \
+        migraphx::auto_register<__VA_ARGS__> x = migraphx::auto_register<__VA_ARGS__>{});
 
 } // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/include/migraphx/compile_src.hpp

@@ -37,8 +37,18 @@ inline namespace MIGRAPHX_INLINE_NS {
 struct src_file
 {
     fs::path path;
-    std::pair<const char*, const char*> content;
-    std::size_t len() const { return content.second - content.first; }
+    std::string_view content;
+
+    src_file() = default;
+    src_file(fs::path file_path, std::string_view file_content)
+        : path{std::move(file_path)}, content{file_content}
+    {
+    }
+
+    explicit src_file(const std::pair<std::string_view, std::string_view>& pair)
+        : path{pair.first}, content{pair.second}
+    {
+    }
 };
 
 struct MIGRAPHX_EXPORT src_compiler

src/include/migraphx/config.hpp

@@ -25,6 +25,7 @@
 #define MIGRAPHX_GUARD_CONFIG_HPP
 
 #include <migraphx/export.h>
+#include <ciso646>
 
 #if !defined(MIGRAPHX_USE_CLANG_TIDY) && !defined(DOXYGEN)
 

src/include/migraphx/filesystem.hpp

@@ -29,6 +29,17 @@
 #if defined(CPPCHECK)
 #define MIGRAPHX_HAS_FILESYSTEM 1
 #define MIGRAPHX_HAS_FILESYSTEM_TS 1
+#elif defined(_WIN32)
+#if _MSC_VER >= 1920
+#define MIGRAPHX_HAS_FILESYSTEM 1
+#define MIGRAPHX_HAS_FILESYSTEM_TS 0
+#elif _MSC_VER >= 1900
+#define MIGRAPHX_HAS_FILESYSTEM 0
+#define MIGRAPHX_HAS_FILESYSTEM_TS 1
+#else
+#define MIGRAPHX_HAS_FILESYSTEM 0
+#define MIGRAPHX_HAS_FILESYSTEM_TS 0
+#endif
 #elif defined(__has_include)
 #if __has_include(<filesystem>) && __cplusplus >= 201703L
 #define MIGRAPHX_HAS_FILESYSTEM 1

src/include/migraphx/float_equal.hpp

@@ -27,9 +27,6 @@
 #include <algorithm>
 #include <cmath>
 #include <numeric>
-#ifdef _MSC_VER
-#include <iso646.h>
-#endif
 
 #include <migraphx/requires.hpp>
 #include <migraphx/config.hpp>

src/include/migraphx/generate.hpp

@@ -48,7 +48,7 @@ constexpr T normalize(unsigned long z)
 template <class T, MIGRAPHX_REQUIRES(is_signed<T>{} and not is_floating_point<T>{})>
 constexpr T normalize(unsigned long z)
 {
-    const auto max      = 1UL << (sizeof(T) * 5);
+    const auto max      = 1ULL << (sizeof(T) * 5);
     const auto half_max = max / 2;
     return half_max - (z % max);
 }
@@ -58,7 +58,7 @@ template <class T,
                             not std::is_same<T, bool>{})>
 constexpr T normalize(unsigned long z)
 {
-    const auto max = 1UL << (sizeof(T) * 5);
+    const auto max = 1ULL << (sizeof(T) * 5);
     return z % max;
 }
 

src/include/migraphx/instruction_ref.hpp

 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -27,12 +27,42 @@
 #include <list>
 #include <functional>
 #include <migraphx/config.hpp>
+#include <migraphx/requires.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 
 struct instruction;
+#if defined(_WIN32) && !defined(NDEBUG)
+struct instruction_ref : std::list<instruction>::iterator
+{
+    using instruction_iter       = std::list<instruction>::iterator;
+    using instruction_const_iter = std::list<instruction>::const_iterator;
+
+    instruction_ref() = default;
+    instruction_ref(const instruction_iter& other) : instruction_iter(other) {}
+
+    template <class T,
+              class U,
+              MIGRAPHX_REQUIRES(std::is_same<T, instruction_ref>{} or
+                                std::is_same<U, instruction_ref>{})>
+    friend bool operator==(const T& x, const U& y)
+    {
+        return x._Unwrapped()._Ptr == y._Unwrapped()._Ptr;
+    }
+
+    template <class T,
+              class U,
+              MIGRAPHX_REQUIRES(std::is_same<T, instruction_ref>{} or
+                                std::is_same<U, instruction_ref>{})>
+    friend bool operator!=(const T& x, const U& y)
+    {
+        return !(x == y);
+    }
+};
+#else
 using instruction_ref = std::list<instruction>::iterator;
+#endif
 
 MIGRAPHX_EXPORT migraphx::instruction* as_address(const instruction_ref& ins) noexcept;
 
@@ -65,4 +95,8 @@ struct equal_to<migraphx::instruction_ref> // NOLINT
 
 } // namespace std
 
+#ifdef _MSC_VER
+#include <migraphx/instruction.hpp>
+#endif
+
 #endif

src/include/migraphx/matcher.hpp

@@ -33,6 +33,7 @@
 #include <migraphx/type_name.hpp>
 #include <migraphx/source_location.hpp>
 #include <migraphx/config.hpp>
+#include <array>
 #include <unordered_map>
 #include <unordered_set>
 

src/include/migraphx/normalize_attributes.hpp

@@ -52,6 +52,7 @@ using dependent_type = typename select_dependent_type<T, Ts...>::type;
  * \param attr_val the normalize_axes attributes from the operator
  * \param prefix error message prefix
  */
+MIGRAPHX_EXPORT
 std::vector<int64_t> normalize_axes(const std::vector<int64_t>& axes,
                                     const shape& input_shape,
                                     const value& attr_val,
@@ -67,6 +68,7 @@ std::vector<int64_t> normalize_axes(const std::vector<int64_t>& axes,
  * \param attr_val the normalize_axes attributes from the operator
  * \param prefix error message prefix
  */
+MIGRAPHX_EXPORT
 std::vector<int64_t> normalize_indices(const std::vector<int64_t>& indices,
                                        const std::vector<int64_t>& axes,
                                        const shape& input_shape,