Merge branch 'develop' into make_generate_script

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

@@ -29,4 +29,4 @@ pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build
 msgpack/[email protected] -DMSGPACK_BUILD_TESTS=Off
 sqlite3@3.17 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCmSoftwarePlatform/composable_kernel@a22e479b8e1557961039db2d5c5ff89cff35e86b -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On
-ROCmSoftwarePlatform/rocMLIR@a48dfb1f163fb0b38369e73e580968b72e85b594 -DBUILD_FAT_LIBROCKCOMPILER=On
+ROCmSoftwarePlatform/rocMLIR@12748a3402c069f733ea7f2ba1f8d8a070b3622a -DBUILD_FAT_LIBROCKCOMPILER=On

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/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/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/pad_calc.hpp

@@ -64,6 +64,7 @@ shape compute_padded_shape(const shape& input,
 
 // Used for dynamic auto padding of pooling operators where padding needs to be computed at
 // evaulation time.
+MIGRAPHX_EXPORT
 shape compute_padded_pool_shape(const shape& input,
                                 const shape& kernel,
                                 const std::vector<std::size_t>& padding,

src/onnx/broadcast_qdq.cpp

@@ -0,0 +1,76 @@
+/*
+ * The MIT License (MIT)
+ *
+ * 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+#include <migraphx/onnx/broadcast_qdq.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+// This method is to prep for quantizelinear or dequantizelinear operation for
+// either the broadcasting of weight-scale or zero-points of qlinearadd operator
+// outputs: operator op (inputs x, broadcasted: scale (float) & zero_pt (8-bit))
+instruction_ref bcast_qdq_instr(const std::string& op_name,
+                                instruction_ref x_in,
+                                instruction_ref arg_fscale,
+                                instruction_ref arg_z_pt,
+                                const onnx_parser::node_info& info)
+{
+    auto in_lens = x_in->get_shape().lens();
+
+    // prep 1: broadcast scale. it can come as a scalar or a 1-D tensor.
+    instruction_ref bcast_scale;
+    if(arg_fscale->get_shape().elements() > 1)
+        bcast_scale = info.add_instruction(
+            migraphx::make_op("broadcast", {{"axis", 0}, {"out_lens", in_lens}}), arg_fscale);
+    else
+        bcast_scale = info.add_instruction(
+            migraphx::make_op("multibroadcast", {{"out_lens", in_lens}}), arg_fscale);
+
+    // prep 2: broadcast zero point. it can come as a scalar or a 1-D tensor.
+    instruction_ref bcast_zero_pt;
+    if(arg_z_pt->get_shape().elements() > 1)
+        bcast_zero_pt = info.add_instruction(
+            migraphx::make_op("broadcast", {{"axis", 0}, {"out_lens", in_lens}}), arg_z_pt);
+    else
+        bcast_zero_pt = info.add_instruction(
+            migraphx::make_op("multibroadcast", {{"out_lens", in_lens}}), arg_z_pt);
+
+    // op_name is either quantizelinear or dequantizelinear:
+    return info.add_instruction(migraphx::make_op(op_name), x_in, bcast_scale, bcast_zero_pt);
+}
+
+// Multibroadcast a scaler..
+instruction_ref bcast_scalar_instr(const migraphx::shape& shape_out,
+                                   instruction_ref arg_in,
+                                   const onnx_parser::node_info& info)
+{
+    auto bcast_instr_out = info.add_instruction(
+        migraphx::make_op("multibroadcast", {{"out_lens", shape_out.lens()}}), arg_in);
+    return bcast_instr_out;
+}
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/onnx/include/migraphx/onnx/broadcast_qdq.hpp

@@ -0,0 +1,56 @@
+/*
+ * The MIT License (MIT)
+ *
+ * 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
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ * THE SOFTWARE.
+ */
+
+#ifndef MIGRAPHX_GUARD_AMDMIGRAPHX_ONNX_BROADCAST_QDQ_HPP
+#define MIGRAPHX_GUARD_AMDMIGRAPHX_ONNX_BROADCAST_QDQ_HPP
+
+#include <string>
+
+#include <migraphx/onnx/op_parser.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/instruction.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+// This method is to prep for quantizelinear or dequantizelinear operation for
+// either the broadcasting of weight-scale or zero-points of qlinearadd operator
+// outputs: operator op (inputs x, broadcasted: scale (float) & zero_pt (8-bit))
+instruction_ref bcast_qdq_instr(const std::string& op_name,
+                                instruction_ref x_in,
+                                instruction_ref arg_fscale,
+                                instruction_ref arg_z_pt,
+                                const onnx_parser::node_info& info);
+
+// Multibroadcast a scaler..
+instruction_ref bcast_scalar_instr(const migraphx::shape& shape_out,
+                                   instruction_ref arg_in,
+                                   const onnx_parser::node_info& info);
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif