Merge branch 'develop' into windows_cxx_compilation

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/msgpack-c@cpp-3.3.0 -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/cpp_generator.cpp

@@ -213,13 +213,13 @@ cpp_generator::function cpp_generator::generate_module(const module& m,
                 ins->get_literal().visit([&](auto v) {
                     assert(v.size() == 1);
                     auto x = v.front();
-                    if(std::isinf(x))
+                    if(std::isinf(static_cast<double>(x)))
                     {
                         string_literal = "__builtin_huge_val()";
                         if(x < 0)
                             string_literal = "-__builtin_huge_val()";
                     }
-                    else if(std::isnan(x))
+                    else if(std::isnan(static_cast<double>(x)))
                         string_literal = "__builtin_nan()";
                     else
                         string_literal = ins->get_literal().to_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/op/convert.hpp

@@ -68,7 +68,7 @@ struct convert : unary<convert>
             auto y = x;
             shape::visit(type, [&](auto as) {
                 // clamping value between target_type's max and min doesn't work for NaNs,
-                if(std::isnan(x))
+                if(std::isnan(static_cast<double>(x)))
                 {
                     y = as.nan();
                 }

src/include/migraphx/op/isnan.hpp

@@ -35,7 +35,7 @@ struct isnan : unary<isnan>
 {
     auto apply() const
     {
-        return [](auto x) { return std::isnan(x); };
+        return [](auto x) { return std::isnan(static_cast<double>(x)); };
     }
 
     std::string name() const { return "isnan"; }

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/include/migraphx/verify.hpp

@@ -90,8 +90,7 @@ struct not_finite_fn
     template <class T>
     bool operator()(T x) const
     {
-        using std::isfinite;
-        return not isfinite(x);
+        return not std::isfinite(static_cast<double>(x));
     }
 };
 static constexpr not_finite_fn not_finite{};
@@ -101,8 +100,7 @@ struct compare_mag_fn
     template <class T, class U>
     bool operator()(T x, U y) const
     {
-        using std::fabs;
-        return fabs(x) < fabs(y);
+        return std::fabs(x) < std::fabs(y);
     }
 };
 static constexpr compare_mag_fn compare_mag{};

src/onnx/broadcast_qdq.cpp

+/*
+ * 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

+/*
+ * 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

src/onnx/parse_qlinearadd.cpp

+/*
+ * 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/op_parser.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/common.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/onnx/checks.hpp>
+#include <migraphx/onnx/broadcast_qdq.hpp>
+#include <migraphx/instruction.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+/*
+ *********************************************************************************
+ *  Reference: see QLinearAdd in                                                 *
+ *  https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md  *
+ *********************************************************************************
+
+  com.microsoft.QLinearAdd
+  Performs element-wise binary addition on 8 bit data types (with Numpy-style broadcasting support).
+
+  C = (A_scale * (A - A_zero_point) + B_scale * (B - B_zero_point))/C_scale + C_zero_point
+
+  Version
+  This version of the operator has been available since version 1 of the 'com.microsoft' operator
+  set.
+
+  Inputs (7 - 8)
+  A : T
+  First operand.
+
+  A_scale : tensor(float)
+  Input A's scale. It's a scalar, which means a per-tensor/layer quantization.
+
+  A_zero_point (optional) : T
+  Input A zero point. Default value is 0 if it's not specified. It's a scalar, which means a
+  per-tensor/layer quantization.
+
+  B : T
+  Second operand.
+
+  B_scale : tensor(float)
+  Input B's scale. It's a scalar, which means a per-tensor/layer quantization.
+
+  B_zero_point (optional) : T
+  Input B zero point. Default value is 0 if it's not specified. It's a scalar, which means a
+  per-tensor/layer quantization.
+
+  C_scale : tensor(float)
+  Output scale. It's a scalar, which means a per-tensor/layer quantization.
+
+  C_zero_point (optional) : T
+
+  Output zero point. Default value is 0 if it's not specified. It's a scalar, which means a
+  per-tensor/layer quantization.
+
+  Outputs
+  C : T
+  Result, has same element type as two inputs
+
+  Type Constraints
+  T : tensor(uint8), tensor(int8)
+  Constrain input and output types to 8 bit signed and unsigned tensors.
+
+*/
+
+struct parse_qlinearadd : op_parser<parse_qlinearadd>
+{
+    std::vector<op_desc> operators() const { return {{"QLinearAdd"}}; }
+
+    // basic type checking for QLinearAdd Operator
+    void check_inputs(const std::vector<instruction_ref>& args) const
+    {
+        if(args.size() < 7)
+            MIGRAPHX_THROW("QLINEARADD: missing inputs");
+
+        const auto& in_a = args[0];
+        const auto& in_b = args[3];
+
+        auto sh_a = in_a->get_shape();
+        auto sh_b = in_b->get_shape();
+
+        auto type_a = sh_a.type();
+        auto type_b = sh_b.type();
+        if(type_a != migraphx::shape::int8_type and type_a != migraphx::shape::uint8_type)
+            MIGRAPHX_THROW("QLINEARADD: unsupported input type");
+        if(type_b != migraphx::shape::int8_type and type_b != migraphx::shape::uint8_type)
+            MIGRAPHX_THROW("QLINEARADD: unsupported input type");
+        if(type_a != type_b)
+            MIGRAPHX_THROW("QLINEARADD: mismatched input types");
+    }
+
+    instruction_ref parse(const op_desc& /* opd */,
+                          const onnx_parser& /*parser*/,
+                          const onnx_parser::node_info& info,
+                          const std::vector<instruction_ref>& args) const
+    {
+        check_inputs(args);
+
+        // A
+        const auto& in_a         = args[0];
+        const auto& in_scale_a   = args[1];
+        const auto& in_zero_pt_a = args[2];
+
+        auto dquant_a = bcast_qdq_instr("dequantizelinear", in_a, in_scale_a, in_zero_pt_a, info);
+
+        // B
+        const auto& in_b         = args[3];
+        const auto& in_scale_b   = args[4];
+        const auto& in_zero_pt_b = args[5];
+        auto dquant_b = bcast_qdq_instr("dequantizelinear", in_b, in_scale_b, in_zero_pt_b, info);
+
+        // C = A + B
+        auto out_c = info.add_common_op("add", dquant_a, dquant_b);
+
+        const auto& in_scale_c = args[6];
+
+        // zero_pt for C is supplied as the last optional argument..
+        if(args.size() == 8)
+            return (bcast_qdq_instr("quantizelinear", out_c, in_scale_c, args[7], info));
+
+        // if no zero_pt: just broadcast the scale..
+        auto bcast_scale_c = bcast_scalar_instr(out_c->get_shape(), in_scale_c, info);
+        return (info.add_instruction(migraphx::make_op("quantizelinear"), out_c, bcast_scale_c));
+    }
+};
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/program.cpp

@@ -347,7 +347,7 @@ void program::finalize()
 template <class T>
 std::string classify(T x)
 {
-    switch(std::fpclassify(x))
+    switch(std::fpclassify(static_cast<double>(x)))
     {
     case FP_INFINITE: return "inf";
     case FP_NAN: return "nan";

src/quantization.cpp

@@ -70,6 +70,10 @@ void quantize_int8(program& prog,
         MIGRAPHX_THROW("QUANTIZE_INT8: only support DOT and CONVOLUTION operation");
     }
 
+    // Run optimize_module() before converting to int8 to const eval and fold in FP32 to
+    // avoid loss of precision.
+    run_passes(prog, {optimize_module{}});
+
     std::shared_ptr<std::vector<std::pair<float, float>>> int8_quant_params =
         std::make_shared<std::vector<std::pair<float, float>>>();
     std::shared_ptr<std::vector<float>> max_abs_vals = std::make_shared<std::vector<float>>();
@@ -143,10 +147,7 @@ void quantize_int8(program& prog,
 
     run_passes(prog,
                {quantize_int8_pass{ins_names, *int8_quant_params},
-                eliminate_common_subexpression{},
-                dead_code_elimination{},
-                simplify_reshapes{},
-                dead_code_elimination{},
+                optimize_module{},
                 simplify_qdq{},
                 dead_code_elimination{}});
 }

src/targets/gpu/CMakeLists.txt

@@ -48,10 +48,18 @@ else()
     set(MIGRAPHX_USE_HIPRTC ON CACHE BOOL "Use hipRTC APIs")
 endif()
 
-include(Embed)
 file(GLOB KERNEL_FILES CONFIGURE_DEPENDS
     ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/*.hpp)
 message(STATUS "KERNEL_FILES: ${KERNEL_FILES}")
+
+if(WIN32)
+    # TODO: re-enable when CK is ported to Windows
+    list(REMOVE_ITEM KERNEL_FILES
+        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck_gemm.hpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/migraphx/kernels/ck.hpp)
+endif()
+
+include(Embed)
 add_embed_library(migraphx_kernels ${KERNEL_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/kernels/include/)
 
 configure_file(device/targets.hpp.in include/migraphx/gpu/device/targets.hpp)

src/targets/gpu/compile_hip.cpp

@@ -248,7 +248,7 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
         {
             if(src.path.extension() != ".cpp")
                 continue;
-            std::cout << std::string(src.content.first, src.len()) << std::endl;
+            std::cout << std::string(src.content) << std::endl;
         }
     }
     auto p      = dynamic_loader::path(&compile_hip_src_with_hiprtc);
@@ -338,7 +338,7 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
         {
             if(src.path.extension() != ".cpp")
                 continue;
-            std::cout << std::string(src.content.first, src.len()) << std::endl;
+            std::cout << std::string(src.content) << std::endl;
         }
     }
 
@@ -359,9 +359,7 @@ bool hip_has_flags(const std::vector<std::string>& flags)
         join_strings(flags, " ") + " -x hip -c --offload-arch=gfx900 --cuda-device-only";
 
     std::string src;
-    src_file input;
-    input.path    = "main.cpp";
-    input.content = std::make_pair(src.data(), src.data() + src.size());
+    src_file input{"main.cpp", src};
 
     try
     {

src/targets/gpu/compile_hip_code_object.cpp

@@ -172,21 +172,17 @@ operation compile_hip_code_object(const std::string& content, hip_compile_option
     assert(options.inputs.size() == options.virtual_inputs.size() or
            options.virtual_inputs.empty());
     std::vector<src_file> srcs = options.additional_src_files;
-    std::transform(migraphx_kernels().begin(),
-                   migraphx_kernels().end(),
-                   std::back_inserter(srcs),
-                   [](auto&& p) {
-                       auto&& name = p.first;
-                       auto&& c    = p.second;
-                       auto path   = name;
-                       return src_file{path, c};
-                   });
-    srcs.push_back(src_file{fs::path{"main.cpp"},
-                            std::make_pair(content.data(), content.data() + content.size())});
+    static auto kernels{::migraphx_kernels()};
+    std::transform(
+        kernels.begin(),
+        kernels.end(),
+        std::back_inserter(srcs),
+        [](const std::pair<std::string_view, std::string_view>& elem) { return src_file{elem}; });
+    srcs.emplace_back("main.cpp", content);
     auto args_hpp =
         generate_args_hpp(options.virtual_inputs.empty() ? options.inputs : options.virtual_inputs);
-    srcs.push_back(src_file{fs::path{"args.hpp"},
-                            std::make_pair(args_hpp.data(), args_hpp.data() + args_hpp.size())});
+    srcs.emplace_back("args.hpp", args_hpp);
+
     options.params += " -DMIGRAPHX_NGLOBAL=" + std::to_string(options.global);
     options.params += " -DMIGRAPHX_NLOCAL=" + std::to_string(options.local);
     options.params += " " + join_strings(compiler_warnings(), " ");

src/targets/gpu/device/targets.hpp.in

@@ -24,7 +24,7 @@
 #ifndef MIGRAPHX_GUARD_DEVICE_TARGETS_CPP
 #define MIGRAPHX_GUARD_DEVICE_TARGETS_CPP
 
-#include <migraphx/config.hpp>
+#include <migraphx/gpu/device/config.hpp>
 #include <string>
 #include <vector>
 
@@ -34,9 +34,13 @@ namespace gpu {
 namespace device {
 #define MIGRAPHX_GPU_TARGETS "@GPU_TARGETS@" // NOLINT
 
+MIGRAPHX_DEVICE_EXPORT
 const std::vector<std::string>& get_targets();
+
+MIGRAPHX_DEVICE_EXPORT
 std::string get_targets_as_string();
 
+MIGRAPHX_DEVICE_EXPORT
 std::string get_device_name();
 
 } // namespace device