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/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/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 bdd7cbd2f65..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()); 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/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 +}