diff --git a/cmake/Embed.cmake b/cmake/Embed.cmake index 5a120e18339..f3ef595c995 100755 --- a/cmake/Embed.cmake +++ b/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(&_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 #include -#include #include -const std::unordered_map>& ${EMBED_NAME}(); +std::unordered_map ${EMBED_NAME}(); ") file(WRITE "${PARSE_SRC}" " #include <${EMBED_NAME}.hpp> ${EXTERNS} -const std::unordered_map>& ${EMBED_NAME}() +std::unordered_map ${EMBED_NAME}() { - static const std::unordered_map> result = {${INIT_KERNELS}}; + static std::unordered_map 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 +extern const char _binary_${SYMBOL}_start[] = { ${ARRAY_VALUES} }; +extern const size_t _binary_${SYMBOL}_length = sizeof(_binary_${SYMBOL}_start); +") endif() endforeach() endfunction() diff --git a/requirements.txt b/requirements.txt index 6e24e0110a4..3a06a254528 100755 --- a/requirements.txt +++ b/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 diff --git a/src/compile_src.cpp b/src/compile_src.cpp index 6466dad8493..2029f734519 100644 --- a/src/compile_src.cpp +++ b/src/compile_src.cpp @@ -46,7 +46,7 @@ std::vector src_compiler::compile(const std::vector& 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(); diff --git a/src/include/migraphx/compile_src.hpp b/src/include/migraphx/compile_src.hpp index 6803a66af65..9baf2bc6cba 100644 --- a/src/include/migraphx/compile_src.hpp +++ b/src/include/migraphx/compile_src.hpp @@ -37,8 +37,18 @@ inline namespace MIGRAPHX_INLINE_NS { struct src_file { fs::path path; - std::pair 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& pair) + : path{pair.first}, content{pair.second} + { + } }; struct MIGRAPHX_EXPORT src_compiler diff --git a/src/include/migraphx/normalize_attributes.hpp b/src/include/migraphx/normalize_attributes.hpp index e88003e6d85..61887af3f59 100644 --- a/src/include/migraphx/normalize_attributes.hpp +++ b/src/include/migraphx/normalize_attributes.hpp @@ -52,6 +52,7 @@ using dependent_type = typename select_dependent_type::type; * \param attr_val the normalize_axes attributes from the operator * \param prefix error message prefix */ +MIGRAPHX_EXPORT std::vector normalize_axes(const std::vector& axes, const shape& input_shape, const value& attr_val, @@ -67,6 +68,7 @@ std::vector normalize_axes(const std::vector& axes, * \param attr_val the normalize_axes attributes from the operator * \param prefix error message prefix */ +MIGRAPHX_EXPORT std::vector normalize_indices(const std::vector& indices, const std::vector& axes, const shape& input_shape, diff --git a/src/include/migraphx/op/allocate.hpp b/src/include/migraphx/op/allocate.hpp index 5d1ca929ffb..33ea6bb2260 100644 --- a/src/include/migraphx/op/allocate.hpp +++ b/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{}; diff --git a/src/include/migraphx/pad_calc.hpp b/src/include/migraphx/pad_calc.hpp index a17c0bc3028..cb5972fb644 100644 --- a/src/include/migraphx/pad_calc.hpp +++ b/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& padding, diff --git a/src/onnx/broadcast_qdq.cpp b/src/onnx/broadcast_qdq.cpp new file mode 100644 index 00000000000..955c6a484df --- /dev/null +++ b/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 + +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 diff --git a/src/onnx/include/migraphx/onnx/broadcast_qdq.hpp b/src/onnx/include/migraphx/onnx/broadcast_qdq.hpp new file mode 100644 index 00000000000..04432b01d86 --- /dev/null +++ b/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 + +#include +#include +#include + +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 diff --git a/src/onnx/parse_qlinearadd.cpp b/src/onnx/parse_qlinearadd.cpp new file mode 100644 index 00000000000..81f00e71d6a --- /dev/null +++ b/src/onnx/parse_qlinearadd.cpp @@ -0,0 +1,154 @@ +/* + * 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 +#include +#include +#include +#include +#include +#include + +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 +{ + std::vector operators() const { return {{"QLinearAdd"}}; } + + // basic type checking for QLinearAdd Operator + void check_inputs(const std::vector& 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& 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 diff --git a/src/quantization.cpp b/src/quantization.cpp index 83f1be38c61..1716c31e19d 100644 --- a/src/quantization.cpp +++ b/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>> int8_quant_params = std::make_shared>>(); std::shared_ptr> max_abs_vals = std::make_shared>(); @@ -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{}}); } diff --git a/src/targets/gpu/CMakeLists.txt b/src/targets/gpu/CMakeLists.txt index d6dcd182f8e..c53860552b7 100644 --- a/src/targets/gpu/CMakeLists.txt +++ b/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) diff --git a/src/targets/gpu/compile_hip.cpp b/src/targets/gpu/compile_hip.cpp index 30565fadfe6..14e64c1e97c 100644 --- a/src/targets/gpu/compile_hip.cpp +++ b/src/targets/gpu/compile_hip.cpp @@ -248,7 +248,7 @@ compile_hip_src(const std::vector& 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& 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& 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 { diff --git a/src/targets/gpu/compile_hip_code_object.cpp b/src/targets/gpu/compile_hip_code_object.cpp index 828093e96b3..74aa9d24138 100644 --- a/src/targets/gpu/compile_hip_code_object.cpp +++ b/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 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& 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(), " "); diff --git a/src/targets/gpu/device/targets.hpp.in b/src/targets/gpu/device/targets.hpp.in index 0d73f7be36d..0a0e19aba6b 100644 --- a/src/targets/gpu/device/targets.hpp.in +++ b/src/targets/gpu/device/targets.hpp.in @@ -24,7 +24,7 @@ #ifndef MIGRAPHX_GUARD_DEVICE_TARGETS_CPP #define MIGRAPHX_GUARD_DEVICE_TARGETS_CPP -#include +#include #include #include @@ -34,9 +34,13 @@ namespace gpu { namespace device { #define MIGRAPHX_GPU_TARGETS "@GPU_TARGETS@" // NOLINT +MIGRAPHX_DEVICE_EXPORT const std::vector& get_targets(); + +MIGRAPHX_DEVICE_EXPORT std::string get_targets_as_string(); +MIGRAPHX_DEVICE_EXPORT std::string get_device_name(); } // namespace device diff --git a/src/targets/gpu/include/migraphx/gpu/compile_hip.hpp b/src/targets/gpu/include/migraphx/gpu/compile_hip.hpp index 3607cd46a40..447cdab126d 100644 --- a/src/targets/gpu/include/migraphx/gpu/compile_hip.hpp +++ b/src/targets/gpu/include/migraphx/gpu/compile_hip.hpp @@ -45,10 +45,7 @@ MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_HIPRTC_WORKAROUNDS); struct hiprtc_src_file { hiprtc_src_file() = default; - hiprtc_src_file(const src_file& s) - : path(s.path.string()), content(s.content.first, s.content.second) - { - } + hiprtc_src_file(const src_file& s) : path(s.path.string()), content(s.content) {} std::string path; std::string content; template diff --git a/src/targets/gpu/jit/ck_gemm.cpp b/src/targets/gpu/jit/ck_gemm.cpp index 65bed54a800..2937f653c09 100644 --- a/src/targets/gpu/jit/ck_gemm.cpp +++ b/src/targets/gpu/jit/ck_gemm.cpp @@ -112,8 +112,7 @@ static std::vector create_ck_headers() std::vector srcs; std::transform( header_strings.begin(), header_strings.end(), std::back_inserter(srcs), [&](auto&& p) { - return src_file{fs::path{p.first}, - {p.second.data(), p.second.data() + p.second.size()}}; + return src_file{p.first, p.second}; }); return srcs; } diff --git a/src/targets/gpu/mlir.cpp b/src/targets/gpu/mlir.cpp index 347f93333ea..fa5bdcc5df6 100644 --- a/src/targets/gpu/mlir.cpp +++ b/src/targets/gpu/mlir.cpp @@ -320,7 +320,10 @@ struct mlir_program MlirType make_tensor(const shape& s) const { - assert(s.standard()); + if(not s.standard()) + MIGRAPHX_THROW("MLIR expects all tensors to be in standard shape"); + if(s.dynamic()) + MIGRAPHX_THROW("MLIR does not support dynamic shapes"); std::vector lens(s.lens().begin(), s.lens().end()); return mlirRankedTensorTypeGet( lens.size(), lens.data(), make_type(s.type()), mlirAttributeGetNull()); diff --git a/test/gpu/jit.cpp b/test/gpu/jit.cpp index 90c19c7c490..b92f1419310 100644 --- a/test/gpu/jit.cpp +++ b/test/gpu/jit.cpp @@ -155,7 +155,7 @@ int main() {} migraphx::src_file make_src_file(const std::string& name, const std::string& content) { - return {name, std::make_pair(content.data(), content.data() + content.size())}; + return {name, content}; } TEST_CASE(simple_compile_hip) diff --git a/test/gpu/stream_sync.cpp b/test/gpu/stream_sync.cpp index da7c00e5a0a..e9d2c97d0b6 100644 --- a/test/gpu/stream_sync.cpp +++ b/test/gpu/stream_sync.cpp @@ -64,7 +64,7 @@ int main() {} migraphx::src_file make_src_file(const std::string& name, const std::string& content) { - return {name, std::make_pair(content.data(), content.data() + content.size())}; + return {name, content}; } hip_stream_ptr get_stream() diff --git a/test/jit.cpp b/test/jit.cpp index a1e5253846e..c269702d6a4 100644 --- a/test/jit.cpp +++ b/test/jit.cpp @@ -48,9 +48,7 @@ compile_function(const std::string& src, const std::string& flags, const std::st migraphx::src_compiler compiler; compiler.flags = flags + "-std=c++14 -fPIC -shared"; compiler.output = "libsimple.so"; - migraphx::src_file f; - f.path = "main.cpp"; - f.content = std::make_pair(src.data(), src.data() + src.size()); + migraphx::src_file f{"main.cpp", src}; auto image = compiler.compile({f}); return migraphx::dynamic_loader{image}.get_function(fname); } diff --git a/test/msgpack.cpp b/test/msgpack.cpp index 8b4e8ce1d42..82912aeac26 100644 --- a/test/msgpack.cpp +++ b/test/msgpack.cpp @@ -97,9 +97,12 @@ TEST_CASE(test_msgpack_bool) TEST_CASE(test_msgpack_float) { - migraphx::value v = 3.0; + // changed all double values in this code to not end with .0 because on msgpack for Windows if + // input type is double and ends with .0 it could be converted to uint64_t or int64_t and the + // goal of these functions is to test double without conversions + migraphx::value v = 3.01; auto buffer = migraphx::to_msgpack(v); - EXPECT(buffer == msgpack_buffer(3.0)); + EXPECT(buffer == msgpack_buffer(3.01)); EXPECT(migraphx::from_msgpack(buffer) == v); } @@ -129,10 +132,10 @@ TEST_CASE(test_msgpack_empty_array) TEST_CASE(test_msgpack_object) { - migraphx::value v = {{"one", 1.0}, {"three", 3.0}, {"two", 2.0}}; + migraphx::value v = {{"one", 1.01}, {"three", 3.01}, {"two", 2.01}}; auto buffer = migraphx::to_msgpack(v); EXPECT(buffer == msgpack_buffer(std::map{ - {"one", 1.0}, {"three", 3.0}, {"two", 2.0}})); + {"one", 1.01}, {"three", 3.01}, {"two", 2.01}})); EXPECT(migraphx::from_msgpack(buffer) == v); } @@ -157,17 +160,17 @@ struct foo TEST_CASE(test_msgpack_object_class) { - migraphx::value v = {{"a", 1.0}, {"b", "abc"}}; + migraphx::value v = {{"a", 1.01}, {"b", "abc"}}; auto buffer = migraphx::to_msgpack(v); - EXPECT(buffer == msgpack_buffer(foo{1.0, "abc"})); + EXPECT(buffer == msgpack_buffer(foo{1.01, "abc"})); EXPECT(migraphx::from_msgpack(buffer) == v); } TEST_CASE(test_msgpack_array_class) { - migraphx::value v = {{{"a", 1.0}, {"b", "abc"}}, {{"a", 3.0}, {"b", "xyz"}}}; + migraphx::value v = {{{"a", 1.01}, {"b", "abc"}}, {{"a", 3.01}, {"b", "xyz"}}}; auto buffer = migraphx::to_msgpack(v); - EXPECT(buffer == msgpack_buffer(std::vector{foo{1.0, "abc"}, foo{3.0, "xyz"}})); + EXPECT(buffer == msgpack_buffer(std::vector{foo{1.01, "abc"}, foo{3.01, "xyz"}})); EXPECT(migraphx::from_msgpack(buffer) == v); } diff --git a/test/multi_target/multitarget_test.cpp b/test/multi_target/multitarget_test.cpp index 6b25405ac87..c375612d760 100644 --- a/test/multi_target/multitarget_test.cpp +++ b/test/multi_target/multitarget_test.cpp @@ -37,7 +37,6 @@ #include #include #include -#include #include #include #include diff --git a/test/onnx/gen_onnx.py b/test/onnx/gen_onnx.py index b1dda835829..dc57e455ca4 100644 --- a/test/onnx/gen_onnx.py +++ b/test/onnx/gen_onnx.py @@ -5096,6 +5096,61 @@ def prelu_brcst_test(): return ([node], [arg0, arg1], [arg_out]) +@onnx_test() +def qlinearadd_test(): + a = helper.make_tensor_value_info('A', TensorProto.UINT8, [64]) + sc_a = helper.make_tensor('A_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_a = helper.make_tensor('A_zero_point', TensorProto.UINT8, [], [0]) + + b = helper.make_tensor_value_info('B', TensorProto.UINT8, [64]) + sc_b = helper.make_tensor('B_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_b = helper.make_tensor('B_zero_point', TensorProto.UINT8, [], + [128]) + + sc_c = helper.make_tensor('C_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_c = helper.make_tensor('C_zero_point', TensorProto.UINT8, [], [64]) + + c = helper.make_tensor_value_info('C', TensorProto.UINT8, [64]) + + node = onnx.helper.make_node( + 'QLinearAdd', + inputs=[ + 'A', 'A_scale', 'A_zero_point', 'B', 'B_scale', 'B_zero_point', + 'C_scale', 'C_zero_point' + ], + outputs=['C'], + ) + return ([node], [a, b], [c], + [sc_a, zero_pt_a, sc_b, zero_pt_b, sc_c, zero_pt_c]) + + +@onnx_test() +def qlinearadd_bcast_test(): + a = helper.make_tensor_value_info('A', TensorProto.INT8, [64]) + sc_a = helper.make_tensor('A_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_a = helper.make_tensor('A_zero_point', TensorProto.INT8, [], [0]) + + b = helper.make_tensor_value_info('B', TensorProto.INT8, [1, 1, 64]) + sc_b = helper.make_tensor('B_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_b = helper.make_tensor('B_zero_point', TensorProto.INT8, [], [32]) + + sc_c = helper.make_tensor('C_scale', TensorProto.FLOAT, [], [0.05]) + zero_pt_c = helper.make_tensor('C_zero_point', TensorProto.INT8, [], [-64]) + + c = helper.make_tensor_value_info('C', TensorProto.INT8, [1, 1, 64]) + + node = onnx.helper.make_node( + 'QLinearAdd', + inputs=[ + 'A', 'A_scale', 'A_zero_point', 'B', 'B_scale', 'B_zero_point', + 'C_scale', 'C_zero_point' + ], + outputs=['C'], + ) + return ([node], [a, b], [c], + [sc_a, zero_pt_a, sc_b, zero_pt_b, sc_c, zero_pt_c]) + + @onnx_test() def quantizelinear_test(): arg0 = helper.make_tensor_value_info('0', TensorProto.FLOAT, [5]) diff --git a/test/onnx/onnx_test.cpp b/test/onnx/onnx_test.cpp index 493b4aee308..54031e8d5f2 100644 --- a/test/onnx/onnx_test.cpp +++ b/test/onnx/onnx_test.cpp @@ -4856,6 +4856,59 @@ TEST_CASE(prelu_brcst_test) EXPECT(p == prog); } +TEST_CASE(qlinearadd_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + + auto a = mm->add_parameter("A", {migraphx::shape::uint8_type, {64}}); + auto b = mm->add_parameter("B", {migraphx::shape::uint8_type, {64}}); + + auto sc_a = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {0.05}}); + auto z_pt_a = mm->add_literal(migraphx::literal{migraphx::shape::uint8_type, {0}}); + + auto sc_b = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {0.05}}); + auto z_pt_b = mm->add_literal(migraphx::literal{migraphx::shape::uint8_type, {128}}); + + auto sc_c = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {0.05}}); + auto z_pt_c = mm->add_literal(migraphx::literal{migraphx::shape::uint8_type, {64}}); + + auto scale_a_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), sc_a); + + auto z_pt_a_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), z_pt_a); + + auto fp_a = + mm->add_instruction(migraphx::make_op("dequantizelinear"), a, scale_a_bcast, z_pt_a_bcast); + + auto scale_b_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), sc_b); + + auto z_pt_b_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), z_pt_b); + + auto fp_b = + mm->add_instruction(migraphx::make_op("dequantizelinear"), b, scale_b_bcast, z_pt_b_bcast); + + auto fp_c = mm->add_instruction(migraphx::make_op("add"), fp_a, fp_b); + + auto scale_c_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), sc_c); + + auto z_pt_c_bcast = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", {64}}}), z_pt_c); + + auto c = + mm->add_instruction(migraphx::make_op("quantizelinear"), fp_c, scale_c_bcast, z_pt_c_bcast); + + mm->add_return({c}); + + auto prog = migraphx::parse_onnx("qlinearadd_test.onnx"); + + EXPECT(p.sort() == prog.sort()); +} + TEST_CASE(quantizelinear_test) { migraphx::program p; diff --git a/test/onnx/qlinearadd_bcast_test.onnx b/test/onnx/qlinearadd_bcast_test.onnx new file mode 100644 index 00000000000..24dd99207c9 Binary files /dev/null and b/test/onnx/qlinearadd_bcast_test.onnx differ diff --git a/test/onnx/qlinearadd_test.onnx b/test/onnx/qlinearadd_test.onnx new file mode 100644 index 00000000000..eaa5500bf19 Binary files /dev/null and b/test/onnx/qlinearadd_test.onnx differ diff --git a/test/onnx/verify_onnx.cpp b/test/onnx/verify_onnx.cpp index e4cc7edcc2b..43226d97157 100644 --- a/test/onnx/verify_onnx.cpp +++ b/test/onnx/verify_onnx.cpp @@ -1245,6 +1245,79 @@ TEST_CASE(nonzero_test) EXPECT(migraphx::verify::verify_rms_range(result_vector, gold)); } +TEST_CASE(qlinearadd_test) +{ + // github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.QLinearAdd + migraphx::program p = migraphx::parse_onnx("qlinearadd_test.onnx"); + p.compile(migraphx::make_target("ref")); + + migraphx::shape a{migraphx::shape::uint8_type, {64}}; + std::vector data_a = {0, 2, 4, 6, 8, 10, 12, 14, 16, 18, 20, 22, 24, + 26, 28, 30, 32, 34, 36, 38, 40, 42, 44, 46, 48, 50, + 52, 54, 56, 58, 60, 62, 64, 66, 68, 70, 72, 74, 76, + 78, 80, 82, 84, 86, 88, 90, 92, 94, 96, 98, 100, 102, + 104, 106, 108, 110, 112, 114, 116, 118, 120, 122, 124, 126}; + + migraphx::shape b{migraphx::shape::uint8_type, {64}}; + std::vector data_b = {128, 126, 124, 122, 120, 118, 116, 114, 112, 110, 108, 106, 104, + 102, 100, 98, 96, 94, 92, 90, 88, 86, 84, 82, 80, 78, + 76, 74, 72, 70, 68, 66, 64, 62, 60, 58, 56, 54, 52, + 50, 48, 46, 44, 42, 40, 38, 36, 34, 32, 30, 28, 26, + 24, 22, 20, 18, 16, 14, 12, 10, 8, 6, 4, 2}; + + migraphx::parameter_map pp; + pp["A"] = migraphx::argument(a, data_a.data()); + pp["B"] = migraphx::argument(b, data_b.data()); + auto result = p.eval(pp).back(); + + std::vector result_vector; + result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); }); + + std::vector gold = {64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, + 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, + 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, + 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64, 64}; + + EXPECT(migraphx::verify::verify_rms_range(result_vector, gold)); +} + +TEST_CASE(qlinearadd_bcast_test) +{ + // github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.QLinearAdd + migraphx::program p = migraphx::parse_onnx("qlinearadd_bcast_test.onnx"); + p.compile(migraphx::make_target("ref")); + + migraphx::shape a{migraphx::shape::int8_type, {64}}; + std::vector data_a = {-64, -62, -60, -58, -56, -54, -52, -50, -48, -46, -44, -42, -40, + -38, -36, -34, -32, -30, -28, -26, -24, -22, -20, -18, -16, -14, + -12, -10, -8, -6, -4, -2, 0, 2, 4, 6, 8, 10, 12, + 14, 16, 18, 20, 22, 24, 26, 28, 30, 32, 34, 36, 38, + 40, 42, 44, 46, 48, 50, 52, 54, 56, 58, 60, 62}; + + migraphx::shape b{migraphx::shape::int8_type, {1, 1, 64}}; + std::vector data_b = {96, 94, 92, 90, 88, 86, 84, 82, 80, 78, 76, 74, 72, + 70, 68, 66, 64, 62, 60, 58, 56, 54, 52, 50, 48, 46, + 44, 42, 40, 38, 36, 34, 32, 30, 28, 26, 24, 22, 20, + 18, 16, 14, 12, 10, 8, 6, 4, 2, 0, -2, -4, -6, + -8, -10, -12, -14, -16, -18, -20, -22, -24, -26, -28, -30}; + + migraphx::parameter_map pp; + pp["A"] = migraphx::argument(a, data_a.data()); + pp["B"] = migraphx::argument(b, data_b.data()); + auto result = p.eval(pp).back(); + + std::vector result_vector; + result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); }); + + std::vector gold = {-64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, + -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, + -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, + -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, + -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64, -64}; + + EXPECT(migraphx::verify::verify_rms_range(result_vector, gold)); +} + TEST_CASE(resize_downsample_f_test) { migraphx::program p = migraphx::parse_onnx("resize_downsample_f_test.onnx");