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

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)
 

requirements.txt

@@ -29,4 +29,4 @@ 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@70eefcf4f263aa5c25f3c9ff0db8f6f199ef0fb9 -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
-ROCmSoftwarePlatform/rocMLIR@a48dfb1f163fb0b38369e73e580968b72e85b594 -DBUILD_FAT_LIBROCKCOMPILER=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/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/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,

src/include/migraphx/op/allocate.hpp

@@ -33,6 +33,19 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
+/**
+ * Static allocate:
+ * No inputs: `allocate()`
+ * `this.s` attribute set to the static output shape of the buffer.
+ *
+ * Dynamic allocate:
+ * One input: `allocate(output_dims)`
+ * `output_dims` are the output buffer dimensions and has a static shape.
+ * Either `this.s` or `this.buf_type` must be set to calculate the dynamic output shape at compute
+ * time. If `this.buf_type` is set, the compute_shape() of allocate at compile time will have
+ * dynamic_dimensions from {0, max_int} with rank = output_dims.ndim(). If `this.s` is set then the
+ * compute_shape() will output `this.s`; `this.s` should be a dynamic shape.
+ */
 struct allocate
 {
     shape s{};

src/include/migraphx/op/nonmaxsuppression.hpp

@@ -24,6 +24,7 @@
 #ifndef MIGRAPHX_GUARD_OPERATORS_NONMAXSUPPRESSION_HPP
 #define MIGRAPHX_GUARD_OPERATORS_NONMAXSUPPRESSION_HPP
 
+#include <array>
 #include <cmath>
 #include <queue>
 #include <cstdint>

src/include/migraphx/op/random_uniform.hpp

@@ -77,11 +77,26 @@ struct random_uniform
             using type = typename decltype(output)::value_type;
             if constexpr(std::is_integral<type>{})
             {
-                // default range for all integer types is
-                // (0, std::uniform_int_distribution<type>::max()).
-                // Todo:  enable different ranges
-                std::uniform_int_distribution<type> dis;
-                std::generate(output.begin(), output.end(), [&] { return dis(gen); });
+#ifdef _MSC_VER
+                // According to the C++ specification, the effect is undefined if the result type
+                // for the generator is not one of short, int, long, long long, unsigned short,
+                // unsigned int, unsigned long, or unsigned long long. See
+                // https://en.cppreference.com/w/cpp/numeric/random/uniform_int_distribution.
+                if constexpr(sizeof(type) == 1)
+                {
+                    std::uniform_int_distribution<int> dis{std::numeric_limits<type>::min(),
+                                                           std::numeric_limits<type>::max()};
+                    std::generate(output.begin(), output.end(), [&] { return dis(gen); });
+                }
+                else
+#endif
+                {
+                    // default range for all integer types is
+                    // (0, std::uniform_int_distribution<type>::max()).
+                    // Todo:  enable different ranges
+                    std::uniform_int_distribution<type> dis;
+                    std::generate(output.begin(), output.end(), [&] { return dis(gen); });
+                }
             }
             else
             {