diff --git a/.github/workflows/ci.yaml b/.github/workflows/ci.yaml index 924b96ea0d1..a3d13bef174 100644 --- a/.github/workflows/ci.yaml +++ b/.github/workflows/ci.yaml @@ -65,6 +65,7 @@ jobs: # GH actions can not update existing cache, as a workaround clear cache and then save it - name: Clear tidy cache before saving + continue-on-error: true if: ${{ steps.tidy_restore.outputs.cache-hit }} shell: bash env: @@ -72,7 +73,6 @@ jobs: run: | gh extension install actions/gh-actions-cache --pin v1.0.1 gh actions-cache delete ${{ steps.tidy_restore.outputs.cache-matched-key }} --confirm - continue-on-error: true - name: Save cache files for tidy uses: actions/cache/save@v3 @@ -124,6 +124,7 @@ jobs: # GH actions can not update existing cache, as a workaround clear cache and then save it - name: Clear cppcheck cache before saving + continue-on-error: true if: ${{ steps.cppcheck_restore.outputs.cache-hit }} shell: bash env: @@ -131,7 +132,6 @@ jobs: run: | gh extension install actions/gh-actions-cache --pin v1.0.1 gh actions-cache delete ${{ steps.cppcheck_restore.outputs.cache-matched-key }} --confirm - continue-on-error: true - name: Save cache files for cppcheck uses: actions/cache/save@v3 @@ -212,6 +212,7 @@ jobs: mkdir build cd build CXX=/opt/rocm/llvm/bin/clang++ CC=/opt/rocm/llvm/bin/clang cmake \ + -DMIGRAPHX_DISABLE_LARGE_BUFFER_TESTS=On \ -DBUILD_DEV=On \ -DCMAKE_CXX_COMPILER_LAUNCHER=/usr/local/bin/ccache \ -DCMAKE_C_COMPILER_LAUNCHER=/usr/local/bin/ccache \ @@ -219,12 +220,12 @@ jobs: make -j$(nproc) tests driver - name: Clear ccache cache before saving + continue-on-error: true if: ${{ steps.ccache_restore.outputs.cache-hit }} shell: bash env: GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} run: | - set +x gh extension install actions/gh-actions-cache --pin v1.0.1 gh actions-cache delete ${{ steps.ccache_restore.outputs.cache-matched-key }} --confirm @@ -365,6 +366,7 @@ jobs: rbuild build -d cget -s gh -T check \ -DCMAKE_BUILD_TYPE=${{matrix.configuration}} \ -DMIGRAPHX_ENABLE_PYTHON=${{matrix.configuration == 'release' && 'On' || 'Off'}} \ + -DMIGRAPHX_DISABLE_LARGE_BUFFER_TESTS=On \ -DBUILD_DEV=On \ -DCMAKE_CXX_FLAGS_DEBUG="-g1 -Os -fdebug-prefix-map=$PWD=. -fdebug-types-section -fno-omit-frame-pointer -fsanitize=undefined -fno-sanitize-recover=undefined" \ -DCMAKE_CXX_FLAGS_CODECOV="-g1 -Og -fdebug-prefix-map=$PWD=. -fdebug-types-section -fprofile-arcs -ftest-coverage -fno-omit-frame-pointer" \ @@ -374,12 +376,12 @@ jobs: # GH actions can not update existing cache, as a workaround clear cache and then save it - name: Clear ccache cache before saving + continue-on-error: true if: ${{ steps.ccache_restore.outputs.cache-hit }} shell: bash env: GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} run: | - set +x gh extension install actions/gh-actions-cache --pin v1.0.1 gh actions-cache delete ${{ steps.ccache_restore.outputs.cache-matched-key }} --confirm @@ -481,6 +483,7 @@ jobs: rbuild build -d cget -s gh -T check \ -DCMAKE_BUILD_TYPE=${{matrix.configuration}} \ -DMIGRAPHX_ENABLE_PYTHON=${{matrix.configuration == 'release' && 'On' || 'Off'}} \ + -DMIGRAPHX_DISABLE_LARGE_BUFFER_TESTS=On \ -DBUILD_DEV=On \ -DCMAKE_CXX_FLAGS_DEBUG="-g1 -Os -fdebug-prefix-map=$PWD=. -fdebug-types-section -fno-omit-frame-pointer -fsanitize=undefined -fno-sanitize-recover=undefined" \ -DCMAKE_CXX_FLAGS_CODECOV="-g1 -Og -fdebug-prefix-map=$PWD=. -fdebug-types-section -fprofile-arcs -ftest-coverage -fno-omit-frame-pointer" \ @@ -491,15 +494,14 @@ jobs: # this is a workaround, with GH actions can not update existing cache - name: Clear ccache cache before saving + continue-on-error: true if: ${{ steps.ccache_restore_fpga.outputs.cache-hit }} shell: bash env: GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} run: | - set +x gh extension install actions/gh-actions-cache gh actions-cache delete ${{ steps.ccache_restore_fpga.outputs.cache-matched-key }} --confirm - continue-on-error: true - name: Save cache files for ccache uses: actions/cache/save@v3 diff --git a/mlir-requirements.txt b/mlir-requirements.txt index 2e3f001355f..02af8d38df6 100644 --- a/mlir-requirements.txt +++ b/mlir-requirements.txt @@ -21,4 +21,4 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN # THE SOFTWARE. ##################################################################################### -ROCmSoftwarePlatform/rocMLIR@3657f509bfed86bb79d5c6e24aa237e48f09f9f3 -DBUILD_FAT_LIBROCKCOMPILER=On +ROCmSoftwarePlatform/rocMLIR@2c519c48eaa278d13e6c40bc0941119826d71512 -DBUILD_FAT_LIBROCKCOMPILER=On diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 62815382990..730d76515b8 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -1,7 +1,7 @@ ##################################################################################### # 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 @@ -28,6 +28,7 @@ include(ROCMInstallTargets) include(ROCMPackageConfigHelpers) include(RegisterOp) include(CheckCXXLinkerFlag) + add_library(migraphx adjust_allocation.cpp @@ -184,6 +185,8 @@ register_migraphx_ops( quant_convolution quant_dot quantizelinear + random_uniform + random_seed recip reduce_max reduce_mean diff --git a/src/driver/CMakeLists.txt b/src/driver/CMakeLists.txt index 294c34ceaa7..995daecb3f5 100755 --- a/src/driver/CMakeLists.txt +++ b/src/driver/CMakeLists.txt @@ -45,6 +45,9 @@ if(NOT WIN32) endif() 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) rocm_install_targets( diff --git a/src/driver/main.cpp b/src/driver/main.cpp index 56b65c5eab1..9ce7e98b397 100644 --- a/src/driver/main.cpp +++ b/src/driver/main.cpp @@ -475,13 +475,15 @@ struct compiler { if(is_offload_copy_set(p) and not co.offload_copy) { - std::cout << "MIGraphX program was likely compiled with offload_copy set, Try " - "passing " - "`--enable-offload-copy` if program run fails.\n"; + std::cout + << "[WARNING]: MIGraphX program was likely compiled with offload_copy " + "set, Try " + "passing " + "`--enable-offload-copy` if program run fails.\n"; } else if(co.offload_copy) { - std::cout << "MIGraphX program was likely compiled without " + std::cout << "[WARNING]: MIGraphX program was likely compiled without " "offload_copy set, Try " "removing " "`--enable-offload-copy` flag if passed to driver, if program run " @@ -802,6 +804,13 @@ int main(int argc, const char* argv[]) auto&& m = get_commands(); auto cmd = args.front(); + + if(cmd == "ort-sha") + { + std::cout << MIGRAPHX_ORT_SHA1 << std::endl; + return 0; + } + if(m.count(cmd) > 0) { m.at(cmd)(argv[0], {args.begin() + 1, args.end()}); diff --git a/src/include/migraphx/check_shapes.hpp b/src/include/migraphx/check_shapes.hpp index ced99e5d593..a273d4627d0 100644 --- a/src/include/migraphx/check_shapes.hpp +++ b/src/include/migraphx/check_shapes.hpp @@ -70,13 +70,19 @@ struct check_shapes check_dynamic(); } - template + template {})> check_shapes(const std::vector& s, const Op& op, const bool d = false) : begin(s.begin()), end(s.end()), name(op.name()), dynamic_allowed(d) { check_dynamic(); } + check_shapes(const std::vector& s, const std::string& n, const bool d = false) + : begin(s.begin()), end(s.end()), name(n), dynamic_allowed(d) + { + check_dynamic(); + } + void check_dynamic() const { if(not dynamic_allowed and this->any_of([&](const shape& s) { return s.dynamic(); })) @@ -228,6 +234,16 @@ struct check_shapes return *this; } + /*! + * Check all shapes have the same layout. + */ + const check_shapes& same_layout() const + { + if(not this->same([](const shape& s) { return find_permutation(s); })) + MIGRAPHX_THROW(prefix() + "Layouts do not match"); + return *this; + } + /*! * Check all shapes are standard. */ diff --git a/src/include/migraphx/op/random_seed.hpp b/src/include/migraphx/op/random_seed.hpp new file mode 100644 index 00000000000..ccd018838c2 --- /dev/null +++ b/src/include/migraphx/op/random_seed.hpp @@ -0,0 +1,72 @@ +/* + * 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_OPERATORS_RANDOM_SEED_HPP +#define MIGRAPHX_GUARD_OPERATORS_RANDOM_SEED_HPP + +#include +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { +namespace op { + +/** + * Generates a random seed for the use of random number generators. Generating the seed + * at runtime guarantees there will be a different random sequence on every execution. + * This operation has no inputs or attributes, and outputs an unsigned integer tensor with + * a single value. + */ +struct random_seed +{ + shape::type_t dtype = shape::type_t::uint64_type; + + template + static auto reflect(Self& self, F f) + { + return pack(f(self.dtype, "dtype")); + } + + std::string name() const { return "random_seed"; } + shape compute_shape(const std::vector& inputs) const + { + check_shapes{inputs, *this}.has(0); + return shape{dtype}; + } + + argument compute(const shape& output_shape, const std::vector&) const + { + argument result(output_shape); + + result.visit([&](auto output) { output.front() = std::random_device{}(); }); + return result; + } +}; + +} // namespace op +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + +#endif diff --git a/src/include/migraphx/op/random_uniform.hpp b/src/include/migraphx/op/random_uniform.hpp new file mode 100644 index 00000000000..2514d33f389 --- /dev/null +++ b/src/include/migraphx/op/random_uniform.hpp @@ -0,0 +1,103 @@ +/* + * 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. + */ + +/** + * Random Uniform distribution operator. Given a shape, populate it with random + * values. Calls to random_uniform using the same randomization seed as a + * literal input will + * always generate the same pseudo-random sequence. + * + * Inputs: (1) randomization seed (any type is allowed) + * (2) output buffer argument to be populated. + * + * Attributes: none + * + * Output: Returns the buffer from input #2. + * + */ +#ifndef MIGRAPHX_GUARD_OPERATORS_RANDOM_UNIFORM_HPP +#define MIGRAPHX_GUARD_OPERATORS_RANDOM_UNIFORM_HPP + +#include +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { +namespace op { + +/** + * random_uniform populates the passed shape with random numbers, in a uniform + * distribution. Range for floating-point data types is (0, 1); + * for integer types it is [0, ] + */ +struct random_uniform +{ + // The random_uniform operation needs the random number generator seed + // to be passed as a runtime input. + + std::string name() const { return "random_uniform"; } + shape compute_shape(std::vector inputs) const + { + check_shapes{inputs, *this, true}.has(2); + + return inputs.at(1); + } + + argument compute(const shape&, std::vector args) const + { + // Output goes into the passed buffer, not the shape output. + auto result = args[1]; + + uint64_t local_seed = args[0].at(0); + std::mt19937 gen(local_seed); + + result.visit([&](auto output) { + using type = typename decltype(output)::value_type; + if constexpr(std::is_integral{}) + { + // default range for all integer types is + // (0, std::uniform_int_distribution::max()). + // Todo: enable different ranges + std::uniform_int_distribution dis; + std::generate(output.begin(), output.end(), [&] { return dis(gen); }); + } + else + { + // default real distribution type is double with range (0, 1); + std::uniform_real_distribution<> dis; + std::generate(output.begin(), output.end(), [&] { return dis(gen); }); + } + }); + return result; + } + + std::ptrdiff_t output_alias(const std::vector&) const { return 1; } +}; + +} // namespace op +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + +#endif diff --git a/src/load_save.cpp b/src/load_save.cpp index d32671522f9..2cad6f48705 100644 --- a/src/load_save.cpp +++ b/src/load_save.cpp @@ -21,6 +21,7 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN * THE SOFTWARE. */ +#include #include #include #include @@ -60,9 +61,29 @@ void save(const program& p, const std::string& filename, const file_options& opt { write_buffer(filename, save_buffer(p, options)); } + +// MIOpen doesn't support serializing fusion plans with Find-2.0 APIs +void print_miopen_warning(const program& p) +{ + auto mods = p.get_modules(); + if(std::any_of(mods.begin(), mods.end(), [](const auto* m) { + return std::any_of(m->begin(), m->end(), [](const instruction& i) { + return i.name() == "gpu::miopen_fusion"; + }); + })) + { + std::cout << "[WARNING]: Program has miopen_fusion instructions for which tuned solutions " + "are not stored inside serialized MIGraphX program. Consider serializing with " + "MIGRAPHX_DISABLE_MIOPEN_FUSION=1 flag set." + << std::endl; + ; + } +} + std::vector save_buffer(const program& p, const file_options& options) { value v = p.to_value(); + print_miopen_warning(p); std::vector buffer; if(options.format == "msgpack") { diff --git a/src/msgpack.cpp b/src/msgpack.cpp index 4fdca67bd95..ed95704bc80 100644 --- a/src/msgpack.cpp +++ b/src/msgpack.cpp @@ -25,6 +25,33 @@ #include #include +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { + +// Leave an extra byte for error checking +constexpr std::size_t msgpack_size_limit = std::numeric_limits::max() - 1; + +template +std::size_t msgpack_chunk_size(const Range& r) +{ + return 1 + (r.size() - 1) / msgpack_size_limit; +} + +template +void msgpack_chunk_for_each(Iterator start, Iterator last, F f) +{ + while(std::distance(start, last) > msgpack_size_limit) + { + auto next = std::next(start, msgpack_size_limit); + f(start, next); + start = next; + } + f(start, last); +} + +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + namespace msgpack { MSGPACK_API_VERSION_NAMESPACE(MSGPACK_DEFAULT_API_NS) { @@ -63,16 +90,31 @@ MSGPACK_API_VERSION_NAMESPACE(MSGPACK_DEFAULT_API_NS) break; } case msgpack::type::BIN: { + // For backwards compatibility v = migraphx::value::binary{o.via.bin.ptr, o.via.bin.size}; break; } case msgpack::type::ARRAY: { - migraphx::value r = migraphx::value::array{}; - std::for_each( - o.via.array.ptr, - o.via.array.ptr + o.via.array.size, - [&](const msgpack::object& so) { r.push_back(so.as()); }); - v = r; + if(o.via.array.size != 0 and o.via.array.ptr->type == msgpack::type::BIN) + { + auto bin = migraphx::value::binary{}; + std::for_each( + o.via.array.ptr, + o.via.array.ptr + o.via.array.size, + [&](const msgpack::object& so) { + bin.insert(bin.end(), so.via.bin.ptr, so.via.bin.ptr + so.via.bin.size); + }); + v = bin; + } + else + { + migraphx::value r = migraphx::value::array{}; + std::for_each( + o.via.array.ptr, + o.via.array.ptr + o.via.array.size, + [&](const msgpack::object& so) { r.push_back(so.as()); }); + v = r; + } break; } case msgpack::type::MAP: { @@ -102,8 +144,12 @@ MSGPACK_API_VERSION_NAMESPACE(MSGPACK_DEFAULT_API_NS) { const auto* data = reinterpret_cast(x.data()); auto size = x.size(); - o.pack_bin(size); - o.pack_bin_body(data, size); + o.pack_array(migraphx::msgpack_chunk_size(x)); + migraphx::msgpack_chunk_for_each( + data, data + size, [&](const char* start, const char* last) { + o.pack_bin(last - start); + o.pack_bin_body(start, last - start); + }); return o; } }; @@ -129,6 +175,8 @@ MSGPACK_API_VERSION_NAMESPACE(MSGPACK_DEFAULT_API_NS) o.pack_array(0); return; } + if(v.size() > migraphx::msgpack_size_limit) + MIGRAPHX_THROW("Size is too large for msgpack"); if(not v.front().get_key().empty()) { o.pack_map(v.size()); diff --git a/src/program.cpp b/src/program.cpp index 8be1dd3e065..7e0dbb8b7a9 100644 --- a/src/program.cpp +++ b/src/program.cpp @@ -624,7 +624,7 @@ std::string get_migraphx_version() program file version is for the data structure or format of the MXR file. Version should be bumped if any changes occur to the format of the MXR file. */ -const int program_file_version = 6; +const int program_file_version = 7; value program::to_value() const { diff --git a/src/targets/gpu/fuse_mlir.cpp b/src/targets/gpu/fuse_mlir.cpp index c33adc560d3..e9481b7d6e7 100644 --- a/src/targets/gpu/fuse_mlir.cpp +++ b/src/targets/gpu/fuse_mlir.cpp @@ -103,7 +103,10 @@ struct mlir_op } if(ins->name() == "@return") { - return ins_shapes[ins->inputs().at(0)].with_type(type); + auto s = ins_shapes[ins->inputs().at(0)].with_type(type); + if(not s.standard()) + MIGRAPHX_THROW("MLIR doesnt support non-standard output"); + return s; } std::vector input_shapes; input_shapes.resize(ins->inputs().size()); @@ -119,6 +122,33 @@ struct mlir_op MIGRAPHX_REGISTER_OP(mlir_op); namespace { +std::tuple> +fuse_input_ops_and_gemm_based_op(module_ref mm, instruction_ref gemm_based_op) +{ + std::vector top_inputs; + std::vector imm_inputs; + size_t input_cnt = 0; + for(instruction_ref input : gemm_based_op->inputs()) + { + std::vector op_stream; + while(contains({"slice", "transpose", "contiguous", "reshape"}, input->name())) + { + op_stream.push_back(input->get_operator()); + input = input->inputs().at(0); + } + top_inputs.push_back(input); + instruction_ref prev_input = + mm->add_parameter("y" + std::to_string(input_cnt++), input->get_shape()); + for(const auto& op : reverse(op_stream)) + { + prev_input = mm->add_instruction(op, {prev_input}); + } + imm_inputs.push_back(prev_input); + } + instruction_ref new_gemm_based_op = + mm->add_instruction(gemm_based_op->get_operator(), imm_inputs); + return {new_gemm_based_op, top_inputs}; +} MIGRAPHX_PRED_MATCHER(is_mlir_conv, instruction_ref ins) { @@ -134,7 +164,7 @@ MIGRAPHX_PRED_MATCHER(is_mlir_conv, instruction_ref ins) return true; } -struct find_mlir_op +struct find_mlir_fused_ops { auto matcher() const { @@ -163,34 +193,6 @@ struct find_mlir_op return ins_map; } - std::tuple> - fuse_input_ops_and_gemm_based_op(module_ref mm, instruction_ref gemm_based_op) const - { - std::vector top_inputs; - std::vector imm_inputs; - size_t input_cnt = 0; - for(instruction_ref input : gemm_based_op->inputs()) - { - std::vector op_stream; - while(contains({"slice", "transpose", "contiguous", "reshape"}, input->name())) - { - op_stream.push_back(input->get_operator()); - input = input->inputs().at(0); - } - top_inputs.push_back(input); - instruction_ref prev_input = - mm->add_parameter("y" + std::to_string(input_cnt++), input->get_shape()); - for(const auto& op : reverse(op_stream)) - { - prev_input = mm->add_instruction(op, {prev_input}); - } - imm_inputs.push_back(prev_input); - } - instruction_ref new_gemm_based_op = - mm->add_instruction(gemm_based_op->get_operator(), imm_inputs); - return {new_gemm_based_op, top_inputs}; - } - // Whitelist supported fusion options, including imposing type constraints // for cases where MLIR only supports an operation (usually a pointwise function) // on particular types. @@ -236,8 +238,7 @@ struct find_mlir_op "log", "recip", "rsqrt", - // There are bugs in MLIR right now for models using sigmoid so disable it for now - // "sigmoid", + "sigmoid", "softmax", "tanh", }; @@ -301,14 +302,95 @@ struct find_mlir_op } }; +struct find_mlir_standalone_convolution_op +{ + auto matcher() const { return match::name("convolution"); } + + void apply(module_pass_manager& mpm, const match::matcher_result& r) const + { + auto conv_based_op = r.result; + // enable only for fp32/fp16/i8 types + if(std::any_of(conv_based_op->inputs().begin(), conv_based_op->inputs().end(), [&](auto i) { + return not contains( + {shape::type_t::float_type, shape::type_t::half_type, shape::type_t::int8_type}, + i->get_shape().type()); + })) + return; + + static size_t counter = 0; + module_ref mm = mpm.create_module("mlir_" + std::to_string(counter++)); + mm->set_bypass(); + auto [anchor_op, top_inputs] = fuse_input_ops_and_gemm_based_op(mm, conv_based_op); + mm->add_return({anchor_op}); + mpm.get_module().replace_instruction( + conv_based_op, mlir_op{conv_based_op->get_operator()}, top_inputs, {mm}); + } +}; + +/** + * @brief Declares a new MIGraphX environment variable which forces to generate + * only specific MLIR operations. + * + * The variable, if defined, forces MIGraphX to use only specific operations + * with MLIR regardless of the underlying GPU architecture. The variable accepts + * a list of operations separated by comma. The variable recognizes the following + * operations: "fused", "convolution". If the variable is not defined MIGraphX + * will decide by itself which operations to delegate to MLIR. The variable is + * intended to be primarily used by rocMLIR developers. + */ +MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_MLIR_USE_SPECIFIC_OPS); +bool is_self_decide() { return string_value_of(MIGRAPHX_MLIR_USE_SPECIFIC_OPS{}, "").empty(); } + +bool is_requested(std::string_view option) +{ + assert(not is_self_decide()); + auto string_value = string_value_of(MIGRAPHX_MLIR_USE_SPECIFIC_OPS{}, ""); + const auto options = split_string(string_value, ','); + return contains(options, option); +} + +bool is_fusion_enabled() +{ + if(is_self_decide()) + { + return true; + } + return is_requested("fused"); +} + +bool is_standalone_convs_enabled(context* ctx) +{ + if(is_self_decide()) + { + if(ctx == nullptr) + { + return false; + } + else + { + const auto& device = ctx->get_current_device(); + const std::string navi_family{"gfx110"}; + return starts_with(device.get_gfx_name(), navi_family); + } + } + return is_requested("convolution"); +} } // namespace -#endif +#endif // MIGRAPHX_MLIR void fuse_mlir::apply(module_pass_manager& mpm) const { #ifdef MIGRAPHX_MLIR - match::find_matches(mpm, find_mlir_op{}); + if(is_fusion_enabled()) + { + match::find_matches(mpm, find_mlir_fused_ops{}); + } + + if(is_standalone_convs_enabled(this->ctx)) + { + match::find_matches(mpm, find_mlir_standalone_convolution_op{}); + } #else (void)mpm; #endif diff --git a/src/targets/gpu/include/migraphx/gpu/convolution.hpp b/src/targets/gpu/include/migraphx/gpu/convolution.hpp index 0b85075c66d..d6680f17ec8 100644 --- a/src/targets/gpu/include/migraphx/gpu/convolution.hpp +++ b/src/targets/gpu/include/migraphx/gpu/convolution.hpp @@ -84,8 +84,10 @@ struct miopen_convolution { check_shapes{inputs, op}.has(4); std::vector conv_inputs(inputs.begin(), inputs.begin() + 2); - check_shapes{conv_inputs, *this}.max_ndims(5).packed_layouts( - {{0, 1, 2}, {0, 1, 2, 3}, {0, 2, 3, 1}, {0, 1, 2, 3, 4}}); + check_shapes{conv_inputs, *this} + .max_ndims(5) + .packed_layouts({{0, 1, 2}, {0, 1, 2, 3}, {0, 2, 3, 1}, {0, 1, 2, 3, 4}}) + .same_layout(); return migraphx::compute_shape(op, conv_inputs); } diff --git a/src/verify_args.cpp b/src/verify_args.cpp index 489e078c963..e19373ca2f8 100644 --- a/src/verify_args.cpp +++ b/src/verify_args.cpp @@ -78,16 +78,6 @@ bool verify_args(const std::string& name, if(verify::range_zero(target)) std::cout << "Target data is all zeros" << std::endl; - // auto mxdiff = max_diff(ref, target); - // std::cout << "Max diff: " << mxdiff << std::endl; - - // auto idx = mismatch_idx(ref, target, float_equal); - // if(idx < verify::range_distance(ref)) - // { - // std::cout << "Mismatch at " << idx << ": " << ref[idx] << " != " << target[idx] - // << std::endl; - // } - auto ref_nan_idx = find_idx(ref, verify::not_finite); if(ref_nan_idx >= 0) std::cout << "Non finite number found in ref at " << ref_nan_idx << ": " @@ -97,7 +87,7 @@ bool verify_args(const std::string& name, if(target_nan_idx >= 0) std::cout << "Non finite number found in target at " << target_nan_idx << ": " << target[target_nan_idx] << std::endl; - // std::cout << std::endl; + std::cout << "MIGraphX verification passed successfully." << std::endl; } }); return passed; diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 1f3da5e75f1..e3fbc6387c5 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -31,6 +31,11 @@ set(CTEST_PARALLEL_LEVEL ${N} CACHE STRING "CTest parallel level") add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -j ${CTEST_PARALLEL_LEVEL} -C ${CMAKE_CFG_INTDIR} --timeout 5000) add_custom_target(tests) +set(MIGRAPHX_DISABLE_LARGE_BUFFER_TESTS Off CACHE BOOL "") +if(MIGRAPHX_DISABLE_LARGE_BUFFER_TESTS) + add_compile_definitions(MIGRAPHX_DISABLE_LARGE_BUFFER_TESTS) +endif() + find_program(MIGRAPHX_GDB gdb) if(MIGRAPHX_GDB) diff --git a/test/check_shapes_test.cpp b/test/check_shapes_test.cpp index 021877ad471..557c9f1b7a8 100644 --- a/test/check_shapes_test.cpp +++ b/test/check_shapes_test.cpp @@ -31,24 +31,39 @@ using migraphx::shape; -bool create_shapes(bool dynamic_allowed) +void create_shapes(bool dynamic_allowed) { - try - { - shape a{shape::int64_type, {3}}; - shape b{shape::float_type, {{3, 6}, {4, 4}}}; - auto op = migraphx::make_op("add"); - migraphx::check_shapes{{a, b}, op, dynamic_allowed}.has(2); - return true; - } - catch(...) - { - return false; - } + shape a{shape::int64_type, {3}}; + shape b{shape::float_type, {{3, 6}, {4, 4}}}; + migraphx::check_shapes{{a, b}, "", dynamic_allowed}.has(2); } -TEST_CASE(allow_dynamic_shape) { EXPECT(create_shapes(true)); } +TEST_CASE(allow_dynamic_shape) +{ + EXPECT(not test::throws([] { create_shapes(true); })); +} + +TEST_CASE(fail_dynamic_shape) +{ + EXPECT(test::throws([] { create_shapes(false); })); +} -TEST_CASE(fail_dynamic_shape) { EXPECT(not create_shapes(false)); } +TEST_CASE(same_layout_fail) +{ + EXPECT(test::throws([] { + shape a{shape::float_type, {2, 3}}; + shape b{shape::float_type, {2, 3}, {1, 2}}; + migraphx::check_shapes{{a, b}, ""}.same_layout(); + })); +} + +TEST_CASE(same_layout_pass) +{ + EXPECT(not test::throws([] { + shape a{shape::float_type, {2, 3}, {1, 2}}; + shape b{shape::float_type, {2, 3}, {1, 2}}; + migraphx::check_shapes{{a, b}, ""}.same_layout(); + })); +} int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/msgpack.cpp b/test/msgpack.cpp index 61c26e9457e..8b4e8ce1d42 100644 --- a/test/msgpack.cpp +++ b/test/msgpack.cpp @@ -25,13 +25,37 @@ #include #include #include +#include #include "test.hpp" +template , T>{})> +void write_msgpack(std::ostream& os, const T& src) +{ + msgpack::pack(os, src); +} +void write_msgpack(std::ostream& os, const std::vector& src) +{ + const auto limit = std::numeric_limits::max() - 1; + std::vector> chunks; + if(src.size() > limit) + { + // Only test two chunks + assert(std::distance(src.begin() + limit, src.end()) < limit); + chunks.emplace_back(src.begin(), src.begin() + limit); + chunks.emplace_back(src.begin() + limit, src.end()); + } + else + { + chunks = {src}; + } + write_msgpack(os, chunks); +} + template std::vector msgpack_buffer(const T& src) { std::stringstream buffer; - msgpack::pack(buffer, src); + write_msgpack(buffer, src); buffer.seekg(0); std::string str = buffer.str(); return std::vector(str.data(), str.data() + str.size()); // NOLINT @@ -147,4 +171,51 @@ TEST_CASE(test_msgpack_array_class) EXPECT(migraphx::from_msgpack(buffer) == v); } +TEST_CASE(test_msgpack_binary) +{ + migraphx::value::binary bin{64}; + std::iota(bin.begin(), bin.end(), 1); + auto buffer = migraphx::to_msgpack(bin); + EXPECT(buffer == msgpack_buffer(bin)); + EXPECT(migraphx::from_msgpack(buffer) == bin); +} + +#ifndef MIGRAPHX_DISABLE_LARGE_BUFFER_TESTS +TEST_CASE(test_msgpack_large_binary1) +{ + const std::size_t n = 4LL * 1024 * 1024 * 1024 + 2; + const char fill_value = 2; + migraphx::value v; + { + std::vector buffer; + { + migraphx::value::binary bin{n}; + std::fill(bin.begin(), bin.begin() + n / 2, fill_value); + std::fill(bin.begin() + n / 2, bin.end(), fill_value + 1); + buffer = migraphx::to_msgpack(std::move(bin)); + } + v = migraphx::from_msgpack(buffer); + } + EXPECT(v.is_binary()); + EXPECT(v.get_binary().size() == n); + EXPECT(std::all_of(v.get_binary().begin(), v.get_binary().begin() + n / 2, [](auto c) { + return c == fill_value; + })); + EXPECT(std::all_of(v.get_binary().begin() + n / 2, v.get_binary().end(), [](auto c) { + return c == fill_value + 1; + })); +} + +TEST_CASE(test_msgpack_binary2) +{ + const std::size_t n = 4LL * 1024 * 1024 * 1024 + 2; + migraphx::value::binary bin{n}; + std::size_t i = 0; + std::generate(bin.begin(), bin.end(), [&] { + i++; + return i % 256; + }); + EXPECT(migraphx::to_msgpack(bin) == msgpack_buffer(bin)); +} +#endif int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/op_shape_test.cpp b/test/op_shape_test.cpp index 00395249a12..07826424206 100644 --- a/test/op_shape_test.cpp +++ b/test/op_shape_test.cpp @@ -2260,6 +2260,20 @@ TEST_CASE(prefix_scan_sum_dyn_2d) } } +TEST_CASE(random_uniform) +{ + std::vector dd{{5, 8}, {3, 7}}; + migraphx::shape s0{migraphx::shape::uint64_type, {1}}; + migraphx::shape s1{migraphx::shape::float_type, dd}; + expect_shape(s1, migraphx::make_op("random_uniform"), s0, s1); +} + +TEST_CASE(random_seed) +{ + migraphx::shape s{migraphx::shape::uint64_type, {1}, {0}}; + expect_shape(s, migraphx::make_op("random_seed")); +} + TEST_CASE(quant_convolution_shape) { migraphx::shape output{migraphx::shape::int32_type, {4, 4, 1, 1}}; diff --git a/test/ref/random_seed.cpp b/test/ref/random_seed.cpp new file mode 100644 index 00000000000..657d1b40958 --- /dev/null +++ b/test/ref/random_seed.cpp @@ -0,0 +1,52 @@ +/* + * 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 + +/** + * Reference test for the random_seed operation + */ +TEST_CASE(random_seed_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + mm->add_instruction(migraphx::make_op("random_seed")); + + p.compile(migraphx::make_target("ref")); + auto result = p.eval({}).back(); + std::vector result_vec1(1); + result.visit([&](auto output) { result_vec1.assign(output.begin(), output.end()); }); + std::vector result_vec2(1); + // Identical calls should give different seeds every time with 1/(2^64) chance of a repeat. + // We don't analyze for true randomness. + result = p.eval({}).back(); + result.visit([&](auto output) { result_vec2.assign(output.begin(), output.end()); }); + EXPECT(result_vec1[0] != result_vec2[0]); +} diff --git a/test/ref/random_uniform.cpp b/test/ref/random_uniform.cpp new file mode 100644 index 00000000000..fa7e62df587 --- /dev/null +++ b/test/ref/random_uniform.cpp @@ -0,0 +1,174 @@ +/* + * 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 + +#include + +/** + * Reference test for the random_uniform operation. Also invokes the random_seed operation. + */ + +TEST_CASE(random_uniform_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + uint64_t seed(0); + size_t sample_size(200); + + // Shape of the random data + migraphx::shape rs{migraphx::shape::float_type, {1, sample_size}}; + + // data tensor must be allocated at this point but does not need to be initialized. + std::vector data(sample_size); + auto input = mm->add_literal(migraphx::literal(rs, data)); + + // Runtime randomization seed + migraphx::shape seed_shape{migraphx::shape::uint64_type, {1}}; + std::vector seed_data{seed}; + auto seed_input = mm->add_literal(migraphx::literal(seed_shape, seed_data)); + + mm->add_instruction(migraphx::make_op("random_uniform"), seed_input, input); + p.compile(migraphx::make_target("ref")); + + // no params_map needed + auto result = p.eval({}).back(); + std::vector result_vec(sample_size); + result.visit([&](auto output) { result_vec.assign(output.begin(), output.end()); }); + + // Compare result with the STL's mt19937 generator + std::mt19937 gen(seed); + std::uniform_real_distribution<> dis(0.0, 1.0); + std::vector rand_samples(sample_size); + std::generate(rand_samples.begin(), rand_samples.end(), [&]() { return dis(gen); }); + EXPECT(migraphx::verify::verify_range(result_vec, rand_samples, 100)); +} + +TEST_CASE(random_uniform_int_test) +{ + // random uniform distribution with an integer type input shape + migraphx::program p; + auto* mm = p.get_main_module(); + float seed(0.1); + size_t sample_size(200); + + // Shape of the random data + migraphx::shape rs{migraphx::shape::uint16_type, {1, sample_size}}; + + // data tensor must be allocated at this point but does not need to be initialized. + std::vector data(sample_size); + auto input = mm->add_literal(migraphx::literal(rs, data)); + + // Runtime randomization seed + migraphx::shape seed_shape{migraphx::shape::float_type, {1}}; + std::vector seed_data{seed}; + auto seed_input = mm->add_literal(migraphx::literal(seed_shape, seed_data)); + + mm->add_instruction(migraphx::make_op("random_uniform"), seed_input, input); + p.compile(migraphx::make_target("ref")); + + migraphx::parameter_map params0; + auto result = p.eval(params0).back(); + std::vector result_vec(sample_size); + result.visit([&](auto output) { result_vec.assign(output.begin(), output.end()); }); + + // Compare result with the STL's mt19937 generator + std::mt19937 gen(seed); + std::uniform_int_distribution dis; + std::vector rand_samples(sample_size); + std::generate(rand_samples.begin(), rand_samples.end(), [&]() { return dis(gen); }); + EXPECT(migraphx::verify::verify_range(result_vec, rand_samples)); +} + +TEST_CASE(random_uniform_dyn_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + uint64_t seed(17); + size_t sample_size(200); + + // Shape of the random data + migraphx::shape rs{migraphx::shape::float_type, {{1, 2}, {2, sample_size + 1}}}; + auto input = mm->add_parameter("Input_1", rs); + + // Runtime randomization seed + migraphx::shape seed_shape{migraphx::shape::uint64_type, {1}}; + auto seed_input = mm->add_parameter("Seed", seed_shape); + + mm->add_instruction(migraphx::make_op("random_uniform", {}), seed_input, input); + p.compile(migraphx::make_target("ref")); + + // Create a dummy input to hold the random data + migraphx::shape input_fixed_shape1{migraphx::shape::float_type, {sample_size}}; + + migraphx::parameter_map params0; + params0["Input_1"] = migraphx::argument(input_fixed_shape1); + + std::vector seed_data = {seed}; + params0["Seed"] = migraphx::argument(seed_shape, seed_data.data()); + auto result = p.eval(params0).back(); + + std::vector result_vec(sample_size); + result.visit([&](auto output) { result_vec.assign(output.begin(), output.end()); }); + + // Compare result with the STL's mt19937 generator + std::mt19937 gen(seed); + std::uniform_real_distribution<> dis(0.0, 1.0); + std::vector rand_samples(sample_size); + std::generate(rand_samples.begin(), rand_samples.end(), [&]() { return dis(gen); }); + EXPECT(migraphx::verify::verify_range(result_vec, rand_samples)); +} + +TEST_CASE(random_uniform_and_seed_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + + size_t sample_size(20000); + + // Shape of the random data + migraphx::shape rs{migraphx::shape::float_type, {{1, 2}, {2, sample_size + 1}}}; + auto input = mm->add_parameter("Input_1", rs); + + // Runtime randomization seed + auto seed_input = mm->add_instruction(migraphx::make_op("random_seed")); + mm->add_instruction(migraphx::make_op("random_uniform"), seed_input, input); + p.compile(migraphx::make_target("ref")); + + // Create a dummy input to hold the random data + migraphx::shape input_fixed_shape1{migraphx::shape::float_type, {sample_size}}; + + migraphx::parameter_map params0; + params0["Input_1"] = migraphx::argument(input_fixed_shape1); + auto result = p.eval(params0).back(); + + result.visit([&](auto output) { EXPECT(output.size() == sample_size); }); + // Do not check the content of the data since it's not repeatable +}