From fc41a6f6e5585bd79989715625940b0277c7a2b9 Mon Sep 17 00:00:00 2001 From: charlie Date: Tue, 30 Apr 2024 20:04:47 -0500 Subject: [PATCH 01/21] initial --- .../include/migraphx/onnx/json_output.hpp | 42 +++++++++++++ src/onnx/json_output.cpp | 62 +++++++++++++++++++ 2 files changed, 104 insertions(+) create mode 100644 src/onnx/include/migraphx/onnx/json_output.hpp create mode 100644 src/onnx/json_output.cpp diff --git a/src/onnx/include/migraphx/onnx/json_output.hpp b/src/onnx/include/migraphx/onnx/json_output.hpp new file mode 100644 index 00000000000..b25a9e3cc5f --- /dev/null +++ b/src/onnx/include/migraphx/onnx/json_output.hpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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_JSON_OUTPUT_HPP +#define MIGRAPHX_GUARD_AMDMIGRAPHX_JSON_OUTPUT_HPP + +#include +#include +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { +namespace onnx { + +void write_program_to_onnx_json(const program& arg, std::string filename); + +} // namespace onnx +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + +#endif diff --git a/src/onnx/json_output.cpp b/src/onnx/json_output.cpp new file mode 100644 index 00000000000..58ea8b1c3e6 --- /dev/null +++ b/src/onnx/json_output.cpp @@ -0,0 +1,62 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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 { + +void write_program_to_onnx_json(const program& prog, std::string filename) +{ + const module* mm = p.get_main_module(); + for(auto ins : iterator_for(*mm)) + { + const auto& name = ins->name(); + if(name == "@literal") + { + + } + else if(name == "@param") + { + + } + else if(name == "@outline") + { + + } + else if(name == "@return") + { + + } + else + { + + } + } +} + +} // namespace onnx +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx From 0796820efcbecc4127a710709937744f2cb2ff78 Mon Sep 17 00:00:00 2001 From: charlie Date: Mon, 28 Oct 2024 15:56:36 -0500 Subject: [PATCH 02/21] something --- src/onnx/json_output.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/onnx/json_output.cpp b/src/onnx/json_output.cpp index 58ea8b1c3e6..b783ad84668 100644 --- a/src/onnx/json_output.cpp +++ b/src/onnx/json_output.cpp @@ -30,7 +30,7 @@ namespace onnx { void write_program_to_onnx_json(const program& prog, std::string filename) { - const module* mm = p.get_main_module(); + const module* mm = prog.get_main_module(); for(auto ins : iterator_for(*mm)) { const auto& name = ins->name(); @@ -50,7 +50,7 @@ void write_program_to_onnx_json(const program& prog, std::string filename) { } - else + else // handle operators and submodules { } From 74870fbe2506877c6eb4e94278c363eb03d48cf2 Mon Sep 17 00:00:00 2001 From: charlie Date: Fri, 6 Dec 2024 20:26:14 -0600 Subject: [PATCH 03/21] more progress --- .../include/migraphx/onnx/json_output.hpp | 5 + src/onnx/json_output.cpp | 125 +++++++++++++++--- 2 files changed, 113 insertions(+), 17 deletions(-) diff --git a/src/onnx/include/migraphx/onnx/json_output.hpp b/src/onnx/include/migraphx/onnx/json_output.hpp index b25a9e3cc5f..4cea4e488ad 100644 --- a/src/onnx/include/migraphx/onnx/json_output.hpp +++ b/src/onnx/include/migraphx/onnx/json_output.hpp @@ -33,6 +33,11 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace onnx { +/** + * Writes the given MIGraphX program to a JSON format resembling the ONNX + * file format's JSON files. Used to allow visualization of MIGraphX graphs + * with ONNX tools such as Netron. + */ void write_program_to_onnx_json(const program& arg, std::string filename); } // namespace onnx diff --git a/src/onnx/json_output.cpp b/src/onnx/json_output.cpp index b783ad84668..761a551642d 100644 --- a/src/onnx/json_output.cpp +++ b/src/onnx/json_output.cpp @@ -24,37 +24,128 @@ #include +#include +#include + +#include +#include +#include +#include + namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace onnx { +using json = nlohmann::json; + void write_program_to_onnx_json(const program& prog, std::string filename) { - const module* mm = prog.get_main_module(); - for(auto ins : iterator_for(*mm)) + json output; + auto prog_value = prog.to_value(); + output["irVersion"] = prog_value.at("version").to(); + output["producerName"] = "AMDMIGraphX"; + output["producerVersion"] = prog_value.at("migraphx_version").to(); + output["graph"] = json::array(); + std::unordered_map names; + for(auto& mod : prog.get_modules()) { - const auto& name = ins->name(); - if(name == "@literal") - { + json graph = { + {"node", json::array()}, + {"initializer", json::array()}, + {"input", json::array()}, + {"output", json::array()}, + {"valueInfo", json::array()} + }; + auto node_arr = graph["node"]; + auto lit_arr = graph["initializer"]; + auto input_arr = graph["input"]; + auto output_arr = graph["output"]; + auto edge_arr = graph["valueInfo"]; + names = mod->print( + [&](auto ins, auto ins_uids) { + const auto& name = ins->name(); + if(name == "@literal") + { + make_onnx_json_literal(ins, ins_uids); + } + else if(name == "@param") + { + make_onnx_json_input(ins, ins_uids); + } + else if(name == "@return") + { + make_onnx_json_output(ins, ins_uids); + } + else + { + node_arr.push_back(make_onnx_json_node(ins, ins_uids)); + const auto& outputs = ins.outputs(); + for(auto out_ins : outputs) + { + node_arr.push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); + } + } + }, + names); + output["graph"].push_back(graph); + } +} - } - else if(name == "@param") +auto make_onnx_json_node(instruction ins, std::unordered_map ins_uids) +{ + json::object node; + json::array input_arr; + for(auto input_ins : ins->inputs()) + { + auto name = input_ins->name(); + if(name == "@literal" or name == "@param") { - + input_arr.push_back(ins_uids.at(input_ins)); } - else if(name == "@outline") + else { - + input_arr.push_back(ins_uids.at(input_ins) + "->" + ins_uids.at(ins)); } - else if(name == "@return") - { + } + json::array output_arr; + for(auto output_ins : ins->outputs()) + { + output_arr.push_back(ins_uids.at(ins) + "->" + ins_uids.at(output_ins)); + } + node["input"] = input_arr; + node["output"] = output_arr; + node["name"] = ins_uids.at(ins); + node["opType"] = ins->name(); + json::object op_attribute_arr; + auto op_value = ins->get_operator()->to_value(); + std::for_each(op_value.begin(), op_value.end(), [](auto v){ + // No idea if this works, magic from migraphx::from_value to get the right type + op_attribute_arr.emplace({value.get_key(), migraphx::from_value(value); + }); + node["attribute"] = op_attribute_arr; + return node; +} - } - else // handle operators and submodules - { +// ONNX graph constant data called "initializer" +void make_onnx_json_literal() +{ + +} + +// ONNX graph edges called "valuetype" +void make_onnx_json_edge() +{ + +} + +void make_onnx_json_input() +{ + +} + +void make_onnx_json_output() +{ - } - } } } // namespace onnx From 25e76a0114284a3fc4248209889b7ab59a607b78 Mon Sep 17 00:00:00 2001 From: charlie Date: Sat, 7 Dec 2024 16:19:20 -0600 Subject: [PATCH 04/21] first draft --- src/CMakeLists.txt | 2 +- src/driver/main.cpp | 4 + .../migraphx/netron_output.hpp} | 17 +- src/netron_output.cpp | 212 ++++++++++++++++++ src/onnx/json_output.cpp | 153 ------------- 5 files changed, 221 insertions(+), 167 deletions(-) rename src/{onnx/include/migraphx/onnx/json_output.hpp => include/migraphx/netron_output.hpp} (70%) create mode 100644 src/netron_output.cpp delete mode 100644 src/onnx/json_output.cpp diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index eda6ea626e4..ecc17aba81b 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -40,7 +40,6 @@ add_library(migraphx common.cpp common_dims.cpp compile_src.cpp - convert_to_json.cpp cpp_generator.cpp dead_code_elimination.cpp dom_info.cpp @@ -73,6 +72,7 @@ add_library(migraphx memory_coloring.cpp module.cpp msgpack.cpp + netron_output.cpp normalize_attributes.cpp normalize_ops.cpp op_enums.cpp diff --git a/src/driver/main.cpp b/src/driver/main.cpp index 04fa0cfe3bc..36859e2178f 100644 --- a/src/driver/main.cpp +++ b/src/driver/main.cpp @@ -56,6 +56,8 @@ #include #include +#include + #include namespace migraphx { @@ -418,6 +420,8 @@ struct loader *os << to_json_string(p.to_value()) << std::endl; else if(type == "binary") write(*os, save_buffer(p)); + else if(type == "netron") + *os << make_netron_output(p) << std::endl; } }; diff --git a/src/onnx/include/migraphx/onnx/json_output.hpp b/src/include/migraphx/netron_output.hpp similarity index 70% rename from src/onnx/include/migraphx/onnx/json_output.hpp rename to src/include/migraphx/netron_output.hpp index 4cea4e488ad..ff5b12617ee 100644 --- a/src/onnx/include/migraphx/onnx/json_output.hpp +++ b/src/include/migraphx/netron_output.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 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 @@ -21,26 +21,17 @@ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN * THE SOFTWARE. */ -#ifndef MIGRAPHX_GUARD_AMDMIGRAPHX_JSON_OUTPUT_HPP -#define MIGRAPHX_GUARD_AMDMIGRAPHX_JSON_OUTPUT_HPP +#ifndef MIGRAPHX_GUARD_RTGLIB_NETRON_OUTPUT_HPP +#define MIGRAPHX_GUARD_RTGLIB_NETRON_OUTPUT_HPP #include #include -#include -#include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { -namespace onnx { -/** - * Writes the given MIGraphX program to a JSON format resembling the ONNX - * file format's JSON files. Used to allow visualization of MIGraphX graphs - * with ONNX tools such as Netron. - */ -void write_program_to_onnx_json(const program& arg, std::string filename); +MIGRAPHX_EXPORT std::string make_netron_output(const std::string& str); -} // namespace onnx } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/netron_output.cpp b/src/netron_output.cpp new file mode 100644 index 00000000000..23d3b3aa95e --- /dev/null +++ b/src/netron_output.cpp @@ -0,0 +1,212 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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 { +namespace { + +using json = nlohmann::json; + +int get_onnx_type(shape::type_t s_type) +{ + switch(s_type) + { + case shape::float_type: return 1; + case shape::uint8_type: return 2; + case shape::int8_type: return 3; + case shape::uint16_type: return 4; + case shape::int16_type: return 5; + case shape::int32_type: return 6; + case shape::int64_type: return 7; + case shape::bool_type: return 9; + case shape::half_type: return 10; + case shape::double_type: return 11; + case shape::uint32_type: return 12; + case shape::uint64_type: return 13; + case shape::bf16_type: return 16; + case shape::fp8e4m3fn_type: return 17; + case shape::fp8e4m3fnuz_type: return 18; + case shape::fp8e5m2_type: return 19; + case shape::fp8e5m2fnuz_type: return 20; + default: { + MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); + } + } +} + +auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) +{ + json node = json::object(); + json input_arr = json::array(); + for(auto input_ins : ins->inputs()) + { + auto name = input_ins->name(); + if(name == "@literal" or name == "@param") + { + input_arr.push_back(ins_uids.at(input_ins)); + } + else + { + input_arr.push_back(ins_uids.at(input_ins) + "->" + ins_uids.at(ins)); + } + } + json output_arr = json::array(); + for(auto output_ins : ins->outputs()) + { + output_arr.push_back(ins_uids.at(ins) + "->" + ins_uids.at(output_ins)); + } + node["input"] = input_arr; + node["output"] = output_arr; + node["name"] = ins_uids.at(ins); + node["opType"] = ins->name(); + json op_attribute_arr = json::object(); + auto op_value = ins->get_operator().to_value(); + std::for_each(op_value.begin(), op_value.end(), [&](auto v){ + op_attribute_arr.emplace(v.get_key(), migraphx::from_value(v)); + }); + node["attribute"] = op_attribute_arr; + return node; +} + +// ONNX graph constant data called "initializer" +auto make_onnx_json_literal(instruction_ref ins, std::unordered_map ins_uids) +{ + json lit = json::object(); + lit["dims"] = ins->get_shape().lens(); + // TODO figure out the data types number + lit["dataType"] = 1; + lit["name"] = ins_uids.at(ins); + // NULL in base64 + lit["rawData"] = "TlVMTA=="; + return lit; +} + +auto make_onnx_json_shape(const shape& s) +{ + json ret = json::object(); + json dim = json::array(); + for(std::size_t len : s.lens()) + { + dim.push_back({"dimValue", len}); + } + return ret; +} + +// ONNX graph edges called "valuetype" +auto make_onnx_json_edge(instruction_ref ins, instruction_ref out_ins, std::unordered_map ins_uids) +{ + json ret = json::object(); + shape ins_shape = ins->get_shape(); + ret["name"] = ins_uids.at(ins) + "->" + ins_uids.at(out_ins); + ret["type"] = { + "tensorType", { + {"elemType", get_onnx_type(ins_shape.type())}, + {"shape", make_onnx_json_shape(ins_shape)} + } + }; + return ret; +} + +auto make_onnx_json_in_out(instruction_ref ins, std::unordered_map ins_uids) +{ + json ret = json::object(); + shape ins_shape = ins->get_shape(); + ret["name"] = ins_uids.at(ins); + ret["type"] = { + "tensorType", { + {"elemType", get_onnx_type(ins_shape.type())}, + {"shape", make_onnx_json_shape(ins_shape)} + } + }; + return ret; +} + +} // namespace + +auto make_netron_output(const program& prog) +{ + json output; + auto prog_value = prog.to_value(); + output["irVersion"] = prog_value.at("version").to(); + output["producerName"] = "AMDMIGraphX"; + output["producerVersion"] = prog_value.at("migraphx_version").to(); + output["graph"] = json::array(); + std::unordered_map names; + for(auto& mod : prog.get_modules()) + { + json graph = {{"node", json::array()}, + {"initializer", json::array()}, + {"input", json::array()}, + {"output", json::array()}, + {"valueInfo", json::array()}}; + auto node_arr = graph["node"]; + auto lit_arr = graph["initializer"]; + auto input_arr = graph["input"]; + auto output_arr = graph["output"]; + auto edge_arr = graph["valueInfo"]; + names = mod->print( + [&](auto ins, auto ins_uids) { + const auto& name = ins->name(); + if(name == "@literal") + { + lit_arr.push_back(make_onnx_json_literal(ins, ins_uids)); + } + else if(name == "@param") + { + input_arr.push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else if(name == "@return") + { + output_arr.push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else + { + node_arr.push_back(make_onnx_json_node(ins, ins_uids)); + const auto& outputs = ins->outputs(); + for(auto out_ins : outputs) + { + node_arr.push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); + } + } + }, + names); + output["graph"].push_back(graph); + } + return output.dump(); +} + +} // namespace onnx +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx diff --git a/src/onnx/json_output.cpp b/src/onnx/json_output.cpp deleted file mode 100644 index 761a551642d..00000000000 --- a/src/onnx/json_output.cpp +++ /dev/null @@ -1,153 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2024 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 { - -using json = nlohmann::json; - -void write_program_to_onnx_json(const program& prog, std::string filename) -{ - json output; - auto prog_value = prog.to_value(); - output["irVersion"] = prog_value.at("version").to(); - output["producerName"] = "AMDMIGraphX"; - output["producerVersion"] = prog_value.at("migraphx_version").to(); - output["graph"] = json::array(); - std::unordered_map names; - for(auto& mod : prog.get_modules()) - { - json graph = { - {"node", json::array()}, - {"initializer", json::array()}, - {"input", json::array()}, - {"output", json::array()}, - {"valueInfo", json::array()} - }; - auto node_arr = graph["node"]; - auto lit_arr = graph["initializer"]; - auto input_arr = graph["input"]; - auto output_arr = graph["output"]; - auto edge_arr = graph["valueInfo"]; - names = mod->print( - [&](auto ins, auto ins_uids) { - const auto& name = ins->name(); - if(name == "@literal") - { - make_onnx_json_literal(ins, ins_uids); - } - else if(name == "@param") - { - make_onnx_json_input(ins, ins_uids); - } - else if(name == "@return") - { - make_onnx_json_output(ins, ins_uids); - } - else - { - node_arr.push_back(make_onnx_json_node(ins, ins_uids)); - const auto& outputs = ins.outputs(); - for(auto out_ins : outputs) - { - node_arr.push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); - } - } - }, - names); - output["graph"].push_back(graph); - } -} - -auto make_onnx_json_node(instruction ins, std::unordered_map ins_uids) -{ - json::object node; - json::array input_arr; - for(auto input_ins : ins->inputs()) - { - auto name = input_ins->name(); - if(name == "@literal" or name == "@param") - { - input_arr.push_back(ins_uids.at(input_ins)); - } - else - { - input_arr.push_back(ins_uids.at(input_ins) + "->" + ins_uids.at(ins)); - } - } - json::array output_arr; - for(auto output_ins : ins->outputs()) - { - output_arr.push_back(ins_uids.at(ins) + "->" + ins_uids.at(output_ins)); - } - node["input"] = input_arr; - node["output"] = output_arr; - node["name"] = ins_uids.at(ins); - node["opType"] = ins->name(); - json::object op_attribute_arr; - auto op_value = ins->get_operator()->to_value(); - std::for_each(op_value.begin(), op_value.end(), [](auto v){ - // No idea if this works, magic from migraphx::from_value to get the right type - op_attribute_arr.emplace({value.get_key(), migraphx::from_value(value); - }); - node["attribute"] = op_attribute_arr; - return node; -} - -// ONNX graph constant data called "initializer" -void make_onnx_json_literal() -{ - -} - -// ONNX graph edges called "valuetype" -void make_onnx_json_edge() -{ - -} - -void make_onnx_json_input() -{ - -} - -void make_onnx_json_output() -{ - -} - -} // namespace onnx -} // namespace MIGRAPHX_INLINE_NS -} // namespace migraphx From e7a40abc06a882f87b3870e6d3f7a192a63722ee Mon Sep 17 00:00:00 2001 From: charlie Date: Sat, 7 Dec 2024 19:15:36 -0600 Subject: [PATCH 05/21] First functional --- src/CMakeLists.txt | 1 + src/driver/main.cpp | 4 + src/include/migraphx/netron_output.hpp | 3 +- src/netron_output.cpp | 145 +++++++++++++++---------- 4 files changed, 94 insertions(+), 59 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index ecc17aba81b..d9e9b92cff1 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -40,6 +40,7 @@ add_library(migraphx common.cpp common_dims.cpp compile_src.cpp + convert_to_json.cpp cpp_generator.cpp dead_code_elimination.cpp dom_info.cpp diff --git a/src/driver/main.cpp b/src/driver/main.cpp index 36859e2178f..876081ef405 100644 --- a/src/driver/main.cpp +++ b/src/driver/main.cpp @@ -168,6 +168,10 @@ struct loader {"--binary"}, ap.help("Print out program in binary format."), ap.set_value("binary")); + ap(output_type, + {"--netron"}, + ap.help("Print out program as Netron readable json."), + ap.set_value("netron")); ap(output, {"--output", "-o"}, ap.help("Output to file.")); } diff --git a/src/include/migraphx/netron_output.hpp b/src/include/migraphx/netron_output.hpp index ff5b12617ee..c07d5f50af9 100644 --- a/src/include/migraphx/netron_output.hpp +++ b/src/include/migraphx/netron_output.hpp @@ -26,11 +26,12 @@ #include #include +#include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { -MIGRAPHX_EXPORT std::string make_netron_output(const std::string& str); +MIGRAPHX_EXPORT std::string make_netron_output(const program& prog); } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/netron_output.cpp b/src/netron_output.cpp index 23d3b3aa95e..175f651f700 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -27,14 +27,13 @@ #include #include -#include +#include #include #include #include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { -namespace onnx { namespace { using json = nlohmann::json; @@ -60,8 +59,9 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e4m3fnuz_type: return 18; case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; - default: { - MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); + case shape::tuple_type: return 0; + default: { + MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); } } } @@ -70,7 +70,7 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_mapinputs()) + for(instruction_ref input_ins : ins->inputs()) { auto name = input_ins->name(); if(name == "@literal" or name == "@param") @@ -83,18 +83,38 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_mapoutputs()) + for(instruction_ref output_ins : ins->outputs()) { - output_arr.push_back(ins_uids.at(ins) + "->" + ins_uids.at(output_ins)); + if(output_ins->name() == "@return") + { + output_arr.push_back(ins_uids.at(output_ins)); + } + else + { + output_arr.push_back(ins_uids.at(ins) + "->" + ins_uids.at(output_ins)); + } } node["input"] = input_arr; node["output"] = output_arr; node["name"] = ins_uids.at(ins); node["opType"] = ins->name(); - json op_attribute_arr = json::object(); + json op_attribute_arr = json::array(); auto op_value = ins->get_operator().to_value(); - std::for_each(op_value.begin(), op_value.end(), [&](auto v){ - op_attribute_arr.emplace(v.get_key(), migraphx::from_value(v)); + std::for_each(op_value.begin(), op_value.end(), [&](auto v) { + std::string attr_key = v.get_key(); + if(attr_key == "code_object") + { + return; + } + else if(attr_key == "symbol_name" or attr_key == "name") + { + node["opType"] = migraphx::from_value(v); + } + // TODO: attributes need to know the type + else + { + op_attribute_arr.push_back({{"name", v.get_key()}, {"ints", {0}}, {"type", "INTS"}}); + } }); node["attribute"] = op_attribute_arr; return node; @@ -106,21 +126,23 @@ auto make_onnx_json_literal(instruction_ref ins, std::unordered_mapget_shape().lens(); // TODO figure out the data types number - lit["dataType"] = 1; + lit["dataType"] = get_onnx_type(ins->get_shape().type()); lit["name"] = ins_uids.at(ins); - // NULL in base64 + // ignoring literal data, setting to "NULL" in base64 lit["rawData"] = "TlVMTA=="; return lit; } +// TODO handle dynamic shapes auto make_onnx_json_shape(const shape& s) { json ret = json::object(); json dim = json::array(); for(std::size_t len : s.lens()) { - dim.push_back({"dimValue", len}); + dim.push_back({{"dimValue", len}}); } + ret["dim"] = dim; return ret; } @@ -130,12 +152,10 @@ auto make_onnx_json_edge(instruction_ref ins, instruction_ref out_ins, std::unor json ret = json::object(); shape ins_shape = ins->get_shape(); ret["name"] = ins_uids.at(ins) + "->" + ins_uids.at(out_ins); - ret["type"] = { - "tensorType", { - {"elemType", get_onnx_type(ins_shape.type())}, - {"shape", make_onnx_json_shape(ins_shape)} - } - }; + json type = {{"tensorType", + {{"elemType", get_onnx_type(ins_shape.type())}, + {"shape", make_onnx_json_shape(ins_shape)}}}}; + ret["type"] = type; return ret; } @@ -144,26 +164,38 @@ auto make_onnx_json_in_out(instruction_ref ins, std::unordered_mapget_shape(); ret["name"] = ins_uids.at(ins); - ret["type"] = { - "tensorType", { - {"elemType", get_onnx_type(ins_shape.type())}, - {"shape", make_onnx_json_shape(ins_shape)} - } - }; + json type = {{"tensorType", + {{"elemType", get_onnx_type(ins_shape.type())}, + {"shape", make_onnx_json_shape(ins_shape)}}}}; + ret["type"] = type; + return ret; +} + +std::unordered_map make_ins_uids(const module& mod) +{ + std::unordered_map ret; + int count = 0; + for(auto ins : iterator_for(mod)) + { + std::string var_name; + var_name = mod.name() + ":"; + var_name.append(ins->name() + ":"); + var_name.append("@" + std::to_string(count)); + count++; + ret.emplace(ins, var_name); + } return ret; } } // namespace -auto make_netron_output(const program& prog) +std::string make_netron_output(const program& prog) { json output; auto prog_value = prog.to_value(); output["irVersion"] = prog_value.at("version").to(); output["producerName"] = "AMDMIGraphX"; output["producerVersion"] = prog_value.at("migraphx_version").to(); - output["graph"] = json::array(); - std::unordered_map names; for(auto& mod : prog.get_modules()) { json graph = {{"node", json::array()}, @@ -171,42 +203,39 @@ auto make_netron_output(const program& prog) {"input", json::array()}, {"output", json::array()}, {"valueInfo", json::array()}}; - auto node_arr = graph["node"]; - auto lit_arr = graph["initializer"]; - auto input_arr = graph["input"]; - auto output_arr = graph["output"]; - auto edge_arr = graph["valueInfo"]; - names = mod->print( - [&](auto ins, auto ins_uids) { - const auto& name = ins->name(); - if(name == "@literal") - { - lit_arr.push_back(make_onnx_json_literal(ins, ins_uids)); - } - else if(name == "@param") - { - input_arr.push_back(make_onnx_json_in_out(ins, ins_uids)); - } - else if(name == "@return") - { - output_arr.push_back(make_onnx_json_in_out(ins, ins_uids)); - } - else + auto ins_uids = make_ins_uids(*mod); + for(auto ins = mod->begin(); ins != mod->end(); ++ins) + { + const auto& name = ins->name(); + if(name == "@literal") + { + graph["initializer"].push_back(make_onnx_json_literal(ins, ins_uids)); + } + else if(name == "@param") + { + graph["input"].push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else if(name == "@return") + { + graph["output"].push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else + { + graph["node"].push_back(make_onnx_json_node(ins, ins_uids)); + const auto& outputs = ins->outputs(); + for(auto out_ins : outputs) { - node_arr.push_back(make_onnx_json_node(ins, ins_uids)); - const auto& outputs = ins->outputs(); - for(auto out_ins : outputs) + if(out_ins->name() != "@return") { - node_arr.push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); + graph["valueInfo"].push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); } } - }, - names); - output["graph"].push_back(graph); + } + } + output["graph"] = graph; } - return output.dump(); + return output.dump(4); } -} // namespace onnx } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx From a3751115c9f70a7ce78fe7d5f492999aa337485c Mon Sep 17 00:00:00 2001 From: charlie Date: Mon, 16 Dec 2024 15:28:29 -0600 Subject: [PATCH 06/21] handle attributes --- src/include/base64.hpp | 594 +++++++++++++++++++++++++++++++++++++++++ src/netron_output.cpp | 33 ++- 2 files changed, 620 insertions(+), 7 deletions(-) create mode 100644 src/include/base64.hpp diff --git a/src/include/base64.hpp b/src/include/base64.hpp new file mode 100644 index 00000000000..cbc01a1136b --- /dev/null +++ b/src/include/base64.hpp @@ -0,0 +1,594 @@ +#ifndef BASE64_HPP_ +#define BASE64_HPP_ + +#include +#include +#include +#include +#include +#include +#include +#include + +#if defined(__cpp_lib_bit_cast) +#include // For std::bit_cast. +#endif + +namespace base64 { + +namespace detail { + +#if defined(__cpp_lib_bit_cast) +using std::bit_cast; +#else +template +std::enable_if_t && + std::is_trivially_copyable_v, + To> +bit_cast(const From& src) noexcept +{ + static_assert(std::is_trivially_constructible_v, + "This implementation additionally requires " + "destination type to be trivially constructible"); + + To dst; + std::memcpy(&dst, &src, sizeof(To)); + return dst; +} +#endif + +inline constexpr char padding_char{'='}; +inline constexpr uint32_t bad_char{0x01FFFFFF}; + +#if !defined(__LITTLE_ENDIAN__) && !defined(__BIG_ENDIAN__) +#if(defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_BIG_ENDIAN__) || \ + (defined(__BYTE_ORDER) && __BYTE_ORDER == __BIG_ENDIAN) || \ + (defined(_BYTE_ORDER) && _BYTE_ORDER == _BIG_ENDIAN) || \ + (defined(BYTE_ORDER) && BYTE_ORDER == BIG_ENDIAN) || \ + (defined(__sun) && defined(__SVR4) && defined(_BIG_ENDIAN)) || defined(__ARMEB__) || \ + defined(__THUMBEB__) || defined(__AARCH64EB__) || defined(_MIBSEB) || defined(__MIBSEB) || \ + defined(__MIBSEB__) || defined(_M_PPC) +#define __BIG_ENDIAN__ +#elif(defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_LITTLE_ENDIAN__) || /* gcc */ \ + (defined(__BYTE_ORDER) && __BYTE_ORDER == __LITTLE_ENDIAN) /* linux header */ \ + || (defined(_BYTE_ORDER) && _BYTE_ORDER == _LITTLE_ENDIAN) || \ + (defined(BYTE_ORDER) && BYTE_ORDER == LITTLE_ENDIAN) /* mingw header */ || \ + (defined(__sun) && defined(__SVR4) && defined(_LITTLE_ENDIAN)) || /* solaris */ \ + defined(__ARMEL__) || defined(__THUMBEL__) || defined(__AARCH64EL__) || defined(_MIPSEL) || \ + defined(__MIPSEL) || defined(__MIPSEL__) || defined(_M_IX86) || defined(_M_X64) || \ + defined(_M_IA64) || /* msvc for intel processors */ \ + defined(_M_ARM) /* msvc code on arm executes in little endian mode */ +#define __LITTLE_ENDIAN__ +#endif +#endif + +#if !defined(__LITTLE_ENDIAN__) & !defined(__BIG_ENDIAN__) +#error "UNKNOWN Platform / endianness. Configure endianness explicitly." +#endif + +#if defined(__LITTLE_ENDIAN__) +std::array constexpr decode_table_0 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x000000f8, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x000000fc, + 0x000000d0, 0x000000d4, 0x000000d8, 0x000000dc, 0x000000e0, 0x000000e4, 0x000000e8, 0x000000ec, + 0x000000f0, 0x000000f4, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00000004, 0x00000008, 0x0000000c, 0x00000010, 0x00000014, 0x00000018, + 0x0000001c, 0x00000020, 0x00000024, 0x00000028, 0x0000002c, 0x00000030, 0x00000034, 0x00000038, + 0x0000003c, 0x00000040, 0x00000044, 0x00000048, 0x0000004c, 0x00000050, 0x00000054, 0x00000058, + 0x0000005c, 0x00000060, 0x00000064, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000068, 0x0000006c, 0x00000070, 0x00000074, 0x00000078, 0x0000007c, 0x00000080, + 0x00000084, 0x00000088, 0x0000008c, 0x00000090, 0x00000094, 0x00000098, 0x0000009c, 0x000000a0, + 0x000000a4, 0x000000a8, 0x000000ac, 0x000000b0, 0x000000b4, 0x000000b8, 0x000000bc, 0x000000c0, + 0x000000c4, 0x000000c8, 0x000000cc, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_1 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000e003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000f003, + 0x00004003, 0x00005003, 0x00006003, 0x00007003, 0x00008003, 0x00009003, 0x0000a003, 0x0000b003, + 0x0000c003, 0x0000d003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00001000, 0x00002000, 0x00003000, 0x00004000, 0x00005000, 0x00006000, + 0x00007000, 0x00008000, 0x00009000, 0x0000a000, 0x0000b000, 0x0000c000, 0x0000d000, 0x0000e000, + 0x0000f000, 0x00000001, 0x00001001, 0x00002001, 0x00003001, 0x00004001, 0x00005001, 0x00006001, + 0x00007001, 0x00008001, 0x00009001, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x0000a001, 0x0000b001, 0x0000c001, 0x0000d001, 0x0000e001, 0x0000f001, 0x00000002, + 0x00001002, 0x00002002, 0x00003002, 0x00004002, 0x00005002, 0x00006002, 0x00007002, 0x00008002, + 0x00009002, 0x0000a002, 0x0000b002, 0x0000c002, 0x0000d002, 0x0000e002, 0x0000f002, 0x00000003, + 0x00001003, 0x00002003, 0x00003003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_2 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00800f00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00c00f00, + 0x00000d00, 0x00400d00, 0x00800d00, 0x00c00d00, 0x00000e00, 0x00400e00, 0x00800e00, 0x00c00e00, + 0x00000f00, 0x00400f00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00400000, 0x00800000, 0x00c00000, 0x00000100, 0x00400100, 0x00800100, + 0x00c00100, 0x00000200, 0x00400200, 0x00800200, 0x00c00200, 0x00000300, 0x00400300, 0x00800300, + 0x00c00300, 0x00000400, 0x00400400, 0x00800400, 0x00c00400, 0x00000500, 0x00400500, 0x00800500, + 0x00c00500, 0x00000600, 0x00400600, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00800600, 0x00c00600, 0x00000700, 0x00400700, 0x00800700, 0x00c00700, 0x00000800, + 0x00400800, 0x00800800, 0x00c00800, 0x00000900, 0x00400900, 0x00800900, 0x00c00900, 0x00000a00, + 0x00400a00, 0x00800a00, 0x00c00a00, 0x00000b00, 0x00400b00, 0x00800b00, 0x00c00b00, 0x00000c00, + 0x00400c00, 0x00800c00, 0x00c00c00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_3 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x003e0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x003f0000, + 0x00340000, 0x00350000, 0x00360000, 0x00370000, 0x00380000, 0x00390000, 0x003a0000, 0x003b0000, + 0x003c0000, 0x003d0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00010000, 0x00020000, 0x00030000, 0x00040000, 0x00050000, 0x00060000, + 0x00070000, 0x00080000, 0x00090000, 0x000a0000, 0x000b0000, 0x000c0000, 0x000d0000, 0x000e0000, + 0x000f0000, 0x00100000, 0x00110000, 0x00120000, 0x00130000, 0x00140000, 0x00150000, 0x00160000, + 0x00170000, 0x00180000, 0x00190000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x001a0000, 0x001b0000, 0x001c0000, 0x001d0000, 0x001e0000, 0x001f0000, 0x00200000, + 0x00210000, 0x00220000, 0x00230000, 0x00240000, 0x00250000, 0x00260000, 0x00270000, 0x00280000, + 0x00290000, 0x002a0000, 0x002b0000, 0x002c0000, 0x002d0000, 0x002e0000, 0x002f0000, 0x00300000, + 0x00310000, 0x00320000, 0x00330000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +// TODO fix decoding tables to avoid the need for different indices in big +// endian? +inline constexpr size_t decidx0{0}; +inline constexpr size_t decidx1{1}; +inline constexpr size_t decidx2{2}; + +#elif defined(__BIG_ENDIAN__) + +std::array constexpr decode_table_0 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00f80000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00fc0000, + 0x00d00000, 0x00d40000, 0x00d80000, 0x00dc0000, 0x00e00000, 0x00e40000, 0x00e80000, 0x00ec0000, + 0x00f00000, 0x00f40000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00040000, 0x00080000, 0x000c0000, 0x00100000, 0x00140000, 0x00180000, + 0x001c0000, 0x00200000, 0x00240000, 0x00280000, 0x002c0000, 0x00300000, 0x00340000, 0x00380000, + 0x003c0000, 0x00400000, 0x00440000, 0x00480000, 0x004c0000, 0x00500000, 0x00540000, 0x00580000, + 0x005c0000, 0x00600000, 0x00640000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00680000, 0x006c0000, 0x00700000, 0x00740000, 0x00780000, 0x007c0000, 0x00800000, + 0x00840000, 0x00880000, 0x008c0000, 0x00900000, 0x00940000, 0x00980000, 0x009c0000, 0x00a00000, + 0x00a40000, 0x00a80000, 0x00ac0000, 0x00b00000, 0x00b40000, 0x00b80000, 0x00bc0000, 0x00c00000, + 0x00c40000, 0x00c80000, 0x00cc0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_1 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0003e000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0003f000, + 0x00034000, 0x00035000, 0x00036000, 0x00037000, 0x00038000, 0x00039000, 0x0003a000, 0x0003b000, + 0x0003c000, 0x0003d000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00001000, 0x00002000, 0x00003000, 0x00004000, 0x00005000, 0x00006000, + 0x00007000, 0x00008000, 0x00009000, 0x0000a000, 0x0000b000, 0x0000c000, 0x0000d000, 0x0000e000, + 0x0000f000, 0x00010000, 0x00011000, 0x00012000, 0x00013000, 0x00014000, 0x00015000, 0x00016000, + 0x00017000, 0x00018000, 0x00019000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x0001a000, 0x0001b000, 0x0001c000, 0x0001d000, 0x0001e000, 0x0001f000, 0x00020000, + 0x00021000, 0x00022000, 0x00023000, 0x00024000, 0x00025000, 0x00026000, 0x00027000, 0x00028000, + 0x00029000, 0x0002a000, 0x0002b000, 0x0002c000, 0x0002d000, 0x0002e000, 0x0002f000, 0x00030000, + 0x00031000, 0x00032000, 0x00033000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_2 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00000f80, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00000fc0, + 0x00000d00, 0x00000d40, 0x00000d80, 0x00000dc0, 0x00000e00, 0x00000e40, 0x00000e80, 0x00000ec0, + 0x00000f00, 0x00000f40, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00000040, 0x00000080, 0x000000c0, 0x00000100, 0x00000140, 0x00000180, + 0x000001c0, 0x00000200, 0x00000240, 0x00000280, 0x000002c0, 0x00000300, 0x00000340, 0x00000380, + 0x000003c0, 0x00000400, 0x00000440, 0x00000480, 0x000004c0, 0x00000500, 0x00000540, 0x00000580, + 0x000005c0, 0x00000600, 0x00000640, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000680, 0x000006c0, 0x00000700, 0x00000740, 0x00000780, 0x000007c0, 0x00000800, + 0x00000840, 0x00000880, 0x000008c0, 0x00000900, 0x00000940, 0x00000980, 0x000009c0, 0x00000a00, + 0x00000a40, 0x00000a80, 0x00000ac0, 0x00000b00, 0x00000b40, 0x00000b80, 0x00000bc0, 0x00000c00, + 0x00000c40, 0x00000c80, 0x00000cc0, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +std::array constexpr decode_table_3 = { + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000003e, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000003f, + 0x00000034, 0x00000035, 0x00000036, 0x00000037, 0x00000038, 0x00000039, 0x0000003a, 0x0000003b, + 0x0000003c, 0x0000003d, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x00000000, 0x00000001, 0x00000002, 0x00000003, 0x00000004, 0x00000005, 0x00000006, + 0x00000007, 0x00000008, 0x00000009, 0x0000000a, 0x0000000b, 0x0000000c, 0x0000000d, 0x0000000e, + 0x0000000f, 0x00000010, 0x00000011, 0x00000012, 0x00000013, 0x00000014, 0x00000015, 0x00000016, + 0x00000017, 0x00000018, 0x00000019, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x0000001a, 0x0000001b, 0x0000001c, 0x0000001d, 0x0000001e, 0x0000001f, 0x00000020, + 0x00000021, 0x00000022, 0x00000023, 0x00000024, 0x00000025, 0x00000026, 0x00000027, 0x00000028, + 0x00000029, 0x0000002a, 0x0000002b, 0x0000002c, 0x0000002d, 0x0000002e, 0x0000002f, 0x00000030, + 0x00000031, 0x00000032, 0x00000033, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, + 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; + +// TODO fix decoding tables to avoid the need for different indices in big +// endian? +inline constexpr size_t decidx0{1}; +inline constexpr size_t decidx1{2}; +inline constexpr size_t decidx2{3}; + +#endif + +std::array constexpr encode_table_0 = { + 'A', 'A', 'A', 'A', 'B', 'B', 'B', 'B', 'C', 'C', 'C', 'C', 'D', 'D', 'D', 'D', 'E', 'E', 'E', + 'E', 'F', 'F', 'F', 'F', 'G', 'G', 'G', 'G', 'H', 'H', 'H', 'H', 'I', 'I', 'I', 'I', 'J', 'J', + 'J', 'J', 'K', 'K', 'K', 'K', 'L', 'L', 'L', 'L', 'M', 'M', 'M', 'M', 'N', 'N', 'N', 'N', 'O', + 'O', 'O', 'O', 'P', 'P', 'P', 'P', 'Q', 'Q', 'Q', 'Q', 'R', 'R', 'R', 'R', 'S', 'S', 'S', 'S', + 'T', 'T', 'T', 'T', 'U', 'U', 'U', 'U', 'V', 'V', 'V', 'V', 'W', 'W', 'W', 'W', 'X', 'X', 'X', + 'X', 'Y', 'Y', 'Y', 'Y', 'Z', 'Z', 'Z', 'Z', 'a', 'a', 'a', 'a', 'b', 'b', 'b', 'b', 'c', 'c', + 'c', 'c', 'd', 'd', 'd', 'd', 'e', 'e', 'e', 'e', 'f', 'f', 'f', 'f', 'g', 'g', 'g', 'g', 'h', + 'h', 'h', 'h', 'i', 'i', 'i', 'i', 'j', 'j', 'j', 'j', 'k', 'k', 'k', 'k', 'l', 'l', 'l', 'l', + 'm', 'm', 'm', 'm', 'n', 'n', 'n', 'n', 'o', 'o', 'o', 'o', 'p', 'p', 'p', 'p', 'q', 'q', 'q', + 'q', 'r', 'r', 'r', 'r', 's', 's', 's', 's', 't', 't', 't', 't', 'u', 'u', 'u', 'u', 'v', 'v', + 'v', 'v', 'w', 'w', 'w', 'w', 'x', 'x', 'x', 'x', 'y', 'y', 'y', 'y', 'z', 'z', 'z', 'z', '0', + '0', '0', '0', '1', '1', '1', '1', '2', '2', '2', '2', '3', '3', '3', '3', '4', '4', '4', '4', + '5', '5', '5', '5', '6', '6', '6', '6', '7', '7', '7', '7', '8', '8', '8', '8', '9', '9', '9', + '9', '+', '+', '+', '+', '/', '/', '/', '/'}; + +std::array constexpr encode_table_1 = { + 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', + 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', + 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', + '5', '6', '7', '8', '9', '+', '/', 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', + 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', + 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', + 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/', 'A', 'B', 'C', 'D', 'E', + 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', + 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', + 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', + '+', '/', 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', + 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', + 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', + '3', '4', '5', '6', '7', '8', '9', '+', '/'}; + +} // namespace detail + +template +inline OutputBuffer encode_into(InputIterator begin, InputIterator end) +{ + typedef std::decay_t input_value_type; + static_assert(std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v); + typedef typename OutputBuffer::value_type output_value_type; + static_assert(std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v); + const size_t binarytextsize = end - begin; + const size_t encodedsize = (binarytextsize / 3 + (binarytextsize % 3 > 0)) << 2; + OutputBuffer encoded(encodedsize, detail::padding_char); + + const uint8_t* bytes = reinterpret_cast(&*begin); + char* currEncoding = reinterpret_cast(&encoded[0]); + + for(size_t i = binarytextsize / 3; i; --i) + { + const uint8_t t1 = *bytes++; + const uint8_t t2 = *bytes++; + const uint8_t t3 = *bytes++; + *currEncoding++ = detail::encode_table_0[t1]; + *currEncoding++ = detail::encode_table_1[((t1 & 0x03) << 4) | ((t2 >> 4) & 0x0F)]; + *currEncoding++ = detail::encode_table_1[((t2 & 0x0F) << 2) | ((t3 >> 6) & 0x03)]; + *currEncoding++ = detail::encode_table_1[t3]; + } + + switch(binarytextsize % 3) + { + case 0: { + break; + } + case 1: { + const uint8_t t1 = bytes[0]; + *currEncoding++ = detail::encode_table_0[t1]; + *currEncoding++ = detail::encode_table_1[(t1 & 0x03) << 4]; + // *currEncoding++ = detail::padding_char; + // *currEncoding++ = detail::padding_char; + break; + } + case 2: { + const uint8_t t1 = bytes[0]; + const uint8_t t2 = bytes[1]; + *currEncoding++ = detail::encode_table_0[t1]; + *currEncoding++ = detail::encode_table_1[((t1 & 0x03) << 4) | ((t2 >> 4) & 0x0F)]; + *currEncoding++ = detail::encode_table_1[(t2 & 0x0F) << 2]; + // *currEncoding++ = detail::padding_char; + break; + } + default: { + throw std::runtime_error{"Invalid base64 encoded data"}; + } + } + + return encoded; +} + +template +inline OutputBuffer encode_into(std::string_view data) +{ + return encode_into(std::begin(data), std::end(data)); +} + +inline std::string to_base64(std::string_view data) +{ + return encode_into(std::begin(data), std::end(data)); +} + +template +inline OutputBuffer decode_into(std::string_view base64Text) +{ + typedef typename OutputBuffer::value_type output_value_type; + static_assert(std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v); + if(base64Text.empty()) + { + return OutputBuffer(); + } + + if((base64Text.size() & 3) != 0) + { + throw std::runtime_error{"Invalid base64 encoded data - Size not divisible by 4"}; + } + + const size_t numPadding = std::count(base64Text.rbegin(), base64Text.rbegin() + 4, '='); + if(numPadding > 2) + { + throw std::runtime_error{"Invalid base64 encoded data - Found more than 2 padding signs"}; + } + + const size_t decodedsize = (base64Text.size() * 3 >> 2) - numPadding; + OutputBuffer decoded(decodedsize, '.'); + + const uint8_t* bytes = reinterpret_cast(&base64Text[0]); + char* currDecoding = reinterpret_cast(&decoded[0]); + + for(size_t i = (base64Text.size() >> 2) - (numPadding != 0); i; --i) + { + const uint8_t t1 = *bytes++; + const uint8_t t2 = *bytes++; + const uint8_t t3 = *bytes++; + const uint8_t t4 = *bytes++; + + const uint32_t d1 = detail::decode_table_0[t1]; + const uint32_t d2 = detail::decode_table_1[t2]; + const uint32_t d3 = detail::decode_table_2[t3]; + const uint32_t d4 = detail::decode_table_3[t4]; + + const uint32_t temp = d1 | d2 | d3 | d4; + + if(temp >= detail::bad_char) + { + throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; + } + + // Use bit_cast instead of union and type punning to avoid + // undefined behaviour risk: + // https://en.wikipedia.org/wiki/Type_punning#Use_of_union + const std::array tempBytes = detail::bit_cast, uint32_t>(temp); + + *currDecoding++ = tempBytes[detail::decidx0]; + *currDecoding++ = tempBytes[detail::decidx1]; + *currDecoding++ = tempBytes[detail::decidx2]; + } + + switch(numPadding) + { + case 0: { + break; + } + case 1: { + const uint8_t t1 = *bytes++; + const uint8_t t2 = *bytes++; + const uint8_t t3 = *bytes++; + + const uint32_t d1 = detail::decode_table_0[t1]; + const uint32_t d2 = detail::decode_table_1[t2]; + const uint32_t d3 = detail::decode_table_2[t3]; + + const uint32_t temp = d1 | d2 | d3; + + if(temp >= detail::bad_char) + { + throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; + } + + // Use bit_cast instead of union and type punning to avoid + // undefined behaviour risk: + // https://en.wikipedia.org/wiki/Type_punning#Use_of_union + const std::array tempBytes = detail::bit_cast, uint32_t>(temp); + *currDecoding++ = tempBytes[detail::decidx0]; + *currDecoding++ = tempBytes[detail::decidx1]; + break; + } + case 2: { + const uint8_t t1 = *bytes++; + const uint8_t t2 = *bytes++; + + const uint32_t d1 = detail::decode_table_0[t1]; + const uint32_t d2 = detail::decode_table_1[t2]; + + const uint32_t temp = d1 | d2; + + if(temp >= detail::bad_char) + { + throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; + } + + const std::array tempBytes = detail::bit_cast, uint32_t>(temp); + *currDecoding++ = tempBytes[detail::decidx0]; + break; + } + default: { + throw std::runtime_error{"Invalid base64 encoded data - Invalid padding number"}; + } + } + + return decoded; +} + +template +inline OutputBuffer decode_into(InputIterator begin, InputIterator end) +{ + typedef std::decay_t input_value_type; + static_assert(std::is_same_v || + std::is_same_v || + std::is_same_v || + std::is_same_v); + std::string_view data(reinterpret_cast(&*begin), end - begin); + return decode_into(data); +} + +inline std::string from_base64(std::string_view data) { return decode_into(data); } + +} // namespace base64 + +#endif // BASE64_HPP_ diff --git a/src/netron_output.cpp b/src/netron_output.cpp index 175f651f700..46406f5ea1e 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -32,6 +32,8 @@ #include #include +#include + namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace { @@ -60,15 +62,24 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; case shape::tuple_type: return 0; - default: { - MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); - } } } +auto make_attribute(migraphx::value val) +{ + json attribute = json::object(); + attribute["name"] = val.get_key(); + auto val_string = val.to(); + val_string = val_string.substr(val_string.find(":") + 1); + attribute["s"] = base64::to_base64(val_string); + attribute["type"] = "STRING"; + return attribute; +} + auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) { json node = json::object(); + // TODO add support for module inputs json input_arr = json::array(); for(instruction_ref input_ins : ins->inputs()) { @@ -77,6 +88,11 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map" + ins_uids.at(ins)); @@ -110,10 +126,9 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map(v); } - // TODO: attributes need to know the type else { - op_attribute_arr.push_back({{"name", v.get_key()}, {"ints", {0}}, {"type", "INTS"}}); + op_attribute_arr.push_back(make_attribute(v)); } }); node["attribute"] = op_attribute_arr; @@ -124,8 +139,7 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) { json lit = json::object(); - lit["dims"] = ins->get_shape().lens(); - // TODO figure out the data types number + lit["dims"] = ins->get_shape().lens(); lit["dataType"] = get_onnx_type(ins->get_shape().type()); lit["name"] = ins_uids.at(ins); // ignoring literal data, setting to "NULL" in base64 @@ -134,6 +148,7 @@ auto make_onnx_json_literal(instruction_ref ins, std::unordered_map Date: Wed, 18 Dec 2024 15:25:45 -0600 Subject: [PATCH 07/21] PR updates, base64 encode only and tests --- src/CMakeLists.txt | 1 + src/base64.cpp | 38 ++ src/include/base64.hpp | 594 -------------------------------- src/include/migraphx/base64.hpp | 39 +++ src/netron_output.cpp | 51 ++- test/base64_test.cpp | 50 +++ 6 files changed, 151 insertions(+), 622 deletions(-) create mode 100644 src/base64.cpp delete mode 100644 src/include/base64.hpp create mode 100644 src/include/migraphx/base64.hpp create mode 100644 test/base64_test.cpp diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index d9e9b92cff1..550bb30bd42 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -37,6 +37,7 @@ add_library(migraphx argument.cpp autocast_fp8.cpp auto_contiguous.cpp + base64.cpp common.cpp common_dims.cpp compile_src.cpp diff --git a/src/base64.cpp b/src/base64.cpp new file mode 100644 index 00000000000..3bf70dd9f33 --- /dev/null +++ b/src/base64.cpp @@ -0,0 +1,38 @@ +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { + +/// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 +std::string b64_encode(const std::string& input) +{ + static const std::string B64chars = + "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"; + const std::size_t len = input.size(); + std::vector res_vec((len + 2) / 3 * 4, '='); + std::size_t j = 0; + std::size_t pad = len % 3; + const size_t last = len - pad; + + for(size_t i = 0; i < last; i += 3) + { + int n = static_cast(input.at(i)) << 16 | static_cast(input.at(i + 1)) << 8 | + input.at(i + 2); + res_vec.at(j++) = B64chars.at(n >> 18); + res_vec[j++] = B64chars.at(n >> 12 & 0x3F); + res_vec[j++] = B64chars.at(n >> 6 & 0x3F); + res_vec[j++] = B64chars.at(n & 0x3F); + } + if(pad) /// Set padding + { + int n = --pad ? static_cast(input.at(last)) << 8 | input.at(last + 1) : input.at(last); + res_vec[j++] = B64chars.at(pad ? n >> 10 & 0x3F : n >> 2); + res_vec[j++] = B64chars.at(pad ? n >> 4 & 0x03F : n << 4 & 0x3F); + res_vec[j++] = pad ? B64chars.at(n << 2 & 0x3F) : '='; + } + return std::string(res_vec.begin(), res_vec.end()); +} + +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx diff --git a/src/include/base64.hpp b/src/include/base64.hpp deleted file mode 100644 index cbc01a1136b..00000000000 --- a/src/include/base64.hpp +++ /dev/null @@ -1,594 +0,0 @@ -#ifndef BASE64_HPP_ -#define BASE64_HPP_ - -#include -#include -#include -#include -#include -#include -#include -#include - -#if defined(__cpp_lib_bit_cast) -#include // For std::bit_cast. -#endif - -namespace base64 { - -namespace detail { - -#if defined(__cpp_lib_bit_cast) -using std::bit_cast; -#else -template -std::enable_if_t && - std::is_trivially_copyable_v, - To> -bit_cast(const From& src) noexcept -{ - static_assert(std::is_trivially_constructible_v, - "This implementation additionally requires " - "destination type to be trivially constructible"); - - To dst; - std::memcpy(&dst, &src, sizeof(To)); - return dst; -} -#endif - -inline constexpr char padding_char{'='}; -inline constexpr uint32_t bad_char{0x01FFFFFF}; - -#if !defined(__LITTLE_ENDIAN__) && !defined(__BIG_ENDIAN__) -#if(defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_BIG_ENDIAN__) || \ - (defined(__BYTE_ORDER) && __BYTE_ORDER == __BIG_ENDIAN) || \ - (defined(_BYTE_ORDER) && _BYTE_ORDER == _BIG_ENDIAN) || \ - (defined(BYTE_ORDER) && BYTE_ORDER == BIG_ENDIAN) || \ - (defined(__sun) && defined(__SVR4) && defined(_BIG_ENDIAN)) || defined(__ARMEB__) || \ - defined(__THUMBEB__) || defined(__AARCH64EB__) || defined(_MIBSEB) || defined(__MIBSEB) || \ - defined(__MIBSEB__) || defined(_M_PPC) -#define __BIG_ENDIAN__ -#elif(defined(__BYTE_ORDER__) && __BYTE_ORDER__ == __ORDER_LITTLE_ENDIAN__) || /* gcc */ \ - (defined(__BYTE_ORDER) && __BYTE_ORDER == __LITTLE_ENDIAN) /* linux header */ \ - || (defined(_BYTE_ORDER) && _BYTE_ORDER == _LITTLE_ENDIAN) || \ - (defined(BYTE_ORDER) && BYTE_ORDER == LITTLE_ENDIAN) /* mingw header */ || \ - (defined(__sun) && defined(__SVR4) && defined(_LITTLE_ENDIAN)) || /* solaris */ \ - defined(__ARMEL__) || defined(__THUMBEL__) || defined(__AARCH64EL__) || defined(_MIPSEL) || \ - defined(__MIPSEL) || defined(__MIPSEL__) || defined(_M_IX86) || defined(_M_X64) || \ - defined(_M_IA64) || /* msvc for intel processors */ \ - defined(_M_ARM) /* msvc code on arm executes in little endian mode */ -#define __LITTLE_ENDIAN__ -#endif -#endif - -#if !defined(__LITTLE_ENDIAN__) & !defined(__BIG_ENDIAN__) -#error "UNKNOWN Platform / endianness. Configure endianness explicitly." -#endif - -#if defined(__LITTLE_ENDIAN__) -std::array constexpr decode_table_0 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x000000f8, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x000000fc, - 0x000000d0, 0x000000d4, 0x000000d8, 0x000000dc, 0x000000e0, 0x000000e4, 0x000000e8, 0x000000ec, - 0x000000f0, 0x000000f4, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00000004, 0x00000008, 0x0000000c, 0x00000010, 0x00000014, 0x00000018, - 0x0000001c, 0x00000020, 0x00000024, 0x00000028, 0x0000002c, 0x00000030, 0x00000034, 0x00000038, - 0x0000003c, 0x00000040, 0x00000044, 0x00000048, 0x0000004c, 0x00000050, 0x00000054, 0x00000058, - 0x0000005c, 0x00000060, 0x00000064, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000068, 0x0000006c, 0x00000070, 0x00000074, 0x00000078, 0x0000007c, 0x00000080, - 0x00000084, 0x00000088, 0x0000008c, 0x00000090, 0x00000094, 0x00000098, 0x0000009c, 0x000000a0, - 0x000000a4, 0x000000a8, 0x000000ac, 0x000000b0, 0x000000b4, 0x000000b8, 0x000000bc, 0x000000c0, - 0x000000c4, 0x000000c8, 0x000000cc, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_1 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000e003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000f003, - 0x00004003, 0x00005003, 0x00006003, 0x00007003, 0x00008003, 0x00009003, 0x0000a003, 0x0000b003, - 0x0000c003, 0x0000d003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00001000, 0x00002000, 0x00003000, 0x00004000, 0x00005000, 0x00006000, - 0x00007000, 0x00008000, 0x00009000, 0x0000a000, 0x0000b000, 0x0000c000, 0x0000d000, 0x0000e000, - 0x0000f000, 0x00000001, 0x00001001, 0x00002001, 0x00003001, 0x00004001, 0x00005001, 0x00006001, - 0x00007001, 0x00008001, 0x00009001, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x0000a001, 0x0000b001, 0x0000c001, 0x0000d001, 0x0000e001, 0x0000f001, 0x00000002, - 0x00001002, 0x00002002, 0x00003002, 0x00004002, 0x00005002, 0x00006002, 0x00007002, 0x00008002, - 0x00009002, 0x0000a002, 0x0000b002, 0x0000c002, 0x0000d002, 0x0000e002, 0x0000f002, 0x00000003, - 0x00001003, 0x00002003, 0x00003003, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_2 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00800f00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00c00f00, - 0x00000d00, 0x00400d00, 0x00800d00, 0x00c00d00, 0x00000e00, 0x00400e00, 0x00800e00, 0x00c00e00, - 0x00000f00, 0x00400f00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00400000, 0x00800000, 0x00c00000, 0x00000100, 0x00400100, 0x00800100, - 0x00c00100, 0x00000200, 0x00400200, 0x00800200, 0x00c00200, 0x00000300, 0x00400300, 0x00800300, - 0x00c00300, 0x00000400, 0x00400400, 0x00800400, 0x00c00400, 0x00000500, 0x00400500, 0x00800500, - 0x00c00500, 0x00000600, 0x00400600, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00800600, 0x00c00600, 0x00000700, 0x00400700, 0x00800700, 0x00c00700, 0x00000800, - 0x00400800, 0x00800800, 0x00c00800, 0x00000900, 0x00400900, 0x00800900, 0x00c00900, 0x00000a00, - 0x00400a00, 0x00800a00, 0x00c00a00, 0x00000b00, 0x00400b00, 0x00800b00, 0x00c00b00, 0x00000c00, - 0x00400c00, 0x00800c00, 0x00c00c00, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_3 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x003e0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x003f0000, - 0x00340000, 0x00350000, 0x00360000, 0x00370000, 0x00380000, 0x00390000, 0x003a0000, 0x003b0000, - 0x003c0000, 0x003d0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00010000, 0x00020000, 0x00030000, 0x00040000, 0x00050000, 0x00060000, - 0x00070000, 0x00080000, 0x00090000, 0x000a0000, 0x000b0000, 0x000c0000, 0x000d0000, 0x000e0000, - 0x000f0000, 0x00100000, 0x00110000, 0x00120000, 0x00130000, 0x00140000, 0x00150000, 0x00160000, - 0x00170000, 0x00180000, 0x00190000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x001a0000, 0x001b0000, 0x001c0000, 0x001d0000, 0x001e0000, 0x001f0000, 0x00200000, - 0x00210000, 0x00220000, 0x00230000, 0x00240000, 0x00250000, 0x00260000, 0x00270000, 0x00280000, - 0x00290000, 0x002a0000, 0x002b0000, 0x002c0000, 0x002d0000, 0x002e0000, 0x002f0000, 0x00300000, - 0x00310000, 0x00320000, 0x00330000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -// TODO fix decoding tables to avoid the need for different indices in big -// endian? -inline constexpr size_t decidx0{0}; -inline constexpr size_t decidx1{1}; -inline constexpr size_t decidx2{2}; - -#elif defined(__BIG_ENDIAN__) - -std::array constexpr decode_table_0 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00f80000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00fc0000, - 0x00d00000, 0x00d40000, 0x00d80000, 0x00dc0000, 0x00e00000, 0x00e40000, 0x00e80000, 0x00ec0000, - 0x00f00000, 0x00f40000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00040000, 0x00080000, 0x000c0000, 0x00100000, 0x00140000, 0x00180000, - 0x001c0000, 0x00200000, 0x00240000, 0x00280000, 0x002c0000, 0x00300000, 0x00340000, 0x00380000, - 0x003c0000, 0x00400000, 0x00440000, 0x00480000, 0x004c0000, 0x00500000, 0x00540000, 0x00580000, - 0x005c0000, 0x00600000, 0x00640000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00680000, 0x006c0000, 0x00700000, 0x00740000, 0x00780000, 0x007c0000, 0x00800000, - 0x00840000, 0x00880000, 0x008c0000, 0x00900000, 0x00940000, 0x00980000, 0x009c0000, 0x00a00000, - 0x00a40000, 0x00a80000, 0x00ac0000, 0x00b00000, 0x00b40000, 0x00b80000, 0x00bc0000, 0x00c00000, - 0x00c40000, 0x00c80000, 0x00cc0000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_1 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0003e000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0003f000, - 0x00034000, 0x00035000, 0x00036000, 0x00037000, 0x00038000, 0x00039000, 0x0003a000, 0x0003b000, - 0x0003c000, 0x0003d000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00001000, 0x00002000, 0x00003000, 0x00004000, 0x00005000, 0x00006000, - 0x00007000, 0x00008000, 0x00009000, 0x0000a000, 0x0000b000, 0x0000c000, 0x0000d000, 0x0000e000, - 0x0000f000, 0x00010000, 0x00011000, 0x00012000, 0x00013000, 0x00014000, 0x00015000, 0x00016000, - 0x00017000, 0x00018000, 0x00019000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x0001a000, 0x0001b000, 0x0001c000, 0x0001d000, 0x0001e000, 0x0001f000, 0x00020000, - 0x00021000, 0x00022000, 0x00023000, 0x00024000, 0x00025000, 0x00026000, 0x00027000, 0x00028000, - 0x00029000, 0x0002a000, 0x0002b000, 0x0002c000, 0x0002d000, 0x0002e000, 0x0002f000, 0x00030000, - 0x00031000, 0x00032000, 0x00033000, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_2 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00000f80, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x00000fc0, - 0x00000d00, 0x00000d40, 0x00000d80, 0x00000dc0, 0x00000e00, 0x00000e40, 0x00000e80, 0x00000ec0, - 0x00000f00, 0x00000f40, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00000040, 0x00000080, 0x000000c0, 0x00000100, 0x00000140, 0x00000180, - 0x000001c0, 0x00000200, 0x00000240, 0x00000280, 0x000002c0, 0x00000300, 0x00000340, 0x00000380, - 0x000003c0, 0x00000400, 0x00000440, 0x00000480, 0x000004c0, 0x00000500, 0x00000540, 0x00000580, - 0x000005c0, 0x00000600, 0x00000640, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000680, 0x000006c0, 0x00000700, 0x00000740, 0x00000780, 0x000007c0, 0x00000800, - 0x00000840, 0x00000880, 0x000008c0, 0x00000900, 0x00000940, 0x00000980, 0x000009c0, 0x00000a00, - 0x00000a40, 0x00000a80, 0x00000ac0, 0x00000b00, 0x00000b40, 0x00000b80, 0x00000bc0, 0x00000c00, - 0x00000c40, 0x00000c80, 0x00000cc0, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -std::array constexpr decode_table_3 = { - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000003e, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x0000003f, - 0x00000034, 0x00000035, 0x00000036, 0x00000037, 0x00000038, 0x00000039, 0x0000003a, 0x0000003b, - 0x0000003c, 0x0000003d, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x00000000, 0x00000001, 0x00000002, 0x00000003, 0x00000004, 0x00000005, 0x00000006, - 0x00000007, 0x00000008, 0x00000009, 0x0000000a, 0x0000000b, 0x0000000c, 0x0000000d, 0x0000000e, - 0x0000000f, 0x00000010, 0x00000011, 0x00000012, 0x00000013, 0x00000014, 0x00000015, 0x00000016, - 0x00000017, 0x00000018, 0x00000019, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x0000001a, 0x0000001b, 0x0000001c, 0x0000001d, 0x0000001e, 0x0000001f, 0x00000020, - 0x00000021, 0x00000022, 0x00000023, 0x00000024, 0x00000025, 0x00000026, 0x00000027, 0x00000028, - 0x00000029, 0x0000002a, 0x0000002b, 0x0000002c, 0x0000002d, 0x0000002e, 0x0000002f, 0x00000030, - 0x00000031, 0x00000032, 0x00000033, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, - 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff, 0x01ffffff}; - -// TODO fix decoding tables to avoid the need for different indices in big -// endian? -inline constexpr size_t decidx0{1}; -inline constexpr size_t decidx1{2}; -inline constexpr size_t decidx2{3}; - -#endif - -std::array constexpr encode_table_0 = { - 'A', 'A', 'A', 'A', 'B', 'B', 'B', 'B', 'C', 'C', 'C', 'C', 'D', 'D', 'D', 'D', 'E', 'E', 'E', - 'E', 'F', 'F', 'F', 'F', 'G', 'G', 'G', 'G', 'H', 'H', 'H', 'H', 'I', 'I', 'I', 'I', 'J', 'J', - 'J', 'J', 'K', 'K', 'K', 'K', 'L', 'L', 'L', 'L', 'M', 'M', 'M', 'M', 'N', 'N', 'N', 'N', 'O', - 'O', 'O', 'O', 'P', 'P', 'P', 'P', 'Q', 'Q', 'Q', 'Q', 'R', 'R', 'R', 'R', 'S', 'S', 'S', 'S', - 'T', 'T', 'T', 'T', 'U', 'U', 'U', 'U', 'V', 'V', 'V', 'V', 'W', 'W', 'W', 'W', 'X', 'X', 'X', - 'X', 'Y', 'Y', 'Y', 'Y', 'Z', 'Z', 'Z', 'Z', 'a', 'a', 'a', 'a', 'b', 'b', 'b', 'b', 'c', 'c', - 'c', 'c', 'd', 'd', 'd', 'd', 'e', 'e', 'e', 'e', 'f', 'f', 'f', 'f', 'g', 'g', 'g', 'g', 'h', - 'h', 'h', 'h', 'i', 'i', 'i', 'i', 'j', 'j', 'j', 'j', 'k', 'k', 'k', 'k', 'l', 'l', 'l', 'l', - 'm', 'm', 'm', 'm', 'n', 'n', 'n', 'n', 'o', 'o', 'o', 'o', 'p', 'p', 'p', 'p', 'q', 'q', 'q', - 'q', 'r', 'r', 'r', 'r', 's', 's', 's', 's', 't', 't', 't', 't', 'u', 'u', 'u', 'u', 'v', 'v', - 'v', 'v', 'w', 'w', 'w', 'w', 'x', 'x', 'x', 'x', 'y', 'y', 'y', 'y', 'z', 'z', 'z', 'z', '0', - '0', '0', '0', '1', '1', '1', '1', '2', '2', '2', '2', '3', '3', '3', '3', '4', '4', '4', '4', - '5', '5', '5', '5', '6', '6', '6', '6', '7', '7', '7', '7', '8', '8', '8', '8', '9', '9', '9', - '9', '+', '+', '+', '+', '/', '/', '/', '/'}; - -std::array constexpr encode_table_1 = { - 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', - 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', - 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', - '5', '6', '7', '8', '9', '+', '/', 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', - 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', - 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', - 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/', 'A', 'B', 'C', 'D', 'E', - 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', - 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', - 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', - '+', '/', 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', - 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', - 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', - '3', '4', '5', '6', '7', '8', '9', '+', '/'}; - -} // namespace detail - -template -inline OutputBuffer encode_into(InputIterator begin, InputIterator end) -{ - typedef std::decay_t input_value_type; - static_assert(std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v); - typedef typename OutputBuffer::value_type output_value_type; - static_assert(std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v); - const size_t binarytextsize = end - begin; - const size_t encodedsize = (binarytextsize / 3 + (binarytextsize % 3 > 0)) << 2; - OutputBuffer encoded(encodedsize, detail::padding_char); - - const uint8_t* bytes = reinterpret_cast(&*begin); - char* currEncoding = reinterpret_cast(&encoded[0]); - - for(size_t i = binarytextsize / 3; i; --i) - { - const uint8_t t1 = *bytes++; - const uint8_t t2 = *bytes++; - const uint8_t t3 = *bytes++; - *currEncoding++ = detail::encode_table_0[t1]; - *currEncoding++ = detail::encode_table_1[((t1 & 0x03) << 4) | ((t2 >> 4) & 0x0F)]; - *currEncoding++ = detail::encode_table_1[((t2 & 0x0F) << 2) | ((t3 >> 6) & 0x03)]; - *currEncoding++ = detail::encode_table_1[t3]; - } - - switch(binarytextsize % 3) - { - case 0: { - break; - } - case 1: { - const uint8_t t1 = bytes[0]; - *currEncoding++ = detail::encode_table_0[t1]; - *currEncoding++ = detail::encode_table_1[(t1 & 0x03) << 4]; - // *currEncoding++ = detail::padding_char; - // *currEncoding++ = detail::padding_char; - break; - } - case 2: { - const uint8_t t1 = bytes[0]; - const uint8_t t2 = bytes[1]; - *currEncoding++ = detail::encode_table_0[t1]; - *currEncoding++ = detail::encode_table_1[((t1 & 0x03) << 4) | ((t2 >> 4) & 0x0F)]; - *currEncoding++ = detail::encode_table_1[(t2 & 0x0F) << 2]; - // *currEncoding++ = detail::padding_char; - break; - } - default: { - throw std::runtime_error{"Invalid base64 encoded data"}; - } - } - - return encoded; -} - -template -inline OutputBuffer encode_into(std::string_view data) -{ - return encode_into(std::begin(data), std::end(data)); -} - -inline std::string to_base64(std::string_view data) -{ - return encode_into(std::begin(data), std::end(data)); -} - -template -inline OutputBuffer decode_into(std::string_view base64Text) -{ - typedef typename OutputBuffer::value_type output_value_type; - static_assert(std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v); - if(base64Text.empty()) - { - return OutputBuffer(); - } - - if((base64Text.size() & 3) != 0) - { - throw std::runtime_error{"Invalid base64 encoded data - Size not divisible by 4"}; - } - - const size_t numPadding = std::count(base64Text.rbegin(), base64Text.rbegin() + 4, '='); - if(numPadding > 2) - { - throw std::runtime_error{"Invalid base64 encoded data - Found more than 2 padding signs"}; - } - - const size_t decodedsize = (base64Text.size() * 3 >> 2) - numPadding; - OutputBuffer decoded(decodedsize, '.'); - - const uint8_t* bytes = reinterpret_cast(&base64Text[0]); - char* currDecoding = reinterpret_cast(&decoded[0]); - - for(size_t i = (base64Text.size() >> 2) - (numPadding != 0); i; --i) - { - const uint8_t t1 = *bytes++; - const uint8_t t2 = *bytes++; - const uint8_t t3 = *bytes++; - const uint8_t t4 = *bytes++; - - const uint32_t d1 = detail::decode_table_0[t1]; - const uint32_t d2 = detail::decode_table_1[t2]; - const uint32_t d3 = detail::decode_table_2[t3]; - const uint32_t d4 = detail::decode_table_3[t4]; - - const uint32_t temp = d1 | d2 | d3 | d4; - - if(temp >= detail::bad_char) - { - throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; - } - - // Use bit_cast instead of union and type punning to avoid - // undefined behaviour risk: - // https://en.wikipedia.org/wiki/Type_punning#Use_of_union - const std::array tempBytes = detail::bit_cast, uint32_t>(temp); - - *currDecoding++ = tempBytes[detail::decidx0]; - *currDecoding++ = tempBytes[detail::decidx1]; - *currDecoding++ = tempBytes[detail::decidx2]; - } - - switch(numPadding) - { - case 0: { - break; - } - case 1: { - const uint8_t t1 = *bytes++; - const uint8_t t2 = *bytes++; - const uint8_t t3 = *bytes++; - - const uint32_t d1 = detail::decode_table_0[t1]; - const uint32_t d2 = detail::decode_table_1[t2]; - const uint32_t d3 = detail::decode_table_2[t3]; - - const uint32_t temp = d1 | d2 | d3; - - if(temp >= detail::bad_char) - { - throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; - } - - // Use bit_cast instead of union and type punning to avoid - // undefined behaviour risk: - // https://en.wikipedia.org/wiki/Type_punning#Use_of_union - const std::array tempBytes = detail::bit_cast, uint32_t>(temp); - *currDecoding++ = tempBytes[detail::decidx0]; - *currDecoding++ = tempBytes[detail::decidx1]; - break; - } - case 2: { - const uint8_t t1 = *bytes++; - const uint8_t t2 = *bytes++; - - const uint32_t d1 = detail::decode_table_0[t1]; - const uint32_t d2 = detail::decode_table_1[t2]; - - const uint32_t temp = d1 | d2; - - if(temp >= detail::bad_char) - { - throw std::runtime_error{"Invalid base64 encoded data - Invalid character"}; - } - - const std::array tempBytes = detail::bit_cast, uint32_t>(temp); - *currDecoding++ = tempBytes[detail::decidx0]; - break; - } - default: { - throw std::runtime_error{"Invalid base64 encoded data - Invalid padding number"}; - } - } - - return decoded; -} - -template -inline OutputBuffer decode_into(InputIterator begin, InputIterator end) -{ - typedef std::decay_t input_value_type; - static_assert(std::is_same_v || - std::is_same_v || - std::is_same_v || - std::is_same_v); - std::string_view data(reinterpret_cast(&*begin), end - begin); - return decode_into(data); -} - -inline std::string from_base64(std::string_view data) { return decode_into(data); } - -} // namespace base64 - -#endif // BASE64_HPP_ diff --git a/src/include/migraphx/base64.hpp b/src/include/migraphx/base64.hpp new file mode 100644 index 00000000000..f574c740c90 --- /dev/null +++ b/src/include/migraphx/base64.hpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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_RTGLIB_BASE64_HPP +#define MIGRAPHX_GUARD_RTGLIB_BASE64_HPP + +#include +#include + +namespace migraphx { +inline namespace MIGRAPHX_INLINE_NS { + +/// encode string to base64 +std::string b64_encode(const std::string& input); + +} // namespace MIGRAPHX_INLINE_NS +} // namespace migraphx + +#endif diff --git a/src/netron_output.cpp b/src/netron_output.cpp index 46406f5ea1e..52338c2561a 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -23,23 +23,17 @@ */ #include - -#include #include - #include #include #include #include - -#include +#include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace { -using json = nlohmann::json; - int get_onnx_type(shape::type_t s_type) { switch(s_type) @@ -62,25 +56,29 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; case shape::tuple_type: return 0; + default: { + MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); + } } } auto make_attribute(migraphx::value val) { - json attribute = json::object(); + value attribute; attribute["name"] = val.get_key(); auto val_string = val.to(); val_string = val_string.substr(val_string.find(":") + 1); - attribute["s"] = base64::to_base64(val_string); + attribute["s"] = b64_encode(val_string); attribute["type"] = "STRING"; return attribute; } +/// Returns a value with the JSON structure needed for a node auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) { - json node = json::object(); + value node; // TODO add support for module inputs - json input_arr = json::array(); + value input_arr; for(instruction_ref input_ins : ins->inputs()) { auto name = input_ins->name(); @@ -98,7 +96,7 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map" + ins_uids.at(ins)); } } - json output_arr = json::array(); + value output_arr; for(instruction_ref output_ins : ins->outputs()) { if(output_ins->name() == "@return") @@ -114,11 +112,11 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_mapname(); - json op_attribute_arr = json::array(); + value op_attribute_arr; auto op_value = ins->get_operator().to_value(); std::for_each(op_value.begin(), op_value.end(), [&](auto v) { std::string attr_key = v.get_key(); - if(attr_key == "code_object") + if(v.is_binary()) { return; } @@ -138,7 +136,7 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) { - json lit = json::object(); + value lit; lit["dims"] = ins->get_shape().lens(); lit["dataType"] = get_onnx_type(ins->get_shape().type()); lit["name"] = ins_uids.at(ins); @@ -151,8 +149,8 @@ auto make_onnx_json_literal(instruction_ref ins, std::unordered_map ins_uids) { - json ret = json::object(); + value ret; shape ins_shape = ins->get_shape(); ret["name"] = ins_uids.at(ins) + "->" + ins_uids.at(out_ins); - json type = {{"tensorType", + value type = {{"tensorType", {{"elemType", get_onnx_type(ins_shape.type())}, {"shape", make_onnx_json_shape(ins_shape)}}}}; ret["type"] = type; @@ -176,10 +174,10 @@ auto make_onnx_json_edge(instruction_ref ins, instruction_ref out_ins, std::unor auto make_onnx_json_in_out(instruction_ref ins, std::unordered_map ins_uids) { - json ret = json::object(); + value ret; shape ins_shape = ins->get_shape(); ret["name"] = ins_uids.at(ins); - json type = {{"tensorType", + value type = {{"tensorType", {{"elemType", get_onnx_type(ins_shape.type())}, {"shape", make_onnx_json_shape(ins_shape)}}}}; ret["type"] = type; @@ -206,18 +204,15 @@ std::unordered_map make_ins_uids(const module& mod std::string make_netron_output(const program& prog) { - json output; + value output; auto prog_value = prog.to_value(); output["irVersion"] = prog_value.at("version").to(); output["producerName"] = "AMDMIGraphX"; output["producerVersion"] = prog_value.at("migraphx_version").to(); for(auto& mod : prog.get_modules()) { - json graph = {{"node", json::array()}, - {"initializer", json::array()}, - {"input", json::array()}, - {"output", json::array()}, - {"valueInfo", json::array()}}; + value graph = { + {"node", {}}, {"initializer", {}}, {"input", {}}, {"output", {}}, {"valueInfo", {}}}; auto ins_uids = make_ins_uids(*mod); for(auto ins = mod->begin(); ins != mod->end(); ++ins) { @@ -253,7 +248,7 @@ std::string make_netron_output(const program& prog) } output["graph"] = graph; } - return output.dump(4); + return to_json_string(output); } } // namespace MIGRAPHX_INLINE_NS diff --git a/test/base64_test.cpp b/test/base64_test.cpp new file mode 100644 index 00000000000..e2e0d14bd7d --- /dev/null +++ b/test/base64_test.cpp @@ -0,0 +1,50 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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 "test.hpp" + +TEST_CASE(base64_encoding) +{ + std::string input_0 = "abc"; + std::string expected_0 = "YWJj"; + EXPECT(migraphx::b64_encode(input_0) == expected_0); + + std::string input_1 = "abcd"; + std::string expected_1 = "YWJjZA=="; + EXPECT(migraphx::b64_encode(input_1) == expected_1); + + std::string input_2 = "convolution"; + std::string expected_2 = "Y29udm9sdXRpb24="; + EXPECT(migraphx::b64_encode(input_2) == expected_2); + + std::string input_3 = "https://www.amd.com/en/products/software/rocm.html"; + std::string expected_3 = "aHR0cHM6Ly93d3cuYW1kLmNvbS9lbi9wcm9kdWN0cy9zb2Z0d2FyZS9yb2NtLmh0bWw="; + EXPECT(migraphx::b64_encode(input_3) == expected_3); + + std::string input_4 = "{1, 3, 7, 9}"; + std::string expected_4 = "ezEsIDMsIDcsIDl9"; + EXPECT(migraphx::b64_encode(input_4) == expected_4); +} + +int main(int argc, const char* argv[]) { test::run(argc, argv); } From 755ce18b4d99f852cc51d9d503b57a11393007ea Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 18 Dec 2024 17:21:33 -0600 Subject: [PATCH 08/21] Fix the bug add RFC tests --- requirements.txt | 3 +- src/base64.cpp | 49 +++++++++++------- test/base64_test.cpp | 118 +++++++++++++++++++++++++++++++++++++++++++ 3 files changed, 152 insertions(+), 18 deletions(-) diff --git a/requirements.txt b/requirements.txt index 4daf866ed52..dfe38237be3 100755 --- a/requirements.txt +++ b/requirements.txt @@ -28,4 +28,5 @@ pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On ROCm/composable_kernel@b7775add2d28251674d81e220cd4a857b90b997a -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On -ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On \ No newline at end of file +ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On +tobiaslocker/base64 diff --git a/src/base64.cpp b/src/base64.cpp index 3bf70dd9f33..fdc9857c251 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -1,38 +1,53 @@ #include #include +#include +#include namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { +namespace{ + +std::array constexpr B64chars{'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', + 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', + 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', + '6', '7', '8', '9', '+', '/'}; + /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -std::string b64_encode(const std::string& input) +const std::string b64_encode(const void* data, const size_t &len) { - static const std::string B64chars = - "ABCDEFGHIJKLMNOPQRSTUVWXYZabcdefghijklmnopqrstuvwxyz0123456789+/"; - const std::size_t len = input.size(); std::vector res_vec((len + 2) / 3 * 4, '='); - std::size_t j = 0; - std::size_t pad = len % 3; + const unsigned char *p = static_cast(data); + size_t j = 0, pad = len % 3; const size_t last = len - pad; - for(size_t i = 0; i < last; i += 3) + #pragma clang diagnostic push + #pragma clang diagnostic ignored "-Wunsafe-buffer-usage" + for (size_t i = 0; i < last; i += 3) { - int n = static_cast(input.at(i)) << 16 | static_cast(input.at(i + 1)) << 8 | - input.at(i + 2); + unsigned n = static_cast(p[i]) << 16 | static_cast(p[i + 1]) << 8 | p[i + 2]; res_vec.at(j++) = B64chars.at(n >> 18); - res_vec[j++] = B64chars.at(n >> 12 & 0x3F); - res_vec[j++] = B64chars.at(n >> 6 & 0x3F); - res_vec[j++] = B64chars.at(n & 0x3F); + res_vec.at(j++) = B64chars.at(n >> 12 & 0x3F); + res_vec.at(j++) = B64chars.at(n >> 6 & 0x3F); + res_vec.at(j++) = B64chars.at(n & 0x3F); } - if(pad) /// Set padding + if (pad) /// Set padding { - int n = --pad ? static_cast(input.at(last)) << 8 | input.at(last + 1) : input.at(last); - res_vec[j++] = B64chars.at(pad ? n >> 10 & 0x3F : n >> 2); - res_vec[j++] = B64chars.at(pad ? n >> 4 & 0x03F : n << 4 & 0x3F); - res_vec[j++] = pad ? B64chars.at(n << 2 & 0x3F) : '='; + unsigned n = --pad ? static_cast(p[last]) << 8 | p[last + 1] : p[last]; + res_vec.at(j++) = B64chars.at(pad ? n >> 10 & 0x3F : n >> 2); + res_vec.at(j++) = B64chars.at(pad ? n >> 4 & 0x03F : n << 4 & 0x3F); + res_vec.at(j++) = pad ? B64chars.at(n << 2 & 0x3F) : '='; } + #pragma clang diagnostic pop return std::string(res_vec.begin(), res_vec.end()); } +} + +std::string b64_encode(const std::string& str) +{ + return b64_encode(str.c_str(), str.size()); +} + } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/test/base64_test.cpp b/test/base64_test.cpp index e2e0d14bd7d..348244ef85c 100644 --- a/test/base64_test.cpp +++ b/test/base64_test.cpp @@ -47,4 +47,122 @@ TEST_CASE(base64_encoding) EXPECT(migraphx::b64_encode(input_4) == expected_4); } +TEST_CASE(base64_RFC_test_vectors) +{ + std::string input_0 = ""; + std::string expected_0 = ""; + EXPECT(migraphx::b64_encode(input_0) == expected_0); + + std::string input_1 = "f"; + std::string expected_1 = "Zg=="; + EXPECT(migraphx::b64_encode(input_1) == expected_1); + + std::string input_2 = "fo"; + std::string expected_2 = "Zm8="; + EXPECT(migraphx::b64_encode(input_2) == expected_2); + + std::string input_3 = "foo"; + std::string expected_3 = "Zm9v"; + EXPECT(migraphx::b64_encode(input_3) == expected_3); + + std::string input_4 = "foob"; + std::string expected_4 = "Zm9vYg=="; + EXPECT(migraphx::b64_encode(input_4) == expected_4); + + std::string input_5 = "fooba"; + std::string expected_5 = "Zm9vYmE="; + EXPECT(migraphx::b64_encode(input_5) == expected_5); + + std::string input_6 = "foobar"; + std::string expected_6 = "Zm9vYmFy"; + EXPECT(migraphx::b64_encode(input_6) == expected_6); +} + +// Following tests altered from https://github.com/tobiaslocker/base64/blob/master/test/base64_tests.cpp +TEST_CASE(base64_encodes_empty) +{ + std::string const expected{}; + std::string const actual{migraphx::b64_encode({})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_three_bytes_zeros) +{ + std::array const input{0x00, 0x00, 0x00}; + auto const expected{"AAAA"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_three_bytes_random) +{ + std::array const input{0xFE, 0xE9, 0x72}; + std::string const expected{"/uly"}; + std::string const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_two_bytes) +{ + std::array const input{0x00, 0x00}; + auto const expected{"AAA="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_one_byte) +{ + std::array const input{0x00}; + auto const expected{"AA=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_four_bytes) { + std::array const input{0x74, 0x68, 0x65, 0x20}; + auto const expected{"dGhlIA=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_five_bytes) { + std::array const input{0x20, 0x62, 0x72, 0x6f, 0x77}; + auto const expected{"IGJyb3c="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_six_bytes) { + std::array const input{0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73}; + auto const expected{"IGp1bXBz"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_BrownFox) { + std::array const input{ + 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x20, 0x62, + 0x72, 0x6f, 0x77, 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, + 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, + 0x65, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; + + auto const expected{ + "dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + +TEST_CASE(base64_encodes_EncodesBrownFastFoxNullInMiddle) { + std::array const input{ + 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x21, 0x20, 0x62, + 0x72, 0x6f, 0x77, 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, + 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, 0x65, 0x00, + 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; + + auto const expected{ + "dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } From d2b16cce4f2e110451ded12f5077eb70e481f202 Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 18 Dec 2024 17:21:47 -0600 Subject: [PATCH 09/21] formatting --- src/base64.cpp | 37 +++++++------ test/base64_test.cpp | 122 ++++++++++++++++++++++--------------------- 2 files changed, 80 insertions(+), 79 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index fdc9857c251..37fd3826956 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -6,48 +6,47 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { -namespace{ +namespace { -std::array constexpr B64chars{'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', - 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', - 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', - '6', '7', '8', '9', '+', '/'}; +std::array constexpr B64chars{ + 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', + 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', + 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', + 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'}; /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -const std::string b64_encode(const void* data, const size_t &len) +const std::string b64_encode(const void* data, const size_t& len) { std::vector res_vec((len + 2) / 3 * 4, '='); - const unsigned char *p = static_cast(data); + const unsigned char* p = static_cast(data); size_t j = 0, pad = len % 3; const size_t last = len - pad; - #pragma clang diagnostic push - #pragma clang diagnostic ignored "-Wunsafe-buffer-usage" - for (size_t i = 0; i < last; i += 3) +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wunsafe-buffer-usage" + for(size_t i = 0; i < last; i += 3) { - unsigned n = static_cast(p[i]) << 16 | static_cast(p[i + 1]) << 8 | p[i + 2]; + unsigned n = + static_cast(p[i]) << 16 | static_cast(p[i + 1]) << 8 | p[i + 2]; res_vec.at(j++) = B64chars.at(n >> 18); res_vec.at(j++) = B64chars.at(n >> 12 & 0x3F); res_vec.at(j++) = B64chars.at(n >> 6 & 0x3F); res_vec.at(j++) = B64chars.at(n & 0x3F); } - if (pad) /// Set padding + if(pad) /// Set padding { - unsigned n = --pad ? static_cast(p[last]) << 8 | p[last + 1] : p[last]; + unsigned n = --pad ? static_cast(p[last]) << 8 | p[last + 1] : p[last]; res_vec.at(j++) = B64chars.at(pad ? n >> 10 & 0x3F : n >> 2); res_vec.at(j++) = B64chars.at(pad ? n >> 4 & 0x03F : n << 4 & 0x3F); res_vec.at(j++) = pad ? B64chars.at(n << 2 & 0x3F) : '='; } - #pragma clang diagnostic pop +#pragma clang diagnostic pop return std::string(res_vec.begin(), res_vec.end()); } -} +} // namespace -std::string b64_encode(const std::string& str) -{ - return b64_encode(str.c_str(), str.size()); -} +std::string b64_encode(const std::string& str) { return b64_encode(str.c_str(), str.size()); } } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/test/base64_test.cpp b/test/base64_test.cpp index 348244ef85c..f818c50712f 100644 --- a/test/base64_test.cpp +++ b/test/base64_test.cpp @@ -64,7 +64,7 @@ TEST_CASE(base64_RFC_test_vectors) std::string input_3 = "foo"; std::string expected_3 = "Zm9v"; EXPECT(migraphx::b64_encode(input_3) == expected_3); - + std::string input_4 = "foob"; std::string expected_4 = "Zm9vYg=="; EXPECT(migraphx::b64_encode(input_4) == expected_4); @@ -78,91 +78,93 @@ TEST_CASE(base64_RFC_test_vectors) EXPECT(migraphx::b64_encode(input_6) == expected_6); } -// Following tests altered from https://github.com/tobiaslocker/base64/blob/master/test/base64_tests.cpp +// Following tests altered from +// https://github.com/tobiaslocker/base64/blob/master/test/base64_tests.cpp TEST_CASE(base64_encodes_empty) { - std::string const expected{}; - std::string const actual{migraphx::b64_encode({})}; - EXPECT(expected == actual); + std::string const expected{}; + std::string const actual{migraphx::b64_encode({})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_three_bytes_zeros) +TEST_CASE(base64_encodes_three_bytes_zeros) { - std::array const input{0x00, 0x00, 0x00}; - auto const expected{"AAAA"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); + std::array const input{0x00, 0x00, 0x00}; + auto const expected{"AAAA"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_three_bytes_random) +TEST_CASE(base64_encodes_three_bytes_random) { - std::array const input{0xFE, 0xE9, 0x72}; - std::string const expected{"/uly"}; - std::string const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); + std::array const input{0xFE, 0xE9, 0x72}; + std::string const expected{"/uly"}; + std::string const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } TEST_CASE(base64_encodes_two_bytes) { - std::array const input{0x00, 0x00}; - auto const expected{"AAA="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); + std::array const input{0x00, 0x00}; + auto const expected{"AAA="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } TEST_CASE(base64_encodes_one_byte) { - std::array const input{0x00}; - auto const expected{"AA=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); + std::array const input{0x00}; + auto const expected{"AA=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_four_bytes) { - std::array const input{0x74, 0x68, 0x65, 0x20}; - auto const expected{"dGhlIA=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); +TEST_CASE(base64_encodes_four_bytes) +{ + std::array const input{0x74, 0x68, 0x65, 0x20}; + auto const expected{"dGhlIA=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_five_bytes) { - std::array const input{0x20, 0x62, 0x72, 0x6f, 0x77}; - auto const expected{"IGJyb3c="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); +TEST_CASE(base64_encodes_five_bytes) +{ + std::array const input{0x20, 0x62, 0x72, 0x6f, 0x77}; + auto const expected{"IGJyb3c="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_six_bytes) { - std::array const input{0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73}; - auto const expected{"IGp1bXBz"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); +TEST_CASE(base64_encodes_six_bytes) +{ + std::array const input{0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73}; + auto const expected{"IGp1bXBz"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_BrownFox) { - std::array const input{ - 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x20, 0x62, - 0x72, 0x6f, 0x77, 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, - 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, - 0x65, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; - - auto const expected{ - "dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); +TEST_CASE(base64_encodes_BrownFox) +{ + std::array const input{ + 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x20, 0x62, 0x72, 0x6f, 0x77, 0x6e, + 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, + 0x20, 0x74, 0x68, 0x65, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; + + auto const expected{"dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } -TEST_CASE(base64_encodes_EncodesBrownFastFoxNullInMiddle) { - std::array const input{ - 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x21, 0x20, 0x62, - 0x72, 0x6f, 0x77, 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, - 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, 0x65, 0x00, - 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; - - auto const expected{ - "dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; - EXPECT(expected == actual); +TEST_CASE(base64_encodes_EncodesBrownFastFoxNullInMiddle) +{ + std::array const input{ + 0x74, 0x68, 0x65, 0x20, 0x71, 0x75, 0x69, 0x63, 0x6b, 0x21, 0x20, 0x62, 0x72, 0x6f, 0x77, + 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, + 0x72, 0x20, 0x74, 0x68, 0x65, 0x00, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; + + auto const expected{"dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; + auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + EXPECT(expected == actual); } int main(int argc, const char* argv[]) { test::run(argc, argv); } From 8bde3efdb25d69bd070cd7a8a6b4738e7b88c1e5 Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 18 Dec 2024 17:24:09 -0600 Subject: [PATCH 10/21] Remove requirements change --- requirements.txt | 1 - 1 file changed, 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index dfe38237be3..d3edaa80070 100755 --- a/requirements.txt +++ b/requirements.txt @@ -29,4 +29,3 @@ msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On ROCm/composable_kernel@b7775add2d28251674d81e220cd4a857b90b997a -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On -tobiaslocker/base64 From 18b373f22a583cca58da618ee5231ae367cd7ee6 Mon Sep 17 00:00:00 2001 From: charlie Date: Wed, 18 Dec 2024 17:24:09 -0600 Subject: [PATCH 11/21] Remove requirements change --- requirements.txt | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/requirements.txt b/requirements.txt index dfe38237be3..72b5429e0c7 100755 --- a/requirements.txt +++ b/requirements.txt @@ -27,6 +27,5 @@ ROCm/half@rocm-5.6.0 pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On -ROCm/composable_kernel@b7775add2d28251674d81e220cd4a857b90b997a -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On -ROCm/rocMLIR@13065c4b3a216e1b13dfb8f746b8a0d421f124e8 -DBUILD_FAT_LIBROCKCOMPILER=On -tobiaslocker/base64 +ROCm/composable_kernel@57cdd70b7cb14e5e3b60cd9a5f96ba8dc343763e -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On +ROCm/rocMLIR@45d830dbbc15fe84c41d95585f526a50719020eb -DBUILD_FAT_LIBROCKCOMPILER=On From f89dfd576aa4b9d912abc29baf21f2c416d22b38 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 09:34:37 -0600 Subject: [PATCH 12/21] fix tidy and another base64 refactor --- src/base64.cpp | 43 ++++++++++++++++++++++++------------------- src/netron_output.cpp | 9 +++++---- 2 files changed, 29 insertions(+), 23 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index 37fd3826956..369513df690 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -7,6 +7,7 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace { +typedef unsigned char byte; std::array constexpr B64chars{ 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', @@ -15,38 +16,42 @@ std::array constexpr B64chars{ 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'}; /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -const std::string b64_encode(const void* data, const size_t& len) +std::string b64_encode(const std::vector buf) { - std::vector res_vec((len + 2) / 3 * 4, '='); - const unsigned char* p = static_cast(data); - size_t j = 0, pad = len % 3; - const size_t last = len - pad; + std::size_t len = buf.size(); + std::vector res_vec((len + 2) / 3 * 4, '='); + std::size_t j = 0; + std::size_t pad_cond = len % 3; + const size_t last = len - pad_cond; -#pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wunsafe-buffer-usage" for(size_t i = 0; i < last; i += 3) { - unsigned n = - static_cast(p[i]) << 16 | static_cast(p[i + 1]) << 8 | p[i + 2]; - res_vec.at(j++) = B64chars.at(n >> 18); - res_vec.at(j++) = B64chars.at(n >> 12 & 0x3F); - res_vec.at(j++) = B64chars.at(n >> 6 & 0x3F); + std::size_t n = static_cast(buf.at(i)) << 16u | + static_cast(buf.at(i + 1)) << 8u | + static_cast(buf.at(i + 2)); + res_vec.at(j++) = B64chars.at(n >> 18u); + res_vec.at(j++) = B64chars.at(n >> 12u & 0x3F); + res_vec.at(j++) = B64chars.at(n >> 6u & 0x3F); res_vec.at(j++) = B64chars.at(n & 0x3F); } - if(pad) /// Set padding + if(pad_cond) /// Set padding { - unsigned n = --pad ? static_cast(p[last]) << 8 | p[last + 1] : p[last]; - res_vec.at(j++) = B64chars.at(pad ? n >> 10 & 0x3F : n >> 2); - res_vec.at(j++) = B64chars.at(pad ? n >> 4 & 0x03F : n << 4 & 0x3F); - res_vec.at(j++) = pad ? B64chars.at(n << 2 & 0x3F) : '='; + std::size_t n = --pad_cond ? static_cast(buf.at(last)) << 8u | + static_cast(buf.at(last + 1)) + : static_cast(buf.at(last)); + res_vec.at(j++) = B64chars.at(pad_cond ? n >> 10u & 0x3F : n >> 2u); + res_vec.at(j++) = B64chars.at(pad_cond ? n >> 4u & 0x03F : n << 4u & 0x3F); + res_vec.at(j++) = pad_cond ? B64chars.at(n << 2u & 0x3F) : '='; } -#pragma clang diagnostic pop return std::string(res_vec.begin(), res_vec.end()); } } // namespace -std::string b64_encode(const std::string& str) { return b64_encode(str.c_str(), str.size()); } +std::string b64_encode(const std::string& str) +{ + return b64_encode(std::vector(str.begin(), str.end())); +} } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/netron_output.cpp b/src/netron_output.cpp index 52338c2561a..a62eda037ab 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -151,10 +151,11 @@ auto make_onnx_json_shape(const shape& s) { value ret; value dim; - for(std::size_t len : s.lens()) - { - dim.push_back({{"dimValue", len}}); - } + auto shape_lens = s.lens(); + std::transform(shape_lens.begin(), + shape_lens.end(), + std::back_inserter(dim), + [](std::size_t len) { return len; }); ret["dim"] = dim; return ret; } From 965644a072b9c1e60e167474c87e80846647de44 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 09:37:54 -0600 Subject: [PATCH 13/21] Licensing --- src/base64.cpp | 23 +++++++++++++++++++++++ src/include/migraphx/netron_output.hpp | 2 +- 2 files changed, 24 insertions(+), 1 deletion(-) diff --git a/src/base64.cpp b/src/base64.cpp index 369513df690..348811b2c60 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -1,3 +1,26 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 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 diff --git a/src/include/migraphx/netron_output.hpp b/src/include/migraphx/netron_output.hpp index c07d5f50af9..fb355a2d9f5 100644 --- a/src/include/migraphx/netron_output.hpp +++ b/src/include/migraphx/netron_output.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 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 From cf9271b4dd0ff9a04a1568b1088605bba89d192b Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 09:54:25 -0600 Subject: [PATCH 14/21] more tidy fixes --- src/base64.cpp | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index 348811b2c60..cd9650b9f2a 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -30,16 +30,16 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace { -typedef unsigned char byte; +using byte = unsigned char; -std::array constexpr B64chars{ +std::array constexpr b64_chars{ 'A', 'B', 'C', 'D', 'E', 'F', 'G', 'H', 'I', 'J', 'K', 'L', 'M', 'N', 'O', 'P', 'Q', 'R', 'S', 'T', 'U', 'V', 'W', 'X', 'Y', 'Z', 'a', 'b', 'c', 'd', 'e', 'f', 'g', 'h', 'i', 'j', 'k', 'l', 'm', 'n', 'o', 'p', 'q', 'r', 's', 't', 'u', 'v', 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'}; /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -std::string b64_encode(const std::vector buf) +std::string b64_encode(const std::vector& buf) { std::size_t len = buf.size(); std::vector res_vec((len + 2) / 3 * 4, '='); @@ -52,19 +52,19 @@ std::string b64_encode(const std::vector buf) std::size_t n = static_cast(buf.at(i)) << 16u | static_cast(buf.at(i + 1)) << 8u | static_cast(buf.at(i + 2)); - res_vec.at(j++) = B64chars.at(n >> 18u); - res_vec.at(j++) = B64chars.at(n >> 12u & 0x3F); - res_vec.at(j++) = B64chars.at(n >> 6u & 0x3F); - res_vec.at(j++) = B64chars.at(n & 0x3F); + res_vec.at(j++) = b64_chars.at(n >> 18u); + res_vec.at(j++) = b64_chars.at(n >> 12u & 0x3Fu); + res_vec.at(j++) = b64_chars.at(n >> 6u & 0x3Fu); + res_vec.at(j++) = b64_chars.at(n & 0x3Fu); } if(pad_cond) /// Set padding { std::size_t n = --pad_cond ? static_cast(buf.at(last)) << 8u | static_cast(buf.at(last + 1)) : static_cast(buf.at(last)); - res_vec.at(j++) = B64chars.at(pad_cond ? n >> 10u & 0x3F : n >> 2u); - res_vec.at(j++) = B64chars.at(pad_cond ? n >> 4u & 0x03F : n << 4u & 0x3F); - res_vec.at(j++) = pad_cond ? B64chars.at(n << 2u & 0x3F) : '='; + res_vec.at(j++) = b64_chars.at(pad_cond ? n >> 10u & 0x3Fu : n >> 2u); + res_vec.at(j++) = b64_chars.at(pad_cond ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); + res_vec.at(j++) = pad_cond ? b64_chars.at(n << 2u & 0x3Fu) : '='; } return std::string(res_vec.begin(), res_vec.end()); } From a195d07166e25660917128f5331677bee4ffab9a Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 11:40:18 -0600 Subject: [PATCH 15/21] Codecov ignore, tidy fixes --- codecov.yml | 1 + src/base64.cpp | 17 +++++++++-------- src/include/migraphx/base64.hpp | 2 +- src/netron_output.cpp | 7 +++---- 4 files changed, 14 insertions(+), 13 deletions(-) diff --git a/codecov.yml b/codecov.yml index 03abe2daeb2..9f2569b0669 100644 --- a/codecov.yml +++ b/codecov.yml @@ -2,3 +2,4 @@ ignore: - "test/" - "src/driver" - "build/" + - "src/netron_output.cpp" diff --git a/src/base64.cpp b/src/base64.cpp index cd9650b9f2a..b5cb00e4563 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -57,16 +57,17 @@ std::string b64_encode(const std::vector& buf) res_vec.at(j++) = b64_chars.at(n >> 6u & 0x3Fu); res_vec.at(j++) = b64_chars.at(n & 0x3Fu); } - if(pad_cond) /// Set padding + // Set padding + if(pad_cond != 0) { - std::size_t n = --pad_cond ? static_cast(buf.at(last)) << 8u | - static_cast(buf.at(last + 1)) - : static_cast(buf.at(last)); - res_vec.at(j++) = b64_chars.at(pad_cond ? n >> 10u & 0x3Fu : n >> 2u); - res_vec.at(j++) = b64_chars.at(pad_cond ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); - res_vec.at(j++) = pad_cond ? b64_chars.at(n << 2u & 0x3Fu) : '='; + std::size_t n = --pad_cond != 0 ? static_cast(buf.at(last)) << 8u | + static_cast(buf.at(last + 1)) + : static_cast(buf.at(last)); + res_vec.at(j++) = b64_chars.at(pad_cond != 0 ? n >> 10u & 0x3Fu : n >> 2u); + res_vec.at(j++) = b64_chars.at(pad_cond != 0 ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); + res_vec.at(j++) = pad_cond != 0 ? b64_chars.at(n << 2u & 0x3Fu) : '='; } - return std::string(res_vec.begin(), res_vec.end()); + return {res_vec.begin(), res_vec.end()}; } } // namespace diff --git a/src/include/migraphx/base64.hpp b/src/include/migraphx/base64.hpp index f574c740c90..7903e4453b7 100644 --- a/src/include/migraphx/base64.hpp +++ b/src/include/migraphx/base64.hpp @@ -31,7 +31,7 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { /// encode string to base64 -std::string b64_encode(const std::string& input); +std::string b64_encode(const std::string& str); } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/netron_output.cpp b/src/netron_output.cpp index a62eda037ab..c864add9d55 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -34,6 +34,7 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { namespace { +// from https://onnx.ai/onnx/intro/concepts.html int get_onnx_type(shape::type_t s_type) { switch(s_type) @@ -56,13 +57,11 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; case shape::tuple_type: return 0; - default: { MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); - } } } -auto make_attribute(migraphx::value val) +auto make_attribute(const migraphx::value val) { value attribute; attribute["name"] = val.get_key(); @@ -115,7 +114,7 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_mapget_operator().to_value(); std::for_each(op_value.begin(), op_value.end(), [&](auto v) { - std::string attr_key = v.get_key(); + const std::string& attr_key = v.get_key(); if(v.is_binary()) { return; From a7587445bb2f23eef67cddcf7a594aa5af9a72a8 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 12:19:38 -0600 Subject: [PATCH 16/21] even more tidy fixes --- src/netron_output.cpp | 78 +++++++++++++++++++++++-------------------- test/base64_test.cpp | 32 +++++++++--------- 2 files changed, 58 insertions(+), 52 deletions(-) diff --git a/src/netron_output.cpp b/src/netron_output.cpp index c864add9d55..85fc1eb9b41 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -61,7 +61,7 @@ int get_onnx_type(shape::type_t s_type) } } -auto make_attribute(const migraphx::value val) +auto make_attribute(const migraphx::value& val) { value attribute; attribute["name"] = val.get_key(); @@ -200,6 +200,46 @@ std::unordered_map make_ins_uids(const module& mod return ret; } +value make_graph(const module* mod) +{ + value graph = { + {"node", {}}, {"initializer", {}}, {"input", {}}, {"output", {}}, {"valueInfo", {}}}; + auto ins_uids = make_ins_uids(*mod); + for(auto ins = mod->begin(); ins != mod->end(); ++ins) + { + const auto& name = ins->name(); + if(name == "@literal") + { + graph["initializer"].push_back(make_onnx_json_literal(ins, ins_uids)); + } + else if(name == "@param") + { + graph["input"].push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else if(name == "@return") + { + graph["output"].push_back(make_onnx_json_in_out(ins, ins_uids)); + } + else if(name.find("hip::hip_allocate_memory") != std::string::npos) + { + continue; + } + else + { + graph["node"].push_back(make_onnx_json_node(ins, ins_uids)); + const auto& outputs = ins->outputs(); + for(auto out_ins : outputs) + { + if(out_ins->name() != "@return") + { + graph["valueInfo"].push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); + } + } + } + } + return graph; +} + } // namespace std::string make_netron_output(const program& prog) @@ -211,41 +251,7 @@ std::string make_netron_output(const program& prog) output["producerVersion"] = prog_value.at("migraphx_version").to(); for(auto& mod : prog.get_modules()) { - value graph = { - {"node", {}}, {"initializer", {}}, {"input", {}}, {"output", {}}, {"valueInfo", {}}}; - auto ins_uids = make_ins_uids(*mod); - for(auto ins = mod->begin(); ins != mod->end(); ++ins) - { - const auto& name = ins->name(); - if(name == "@literal") - { - graph["initializer"].push_back(make_onnx_json_literal(ins, ins_uids)); - } - else if(name == "@param") - { - graph["input"].push_back(make_onnx_json_in_out(ins, ins_uids)); - } - else if(name == "@return") - { - graph["output"].push_back(make_onnx_json_in_out(ins, ins_uids)); - } - else if(name.find("hip::hip_allocate_memory") != std::string::npos) - { - continue; - } - else - { - graph["node"].push_back(make_onnx_json_node(ins, ins_uids)); - const auto& outputs = ins->outputs(); - for(auto out_ins : outputs) - { - if(out_ins->name() != "@return") - { - graph["valueInfo"].push_back(make_onnx_json_edge(ins, out_ins, ins_uids)); - } - } - } - } + auto graph = make_graph(mod); output["graph"] = graph; } return to_json_string(output); diff --git a/test/base64_test.cpp b/test/base64_test.cpp index f818c50712f..9448c54eb75 100644 --- a/test/base64_test.cpp +++ b/test/base64_test.cpp @@ -90,8 +90,8 @@ TEST_CASE(base64_encodes_empty) TEST_CASE(base64_encodes_three_bytes_zeros) { std::array const input{0x00, 0x00, 0x00}; - auto const expected{"AAAA"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"AAAA"}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -106,40 +106,40 @@ TEST_CASE(base64_encodes_three_bytes_random) TEST_CASE(base64_encodes_two_bytes) { std::array const input{0x00, 0x00}; - auto const expected{"AAA="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"AAA="}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } TEST_CASE(base64_encodes_one_byte) { std::array const input{0x00}; - auto const expected{"AA=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"AA=="}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } TEST_CASE(base64_encodes_four_bytes) { std::array const input{0x74, 0x68, 0x65, 0x20}; - auto const expected{"dGhlIA=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"dGhlIA=="}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } TEST_CASE(base64_encodes_five_bytes) { std::array const input{0x20, 0x62, 0x72, 0x6f, 0x77}; - auto const expected{"IGJyb3c="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"IGJyb3c="}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } TEST_CASE(base64_encodes_six_bytes) { std::array const input{0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73}; - auto const expected{"IGp1bXBz"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"IGp1bXBz"}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -150,8 +150,8 @@ TEST_CASE(base64_encodes_BrownFox) 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, 0x65, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; - auto const expected{"dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -162,8 +162,8 @@ TEST_CASE(base64_encodes_EncodesBrownFastFoxNullInMiddle) 0x6e, 0x20, 0x66, 0x6f, 0x78, 0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73, 0x20, 0x6f, 0x76, 0x65, 0x72, 0x20, 0x74, 0x68, 0x65, 0x00, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; - auto const expected{"dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; - auto const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string expected{"dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; + std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } From 7a2ee8d4e98ae0f31539930893f3ee8266098960 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 16:35:22 -0600 Subject: [PATCH 17/21] PR comments --- src/base64.cpp | 27 ++++++------ src/include/migraphx/base64.hpp | 2 +- src/netron_output.cpp | 4 +- test/base64_test.cpp | 74 ++++++++++----------------------- 4 files changed, 37 insertions(+), 70 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index b5cb00e4563..da0a7200b1d 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -39,13 +39,13 @@ std::array constexpr b64_chars{ 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'}; /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -std::string b64_encode(const std::vector& buf) +std::string encode(const std::string& buf) { std::size_t len = buf.size(); - std::vector res_vec((len + 2) / 3 * 4, '='); + std::string res_vec((len + 2) / 3 * 4, '='); std::size_t j = 0; - std::size_t pad_cond = len % 3; - const size_t last = len - pad_cond; + std::size_t remaining = len % 3; + const size_t last = len - remaining; for(size_t i = 0; i < last; i += 3) { @@ -58,24 +58,21 @@ std::string b64_encode(const std::vector& buf) res_vec.at(j++) = b64_chars.at(n & 0x3Fu); } // Set padding - if(pad_cond != 0) + if(remaining != 0) { - std::size_t n = --pad_cond != 0 ? static_cast(buf.at(last)) << 8u | - static_cast(buf.at(last + 1)) - : static_cast(buf.at(last)); - res_vec.at(j++) = b64_chars.at(pad_cond != 0 ? n >> 10u & 0x3Fu : n >> 2u); - res_vec.at(j++) = b64_chars.at(pad_cond != 0 ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); - res_vec.at(j++) = pad_cond != 0 ? b64_chars.at(n << 2u & 0x3Fu) : '='; + std::size_t n = --remaining != 0 ? static_cast(buf.at(last)) << 8u | + static_cast(buf.at(last + 1)) + : static_cast(buf.at(last)); + res_vec.at(j++) = b64_chars.at(remaining != 0 ? n >> 10u & 0x3Fu : n >> 2u); + res_vec.at(j++) = b64_chars.at(remaining != 0 ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); + res_vec.at(j++) = remaining != 0 ? b64_chars.at(n << 2u & 0x3Fu) : '='; } return {res_vec.begin(), res_vec.end()}; } } // namespace -std::string b64_encode(const std::string& str) -{ - return b64_encode(std::vector(str.begin(), str.end())); -} +std::string base64_encode(const std::string& str) { return encode(str); } } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/include/migraphx/base64.hpp b/src/include/migraphx/base64.hpp index 7903e4453b7..36035430826 100644 --- a/src/include/migraphx/base64.hpp +++ b/src/include/migraphx/base64.hpp @@ -31,7 +31,7 @@ namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { /// encode string to base64 -std::string b64_encode(const std::string& str); +std::string base64_encode(const std::string& str); } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/netron_output.cpp b/src/netron_output.cpp index 85fc1eb9b41..a9ecdde0ecc 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -57,7 +57,7 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; case shape::tuple_type: return 0; - MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); + default: MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); } } @@ -67,7 +67,7 @@ auto make_attribute(const migraphx::value& val) attribute["name"] = val.get_key(); auto val_string = val.to(); val_string = val_string.substr(val_string.find(":") + 1); - attribute["s"] = b64_encode(val_string); + attribute["s"] = base64_encode(val_string); attribute["type"] = "STRING"; return attribute; } diff --git a/test/base64_test.cpp b/test/base64_test.cpp index 9448c54eb75..65b8c3a8800 100644 --- a/test/base64_test.cpp +++ b/test/base64_test.cpp @@ -26,72 +26,42 @@ TEST_CASE(base64_encoding) { - std::string input_0 = "abc"; - std::string expected_0 = "YWJj"; - EXPECT(migraphx::b64_encode(input_0) == expected_0); + EXPECT(migraphx::base64_encode("abc") == "YWJj"); - std::string input_1 = "abcd"; - std::string expected_1 = "YWJjZA=="; - EXPECT(migraphx::b64_encode(input_1) == expected_1); + EXPECT(migraphx::base64_encode("abcd") == "YWJjZA=="); - std::string input_2 = "convolution"; - std::string expected_2 = "Y29udm9sdXRpb24="; - EXPECT(migraphx::b64_encode(input_2) == expected_2); + EXPECT(migraphx::base64_encode("convolution") == "Y29udm9sdXRpb24="); - std::string input_3 = "https://www.amd.com/en/products/software/rocm.html"; - std::string expected_3 = "aHR0cHM6Ly93d3cuYW1kLmNvbS9lbi9wcm9kdWN0cy9zb2Z0d2FyZS9yb2NtLmh0bWw="; - EXPECT(migraphx::b64_encode(input_3) == expected_3); + EXPECT(migraphx::base64_encode("https://www.amd.com/en/products/software/rocm.html") == + "aHR0cHM6Ly93d3cuYW1kLmNvbS9lbi9wcm9kdWN0cy9zb2Z0d2FyZS9yb2NtLmh0bWw="); - std::string input_4 = "{1, 3, 7, 9}"; - std::string expected_4 = "ezEsIDMsIDcsIDl9"; - EXPECT(migraphx::b64_encode(input_4) == expected_4); + EXPECT(migraphx::base64_encode("{1, 3, 7, 9}") == "ezEsIDMsIDcsIDl9"); } TEST_CASE(base64_RFC_test_vectors) { - std::string input_0 = ""; - std::string expected_0 = ""; - EXPECT(migraphx::b64_encode(input_0) == expected_0); + EXPECT(migraphx::base64_encode("") == ""); - std::string input_1 = "f"; - std::string expected_1 = "Zg=="; - EXPECT(migraphx::b64_encode(input_1) == expected_1); + EXPECT(migraphx::base64_encode("f") == "Zg=="); - std::string input_2 = "fo"; - std::string expected_2 = "Zm8="; - EXPECT(migraphx::b64_encode(input_2) == expected_2); + EXPECT(migraphx::base64_encode("fo") == "Zm8="); - std::string input_3 = "foo"; - std::string expected_3 = "Zm9v"; - EXPECT(migraphx::b64_encode(input_3) == expected_3); + EXPECT(migraphx::base64_encode("foo") == "Zm9v"); - std::string input_4 = "foob"; - std::string expected_4 = "Zm9vYg=="; - EXPECT(migraphx::b64_encode(input_4) == expected_4); + EXPECT(migraphx::base64_encode("foob") == "Zm9vYg=="); - std::string input_5 = "fooba"; - std::string expected_5 = "Zm9vYmE="; - EXPECT(migraphx::b64_encode(input_5) == expected_5); + EXPECT(migraphx::base64_encode("fooba") == "Zm9vYmE="); - std::string input_6 = "foobar"; - std::string expected_6 = "Zm9vYmFy"; - EXPECT(migraphx::b64_encode(input_6) == expected_6); + EXPECT(migraphx::base64_encode("foobar") == "Zm9vYmFy"); } // Following tests altered from // https://github.com/tobiaslocker/base64/blob/master/test/base64_tests.cpp -TEST_CASE(base64_encodes_empty) -{ - std::string const expected{}; - std::string const actual{migraphx::b64_encode({})}; - EXPECT(expected == actual); -} - TEST_CASE(base64_encodes_three_bytes_zeros) { std::array const input{0x00, 0x00, 0x00}; std::string expected{"AAAA"}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -99,7 +69,7 @@ TEST_CASE(base64_encodes_three_bytes_random) { std::array const input{0xFE, 0xE9, 0x72}; std::string const expected{"/uly"}; - std::string const actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string const actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -107,7 +77,7 @@ TEST_CASE(base64_encodes_two_bytes) { std::array const input{0x00, 0x00}; std::string expected{"AAA="}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -115,7 +85,7 @@ TEST_CASE(base64_encodes_one_byte) { std::array const input{0x00}; std::string expected{"AA=="}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -123,7 +93,7 @@ TEST_CASE(base64_encodes_four_bytes) { std::array const input{0x74, 0x68, 0x65, 0x20}; std::string expected{"dGhlIA=="}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -131,7 +101,7 @@ TEST_CASE(base64_encodes_five_bytes) { std::array const input{0x20, 0x62, 0x72, 0x6f, 0x77}; std::string expected{"IGJyb3c="}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -139,7 +109,7 @@ TEST_CASE(base64_encodes_six_bytes) { std::array const input{0x20, 0x6a, 0x75, 0x6d, 0x70, 0x73}; std::string expected{"IGp1bXBz"}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -151,7 +121,7 @@ TEST_CASE(base64_encodes_BrownFox) 0x20, 0x74, 0x68, 0x65, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; std::string expected{"dGhlIHF1aWNrIGJyb3duIGZveCBqdW1wcyBvdmVyIHRoZSBsYXp5IGRvZw=="}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } @@ -163,7 +133,7 @@ TEST_CASE(base64_encodes_EncodesBrownFastFoxNullInMiddle) 0x72, 0x20, 0x74, 0x68, 0x65, 0x00, 0x20, 0x6c, 0x61, 0x7a, 0x79, 0x20, 0x64, 0x6f, 0x67}; std::string expected{"dGhlIHF1aWNrISBicm93biBmb3gganVtcHMgb3ZlciB0aGUAIGxhenkgZG9n"}; - std::string actual{migraphx::b64_encode({input.begin(), input.end()})}; + std::string actual{migraphx::base64_encode({input.begin(), input.end()})}; EXPECT(expected == actual); } From 211b2d78ef6dc413b3db47eb427836a9e3d4e606 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 16:41:22 -0600 Subject: [PATCH 18/21] fix back to unsigned char (byte) --- src/base64.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index da0a7200b1d..1d5b888ee71 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -39,10 +39,10 @@ std::array constexpr b64_chars{ 'w', 'x', 'y', 'z', '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', '+', '/'}; /// base64 encoder snippet altered from https://stackoverflow.com/a/37109258 -std::string encode(const std::string& buf) +std::string encode(const std::vector& buf) { std::size_t len = buf.size(); - std::string res_vec((len + 2) / 3 * 4, '='); + std::vector res_vec((len + 2) / 3 * 4, '='); std::size_t j = 0; std::size_t remaining = len % 3; const size_t last = len - remaining; @@ -72,7 +72,10 @@ std::string encode(const std::string& buf) } // namespace -std::string base64_encode(const std::string& str) { return encode(str); } +std::string base64_encode(const std::string& str) +{ + return encode(std::vector(str.begin(), str.end())); +} } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx From 18019ed7f345da40318fa3006e8220cc1c7aa7bd Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 16:52:09 -0600 Subject: [PATCH 19/21] move throw --- src/netron_output.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/netron_output.cpp b/src/netron_output.cpp index a9ecdde0ecc..bba99370a95 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -57,8 +57,8 @@ int get_onnx_type(shape::type_t s_type) case shape::fp8e5m2_type: return 19; case shape::fp8e5m2fnuz_type: return 20; case shape::tuple_type: return 0; - default: MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); } + MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); } auto make_attribute(const migraphx::value& val) From a062ef8ef2369c2e46465fa30792a5d67864b9c5 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 17:19:13 -0600 Subject: [PATCH 20/21] formatting --- src/base64.cpp | 2 +- src/netron_output.cpp | 55 +++++++++++++++++++++++-------------------- 2 files changed, 31 insertions(+), 26 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index 1d5b888ee71..15e6542529a 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -43,7 +43,7 @@ std::string encode(const std::vector& buf) { std::size_t len = buf.size(); std::vector res_vec((len + 2) / 3 * 4, '='); - std::size_t j = 0; + std::size_t j = 0; std::size_t remaining = len % 3; const size_t last = len - remaining; diff --git a/src/netron_output.cpp b/src/netron_output.cpp index bba99370a95..64403dd7090 100644 --- a/src/netron_output.cpp +++ b/src/netron_output.cpp @@ -39,24 +39,24 @@ int get_onnx_type(shape::type_t s_type) { switch(s_type) { - case shape::float_type: return 1; - case shape::uint8_type: return 2; - case shape::int8_type: return 3; - case shape::uint16_type: return 4; - case shape::int16_type: return 5; - case shape::int32_type: return 6; - case shape::int64_type: return 7; - case shape::bool_type: return 9; - case shape::half_type: return 10; - case shape::double_type: return 11; - case shape::uint32_type: return 12; - case shape::uint64_type: return 13; - case shape::bf16_type: return 16; - case shape::fp8e4m3fn_type: return 17; - case shape::fp8e4m3fnuz_type: return 18; - case shape::fp8e5m2_type: return 19; - case shape::fp8e5m2fnuz_type: return 20; - case shape::tuple_type: return 0; + case shape::float_type: return 1; + case shape::uint8_type: return 2; + case shape::int8_type: return 3; + case shape::uint16_type: return 4; + case shape::int16_type: return 5; + case shape::int32_type: return 6; + case shape::int64_type: return 7; + case shape::bool_type: return 9; + case shape::half_type: return 10; + case shape::double_type: return 11; + case shape::uint32_type: return 12; + case shape::uint64_type: return 13; + case shape::bf16_type: return 16; + case shape::fp8e4m3fn_type: return 17; + case shape::fp8e4m3fnuz_type: return 18; + case shape::fp8e5m2_type: return 19; + case shape::fp8e5m2fnuz_type: return 20; + case shape::tuple_type: return 0; } MIGRAPHX_THROW("MIGraphX type " + std::to_string(s_type) + " not supported"); } @@ -73,7 +73,8 @@ auto make_attribute(const migraphx::value& val) } /// Returns a value with the JSON structure needed for a node -auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) +auto make_onnx_json_node(instruction_ref ins, + std::unordered_map ins_uids) { value node; // TODO add support for module inputs @@ -133,12 +134,13 @@ auto make_onnx_json_node(instruction_ref ins, std::unordered_map ins_uids) +auto make_onnx_json_literal(instruction_ref ins, + std::unordered_map ins_uids) { value lit; lit["dims"] = ins->get_shape().lens(); lit["dataType"] = get_onnx_type(ins->get_shape().type()); - lit["name"] = ins_uids.at(ins); + lit["name"] = ins_uids.at(ins); // ignoring literal data, setting to "NULL" in base64 lit["rawData"] = "TlVMTA=="; return lit; @@ -160,11 +162,13 @@ auto make_onnx_json_shape(const shape& s) } // ONNX graph edges called "valuetype" -auto make_onnx_json_edge(instruction_ref ins, instruction_ref out_ins, std::unordered_map ins_uids) +auto make_onnx_json_edge(instruction_ref ins, + instruction_ref out_ins, + std::unordered_map ins_uids) { value ret; shape ins_shape = ins->get_shape(); - ret["name"] = ins_uids.at(ins) + "->" + ins_uids.at(out_ins); + ret["name"] = ins_uids.at(ins) + "->" + ins_uids.at(out_ins); value type = {{"tensorType", {{"elemType", get_onnx_type(ins_shape.type())}, {"shape", make_onnx_json_shape(ins_shape)}}}}; @@ -172,11 +176,12 @@ auto make_onnx_json_edge(instruction_ref ins, instruction_ref out_ins, std::unor return ret; } -auto make_onnx_json_in_out(instruction_ref ins, std::unordered_map ins_uids) +auto make_onnx_json_in_out(instruction_ref ins, + std::unordered_map ins_uids) { value ret; shape ins_shape = ins->get_shape(); - ret["name"] = ins_uids.at(ins); + ret["name"] = ins_uids.at(ins); value type = {{"tensorType", {{"elemType", get_onnx_type(ins_shape.type())}, {"shape", make_onnx_json_shape(ins_shape)}}}}; From 5e92ff92e6beebfb53c66892875b9d6bfefa3e60 Mon Sep 17 00:00:00 2001 From: charlie Date: Thu, 19 Dec 2024 17:22:33 -0600 Subject: [PATCH 21/21] pacify cppcheck --- src/base64.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/src/base64.cpp b/src/base64.cpp index 15e6542529a..0d08b1f6220 100644 --- a/src/base64.cpp +++ b/src/base64.cpp @@ -60,12 +60,12 @@ std::string encode(const std::vector& buf) // Set padding if(remaining != 0) { - std::size_t n = --remaining != 0 ? static_cast(buf.at(last)) << 8u | - static_cast(buf.at(last + 1)) - : static_cast(buf.at(last)); - res_vec.at(j++) = b64_chars.at(remaining != 0 ? n >> 10u & 0x3Fu : n >> 2u); - res_vec.at(j++) = b64_chars.at(remaining != 0 ? n >> 4u & 0x03Fu : n << 4u & 0x3Fu); - res_vec.at(j++) = remaining != 0 ? b64_chars.at(n << 2u & 0x3Fu) : '='; + std::size_t n = --remaining == 0 ? static_cast(buf.at(last)) + : static_cast(buf.at(last)) << 8u | + static_cast(buf.at(last + 1)); + res_vec.at(j++) = b64_chars.at(remaining == 0 ? n >> 2u : n >> 10u & 0x3Fu); + res_vec.at(j++) = b64_chars.at(remaining == 0 ? n << 4u & 0x3Fu : n >> 4u & 0x03Fu); + res_vec.at(j++) = remaining == 0 ? '=' : b64_chars.at(n << 2u & 0x3Fu); } return {res_vec.begin(), res_vec.end()}; }