From e8e9dfd9fb8d0288fa6a289728e065c515320567 Mon Sep 17 00:00:00 2001 From: Dmitrii Zarukin Date: Tue, 7 Jan 2025 15:56:51 -0800 Subject: [PATCH 01/16] benchdnn: graph: styling --- tests/benchdnn/graph/custom_driver.cpp | 15 ++++++++------- tests/benchdnn/graph/custom_driver.hpp | 7 +++---- tests/benchdnn/graph/utils.cpp | 7 ++++--- 3 files changed, 15 insertions(+), 14 deletions(-) diff --git a/tests/benchdnn/graph/custom_driver.cpp b/tests/benchdnn/graph/custom_driver.cpp index cf0111266ce..d82b7c44864 100644 --- a/tests/benchdnn/graph/custom_driver.cpp +++ b/tests/benchdnn/graph/custom_driver.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2023-2024 Intel Corporation +* Copyright 2023-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -257,15 +257,16 @@ void init_memory_args(dnn_mem_map_t &mem_map, const prb_t *prb, if (prb->arg_mds_.find(exec_arg) == prb->arg_mds_.end()) { assert(!"missing required args"); SAFE_V(FAIL); - }; + } auto arg_mds_ = prb->arg_mds_.find(exec_arg)->second; - dnnl_dims_t dnnl_dims; - auto dim = ::std::get<1>(arg_mds_); - for (size_t i = 0; i < dim.size(); i++) { - dnnl_dims[i] = dim[i]; + dnnl_dims_t dnnl_dims {}; + auto dims = ::std::get<1>(arg_mds_); + for (size_t i = 0; i < dims.size(); i++) { + dnnl_dims[i] = dims[i]; } + mem_map.emplace(exec_arg, - dnn_mem_t(static_cast(dim.size()), dnnl_dims, + dnn_mem_t(static_cast(dims.size()), dnnl_dims, std::get<2>(arg_mds_), ::std::get<0>(arg_mds_), test_engine)); } diff --git a/tests/benchdnn/graph/custom_driver.hpp b/tests/benchdnn/graph/custom_driver.hpp index a843501e60e..3ccf903e190 100644 --- a/tests/benchdnn/graph/custom_driver.hpp +++ b/tests/benchdnn/graph/custom_driver.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2023-2024 Intel Corporation +* Copyright 2023-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -49,14 +49,13 @@ struct settings_t { }; struct prb_t { - prb_t(const settings_t &s) { - arg_mds_ = s.arg_mds_; - alg = s.alg; + prb_t(const settings_t &s) : arg_mds_(s.arg_mds_), alg(s.alg) { switch (alg) { case TRANSPOSE: order = s.order; break; default: break; } } + ::std::unordered_map arg_mds_; ::std::vector order; attr_t attr; diff --git a/tests/benchdnn/graph/utils.cpp b/tests/benchdnn/graph/utils.cpp index 66477ca1758..75a99010083 100644 --- a/tests/benchdnn/graph/utils.cpp +++ b/tests/benchdnn/graph/utils.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022-2024 Intel Corporation +* Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -1255,8 +1255,9 @@ dnnl_data_type_t convert_dt(const dnnl::graph::logical_tensor::data_type dt) { case graph_dt::s32: return dnnl_s32; case graph_dt::s8: return dnnl_s8; case graph_dt::u8: return dnnl_u8; - // use u8 instead of boolean in the reference path - // dnn_graph_mem_t will use the data type from the logical tensor and the u8 data handle + // Use `u8` instead of `boolean` in the reference path. + // `dnn_graph_mem_t` will use the data type from the logical tensor and + // the `u8` data handle. case graph_dt::boolean: return dnnl_u8; case graph_dt::f8_e5m2: return dnnl_f8_e5m2; case graph_dt::f8_e4m3: return dnnl_f8_e4m3; From 07589b9571864519f789e9197d45c92d4d8ad82c Mon Sep 17 00:00:00 2001 From: Dmitrii Zarukin Date: Tue, 7 Jan 2025 17:42:37 -0800 Subject: [PATCH 02/16] benchdnn: graph: move custom op on f32 completely --- tests/benchdnn/graph/graph_memory.cpp | 25 +++++++++++------------- tests/benchdnn/graph/setting_handler.cpp | 15 +++----------- 2 files changed, 14 insertions(+), 26 deletions(-) diff --git a/tests/benchdnn/graph/graph_memory.cpp b/tests/benchdnn/graph/graph_memory.cpp index 87264af619c..b7752c415a4 100644 --- a/tests/benchdnn/graph/graph_memory.cpp +++ b/tests/benchdnn/graph/graph_memory.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2023-2024 Intel Corporation +* Copyright 2023-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -42,18 +42,13 @@ size_t get_benchdnn_device_limit() { return benchdnn_device_limit; } +// Constructs memories for all inputs and outputs needed for comparison. dnn_graph_mem_t::dnn_graph_mem_t(const dnn_mem_t &mem, const deserialized_lt <, const bool is_op_input, const bool is_fake_output) { - - // Init memory for all inputs and outputs that needs comparison const auto &prim_dt = mem.dt(); - const auto &graph_dt = static_cast(lt.get_data_type()); - const bool is_boolean - = lt.get_data_type() == logical_tensor::data_type::boolean; - - // Use data type from graph path to represent boolean - const auto &c_data_type = is_boolean ? prim_dt : graph_dt; + // Conversion from graph types to dnnl types + boolean to u8. + const auto &graph_dt = convert_dt(lt.get_data_type()); // Get memory tag of primitive memory int ndims = mem.ndims(); @@ -77,7 +72,7 @@ dnn_graph_mem_t::dnn_graph_mem_t(const dnn_mem_t &mem, // otherwise use shape & tag from ref path side // Create memory for graph path - const auto data_type = static_cast(c_data_type); + const auto data_type = static_cast(graph_dt); if (is_op_input) { if (graph_dims_.empty()) graph_dims_.push_back(1); if (graph_strides_.empty()) graph_strides_.push_back(1); @@ -93,12 +88,14 @@ dnn_graph_mem_t::dnn_graph_mem_t(const dnn_mem_t &mem, std::memcpy(graph_data_handle, prim_data_handle, graph_mem.size()); }; - // Not do reorder for boolean data tensor - if (!is_boolean && prim_dt != c_data_type) { - dnn_mem_t c_mem(ndims, mem.dims(), c_data_type, mtag, g_eng.get()); + if (prim_dt != graph_dt) { + // Call a reorder (for data conversion) when reference memory + // doesn't coincide with the graph memory... + dnn_mem_t c_mem(ndims, mem.dims(), graph_dt, mtag, g_eng.get()); SAFE_V(c_mem.reorder(mem)); prim_to_graph_memcpy(mem_, c_mem); } else { + // ... otherwise, perform a plain memcpy. prim_to_graph_memcpy(mem_, mem); } } else { @@ -106,7 +103,7 @@ dnn_graph_mem_t::dnn_graph_mem_t(const dnn_mem_t &mem, dnnl::memory::desc md(graph_dims_, data_type, graph_strides_); mem_ = dnn_mem_t(md.get(), g_eng.get()); } else { - mem_ = dnn_mem_t(mem.md_, c_data_type, mtag, g_eng.get()); + mem_ = dnn_mem_t(mem.md_, graph_dt, mtag, g_eng.get()); } } } diff --git a/tests/benchdnn/graph/setting_handler.cpp b/tests/benchdnn/graph/setting_handler.cpp index 6a8471e3999..697b904e8a4 100644 --- a/tests/benchdnn/graph/setting_handler.cpp +++ b/tests/benchdnn/graph/setting_handler.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022-2024 Intel Corporation +* Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -169,20 +169,13 @@ ::custom::settings_t get_setting( res->state = res_state_t::INVALID_ARGUMENTS; return op_setting; } - // Select op has boolean weights. It requires special handling for dt - // conversion because custom driver prb values directly translate into graph - // objects, there's no intermediate primitive layer that can be instructed - // to have f32 data type. - const bool op_is_select = opkind == ::graph::op::kind::Select; for (size_t i = 0; i < base_op_ref.in_lts_.size(); i++) { const auto arg = get_prim_arg_name_from_graph_op_input_offset( opkind, static_cast(i)); const auto < = base_op_ref.in_lts_[i]; auto dim = lt.shape_; - const auto orig_dt = convert_dt(lt.get_data_type()); - const auto dt - = op_is_select && arg == DNNL_ARG_WEIGHTS ? orig_dt : dnnl_f32; + const auto dt = dnnl_f32; auto tag = strides2memory_tag(lt.stride_.size(), lt.stride_, false); // 0-dim means scalar input in graph, extend to 1-dim to match behavior. @@ -197,9 +190,7 @@ ::custom::settings_t get_setting( opkind, static_cast(i)); const auto < = base_op_ref.out_lts_[i]; auto dim = lt.shape_; - const auto orig_dt = convert_dt(lt.get_data_type()); - const auto dt - = op_is_select && arg == DNNL_ARG_WEIGHTS ? orig_dt : dnnl_f32; + const auto dt = dnnl_f32; auto tag = strides2memory_tag(lt.stride_.size(), lt.stride_, false); // 0-dim means scalar input in graph, extend to 1-dim to match behavior. From 173c98326764f7f7a910db486c0ccb3a32605adc Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Sun, 29 Dec 2024 22:19:23 -0800 Subject: [PATCH 03/16] graph: api: support GenIndex and GreaterEqual operations --- include/oneapi/dnnl/dnnl_graph.hpp | 4 +++- include/oneapi/dnnl/dnnl_graph_types.h | 4 +++- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/include/oneapi/dnnl/dnnl_graph.hpp b/include/oneapi/dnnl/dnnl_graph.hpp index a9732cfee2f..8a7aa34eafb 100644 --- a/include/oneapi/dnnl/dnnl_graph.hpp +++ b/include/oneapi/dnnl/dnnl_graph.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -846,6 +846,8 @@ class op : public op_handle { TanhBackward = dnnl_graph_op_tanh_backward, TypeCast = dnnl_graph_op_type_cast, Wildcard = dnnl_graph_op_wildcard, + GenIndex = dnnl_graph_op_gen_index, + GreaterEqual = dnnl_graph_op_greater_equal, // Sentinel LastSymbol = dnnl_graph_op_last_symbol, }; diff --git a/include/oneapi/dnnl/dnnl_graph_types.h b/include/oneapi/dnnl/dnnl_graph_types.h index 1cb049653a6..4aeb4d6bd87 100644 --- a/include/oneapi/dnnl/dnnl_graph_types.h +++ b/include/oneapi/dnnl/dnnl_graph_types.h @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2020-2024 Intel Corporation + * Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -256,6 +256,8 @@ typedef enum { dnnl_graph_op_select, dnnl_graph_op_pow, dnnl_graph_op_group_norm, + dnnl_graph_op_gen_index, + dnnl_graph_op_greater_equal, dnnl_graph_op_last_symbol, } dnnl_graph_op_kind_t; From 2c64f5531f79bdc5c6e4f98497059571e56399a1 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Sun, 29 Dec 2024 22:22:29 -0800 Subject: [PATCH 04/16] graph: interface: support GenIndex and GreaterEqual operations --- src/graph/interface/c_types_map.hpp | 4 +++- src/graph/interface/op.hpp | 4 +++- src/graph/interface/op_def.hpp | 30 ++++++++++++++++++++++++++++- src/graph/interface/opset.hpp | 4 +++- 4 files changed, 38 insertions(+), 4 deletions(-) diff --git a/src/graph/interface/c_types_map.hpp b/src/graph/interface/c_types_map.hpp index 37c1a91685c..c07ae3ff61a 100644 --- a/src/graph/interface/c_types_map.hpp +++ b/src/graph/interface/c_types_map.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -156,6 +156,8 @@ const op_kind_t End = dnnl_graph_op_end; const op_kind_t Exp = dnnl_graph_op_exp; const op_kind_t GELU = dnnl_graph_op_gelu; const op_kind_t GELUBackward = dnnl_graph_op_gelu_backward; +const op_kind_t GenIndex = dnnl_graph_op_gen_index; +const op_kind_t GreaterEqual = dnnl_graph_op_greater_equal; const op_kind_t GroupNorm = dnnl_graph_op_group_norm; const op_kind_t HardSigmoid = dnnl_graph_op_hard_sigmoid; const op_kind_t HardSigmoidBackward = dnnl_graph_op_hard_sigmoid_backward; diff --git a/src/graph/interface/op.hpp b/src/graph/interface/op.hpp index 406c0449fc9..ec5312b6fc2 100644 --- a/src/graph/interface/op.hpp +++ b/src/graph/interface/op.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -394,6 +394,8 @@ struct dnnl_graph_op : public std::enable_shared_from_this { CASE(Exp); CASE(GELU); CASE(GELUBackward); + CASE(GenIndex); + CASE(GreaterEqual); CASE(GroupNorm); CASE(HardSigmoid); CASE(HardSigmoidBackward); diff --git a/src/graph/interface/op_def.hpp b/src/graph/interface/op_def.hpp index 45947c67933..7d3c03aabae 100644 --- a/src/graph/interface/op_def.hpp +++ b/src/graph/interface/op_def.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -436,6 +436,34 @@ DNNL_GRAPH_OP_SCHEMA(GELUBackward, 1, "T", {data_type::f32, data_type::bf16, data_type::f16}) .set_shape_inference_function(infer_identity_output_shape)) +DNNL_GRAPH_OP_SCHEMA(GenIndex, 1, + op_schema_t() + .set_num_inputs(1) + .set_num_outputs(1) + .set_input(0, "src", "T1") + .set_output(0, "dst", "T2") + .set_attr(op_attr::axis, true, attribute_kind::i) + .set_type_constraints( + "T1", {data_type::f32, data_type::bf16, data_type::f16}) + .set_type_constraints("T2", {data_type::s32}) + .set_shape_inference_function(infer_identity_output_shape)) + +DNNL_GRAPH_OP_SCHEMA(GreaterEqual, 1, + op_schema_t() + .set_num_inputs(2) + .set_num_outputs(1) + .set_input(0, "src_0", "T1") + .set_input(1, "src_1", "T1") + .set_output(0, "dst", "T2") + .set_attr(op_attr::auto_broadcast, false, attribute_kind::s, + "numpy", {"none", "numpy"}) + .set_type_constraints("T1", + {data_type::f32, data_type::bf16, data_type::f16, + data_type::s32}) + .set_type_constraints("T2", {data_type::boolean}) + .set_shape_inference_function( + infer_elemwise_arithmetic_output_shape)) + DNNL_GRAPH_OP_SCHEMA(GroupNorm, 1, op_schema_t() .set_inputs_option(op_schema_t::param_num_option::optional) diff --git a/src/graph/interface/opset.hpp b/src/graph/interface/opset.hpp index a0d2eb08a65..99588c4c2ad 100644 --- a/src/graph/interface/opset.hpp +++ b/src/graph/interface/opset.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -65,6 +65,8 @@ class opset_v1_t { fn(get_op_schema()); fn(get_op_schema()); fn(get_op_schema()); + fn(get_op_schema()); + fn(get_op_schema()); fn(get_op_schema()); fn(get_op_schema()); fn(get_op_schema Date: Wed, 25 Dec 2024 05:42:41 +0000 Subject: [PATCH 05/16] graph: backend: dnnl: add internal op of dnnl_gen_index --- src/graph/backend/dnnl/dnnl_op_def.hpp | 17 +++++- src/graph/backend/dnnl/dnnl_opset.hpp | 3 +- src/graph/backend/dnnl/internal_ops.hpp | 5 +- src/graph/backend/dnnl/layout_propagator.cpp | 16 +++++- src/graph/backend/dnnl/layout_propagator.hpp | 3 +- src/graph/backend/dnnl/op_executable.cpp | 31 ++++++++++- src/graph/backend/dnnl/op_executable.hpp | 55 +++++++++++++++++++- src/graph/backend/dnnl/passes/lower.cpp | 11 +++- src/graph/backend/dnnl/utils.hpp | 11 +++- 9 files changed, 142 insertions(+), 10 deletions(-) diff --git a/src/graph/backend/dnnl/dnnl_op_def.hpp b/src/graph/backend/dnnl/dnnl_op_def.hpp index afdd51a37a9..148efc50817 100644 --- a/src/graph/backend/dnnl/dnnl_op_def.hpp +++ b/src/graph/backend/dnnl/dnnl_op_def.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2021-2024 Intel Corporation +* Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -774,6 +774,21 @@ DNNL_GRAPH_OP_SCHEMA(dnnl_eltwise_bwd, 1, executable_creator) .SET_ARG_INDICES_GETTER(eltwise_bwd_executable_t)) +DNNL_GRAPH_OP_SCHEMA(dnnl_gen_index, 1, + op_schema_t() + .set_num_inputs(1) + .set_num_outputs(1) + .set_input(0, "input") + .set_output(0, "output") + // Attributes inherited from front GenIndex ops + .set_attr(op_attr::axis, true, attribute_kind::i) + // Analysis rules + .set_shape_inference_function(infer_identity_output_shape) + .SET_LAYOUT_PROPAGATOR(layout_propagator_for_gen_index) + .SET_EXECUTABLE_CREATOR( + executable_creator) + .SET_ARG_INDICES_GETTER(genindex_executable_t)) + DNNL_GRAPH_OP_SCHEMA(dnnl_shuffle, 1, op_schema_t() .set_num_inputs(1) diff --git a/src/graph/backend/dnnl/dnnl_opset.hpp b/src/graph/backend/dnnl/dnnl_opset.hpp index 3a8beb8e2f4..9f1595a0920 100644 --- a/src/graph/backend/dnnl/dnnl_opset.hpp +++ b/src/graph/backend/dnnl/dnnl_opset.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2021-2024 Intel Corporation +* Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -72,6 +72,7 @@ class dnnl_opset_t { fn(get_op_schema()); fn(get_op_schema()); + fn(get_op_schema()); fn(get_op_schema()); fn(get_op_schema()); fn(get_op_schema()); diff --git a/src/graph/backend/dnnl/internal_ops.hpp b/src/graph/backend/dnnl/internal_ops.hpp index 95e6c86e03a..c8d73e9f56c 100644 --- a/src/graph/backend/dnnl/internal_ops.hpp +++ b/src/graph/backend/dnnl/internal_ops.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2021-2024 Intel Corporation +* Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -77,7 +77,8 @@ namespace op_kind { X(dnnl_reorder, Dnnl_reorder) \ X(dnnl_convtranspose_bwd_data, Dnnl_convtranspose_bwd_data) \ X(dnnl_convtranspose_bwd_weights, Dnnl_convtranspose_bwd_weights) \ - X(dnnl_groupnorm, Dnnl_groupnorm) + X(dnnl_groupnorm, Dnnl_groupnorm) \ + X(dnnl_gen_index, Dnnl_gen_index) enum kind_t { kDNNL_INTERNAL_OP_STARTER = 0x1234, diff --git a/src/graph/backend/dnnl/layout_propagator.cpp b/src/graph/backend/dnnl/layout_propagator.cpp index ab4902d7000..dd1aa5de5ce 100644 --- a/src/graph/backend/dnnl/layout_propagator.cpp +++ b/src/graph/backend/dnnl/layout_propagator.cpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2022-2024 Intel Corporation + * Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -1510,6 +1510,20 @@ status_t layout_propagator_for_add_zps(std::shared_ptr &op, return status::invalid_graph_op; } +status_t layout_propagator_for_gen_index(std::shared_ptr &op, + const dnnl::engine &p_engine, fusion_info_mgr_t &mgr, + pd_cache_t &pd_cache, subgraph_rewriter_t &rewriter) { + UNUSED(p_engine); + UNUSED(mgr); + UNUSED(pd_cache); + UNUSED(rewriter); + auto src_md = make_dnnl_memory_desc( + op->get_input_value(0)->get_logical_tensor()); + value_ptr dst_val = op->get_output_value(0); + status_t status = fill_layout_info(dst_val, src_md); + return status; +} + status_t layout_propagator_for_groupnorm(op_ptr &op, const dnnl::engine &p_engine, fusion_info_mgr_t &mgr, pd_cache_t &pd_cache, subgraph_rewriter_t &rewriter) { diff --git a/src/graph/backend/dnnl/layout_propagator.hpp b/src/graph/backend/dnnl/layout_propagator.hpp index 5d7e1f62948..9b231ad6da3 100644 --- a/src/graph/backend/dnnl/layout_propagator.hpp +++ b/src/graph/backend/dnnl/layout_propagator.hpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2022-2024 Intel Corporation + * Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -91,6 +91,7 @@ DECLARE_LAYOUT_PROPAGATOR(constant_filler); DECLARE_LAYOUT_PROPAGATOR(sub_zps); DECLARE_LAYOUT_PROPAGATOR(add_zps); DECLARE_LAYOUT_PROPAGATOR(groupnorm); +DECLARE_LAYOUT_PROPAGATOR(gen_index); #undef DECLARE_LAYOUT_PROPAGATOR diff --git a/src/graph/backend/dnnl/op_executable.cpp b/src/graph/backend/dnnl/op_executable.cpp index 94e9d43b3aa..55b0c7e9f96 100644 --- a/src/graph/backend/dnnl/op_executable.cpp +++ b/src/graph/backend/dnnl/op_executable.cpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2022-2024 Intel Corporation + * Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,11 +25,14 @@ #include +#include "common/dnnl_thread.hpp" + #include "graph/backend/dnnl/common.hpp" #include "graph/backend/dnnl/dnnl_constant_tensor_cache.hpp" #include "graph/backend/dnnl/fusion_info.hpp" #include "graph/backend/dnnl/internal_attrs.hpp" #include "graph/backend/dnnl/op_executable.hpp" +#include "graph/backend/dnnl/utils.hpp" namespace dnnl { namespace impl { @@ -1770,6 +1773,20 @@ groupnorm_executable_t::desc_t groupnorm_executable_t::create_desc( return {pd, false}; } +void genindex_executable_t ::execute(const stream &stream, + const std::unordered_map &args) const { + auto &output = args.find(DNNL_ARG_DST)->second; + auto output_ptr = static_cast(output.get_data_handle()); + dnnl::impl::parallel_nd(nelems_, [&](dim_t i) { + dims_t input_dims; // decomposition for physical offsets + dnnl::impl::utils::l_dims_by_l_offset( + input_dims, i, output_dims_, ndims_); + auto offset + = utils::offset_compute(output_strides_, input_dims, ndims_); + output_ptr[offset] = input_dims[axis_]; + }); +} + static void get_arg_indices_for_post_ops(const op_t *op, fusion_info_mgr_t &mgr, arg_indices_t &indices, size_t &base_index) { const fusion_info_t &fusion_info @@ -2363,6 +2380,18 @@ arg_indices_t groupnorm_executable_t::get_arg_indices( return get_arg_indices_for_lnorm_and_gnorm(op, mgr); } +arg_indices_t genindex_executable_t::get_arg_indices( + const op_t *op, fusion_info_mgr_t &mgr) { + UNUSED(op); + UNUSED(mgr); + + arg_indices_t arg_indices; + arg_indices.insert({DNNL_ARG_SRC, indices_t {input, 0}}); + arg_indices.insert({DNNL_ARG_DST, indices_t {output, 0}}); + + return arg_indices; +} + } // namespace dnnl_impl } // namespace graph } // namespace impl diff --git a/src/graph/backend/dnnl/op_executable.hpp b/src/graph/backend/dnnl/op_executable.hpp index 6fa36e22cc1..ef199cf9c69 100644 --- a/src/graph/backend/dnnl/op_executable.hpp +++ b/src/graph/backend/dnnl/op_executable.hpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2021-2024 Intel Corporation + * Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -2466,6 +2466,59 @@ struct groupnorm_executable_t : public op_executable_t { dnnl::group_normalization_forward prim_; }; +struct genindex_executable_t : public op_executable_t { + DECLARE_ARG_INDICES_GETTER; + + genindex_executable_t(std::shared_ptr &op, + const dnnl::engine &p_engine, fusion_info_mgr_t &mgr, + pd_cache_t &pd_cache) { + if (p_engine.get_kind() == engine::kind::gpu) { + assertm(false, + "genindex opexcutable is unimplemented " + "under SYCL and OCL " + "runtime!"); + throw std::runtime_error("Unimplement"); + } + using ltw = logical_tensor_wrapper_t; + const auto &input_lt = op->get_input_value(0)->get_logical_tensor(); + nelems_ = ltw(input_lt).nelems(); + ndims_ = ltw(input_lt).ndims(); + axis_ = op->get_attr(dnnl::impl::graph::op_attr::axis); + const auto &output_lt = op->get_output_value(0)->get_logical_tensor(); + for (int i = 0; i < ndims_; i++) { + output_dims_[i] = output_lt.dims[i]; + output_strides_[i] = output_lt.layout.strides[i]; + } + } + + void execute(const stream &stream, + const std::unordered_map &args) const override; + +#ifdef DNNL_WITH_SYCL + ::sycl::event execute_sycl(const stream &stream, + const std::unordered_map &args, + const std::vector<::sycl::event> &deps = {}) const override { + execute(stream, args); + return {}; + } +#endif + +#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_OCL + cl_event execute_ocl(const stream &stream, + const std::unordered_map &args, + const std::vector &deps = {}) const override { + assertm(false, + "genindex op excutable is unimplemented " + "under OCL runtime!"); + return {}; + } +#endif + +private: + int axis_, nelems_, ndims_; + dims_t output_dims_, output_strides_; +}; + } // namespace dnnl_impl } // namespace graph } // namespace impl diff --git a/src/graph/backend/dnnl/passes/lower.cpp b/src/graph/backend/dnnl/passes/lower.cpp index dca5d32909e..d1143cbbc0b 100644 --- a/src/graph/backend/dnnl/passes/lower.cpp +++ b/src/graph/backend/dnnl/passes/lower.cpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2022-2024 Intel Corporation + * Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -771,6 +771,14 @@ static status_t select_handler( return status::success; } +static status_t gen_index_handler( + const std::shared_ptr &op, subgraph_rewriter_t &rewriter) { + auto new_op = std::make_shared(op_kind::dnnl_gen_index); + new_op->merge_attributes(op->get_attributes()); + rewriter.replace_op(op, new_op); + return status::success; +} + #define ITEM(kind, func) \ { \ graph::op_kind::kind, handler_func { (func) } \ @@ -881,6 +889,7 @@ static const std::unordered_map handler_table { ITEM(Concat, common_handler), ITEM(SquaredDifference, squared_difference_handler), ITEM(Select, select_handler), + ITEM(GenIndex, gen_index_handler), // utility ITEM(Wildcard, dummy_handler), ITEM(End, dummy_handler), diff --git a/src/graph/backend/dnnl/utils.hpp b/src/graph/backend/dnnl/utils.hpp index b4b7de7d01b..f36c60d8465 100644 --- a/src/graph/backend/dnnl/utils.hpp +++ b/src/graph/backend/dnnl/utils.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2022 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -123,6 +123,15 @@ inline bool all_zero(const std::vector &vec) { return no_zero_pos == vec.end(); } +inline dim_t offset_compute( + const dims_t &strides, const dims_t &idx, int ndims) { + dim_t off = 0; + for (int i = 0; i < ndims; i++) { + off += idx[i] * strides[i]; + } + return off; +} + } // namespace utils } // namespace dnnl_impl } // namespace graph From 2290e02045e470939277bdedb942972f60f3a6ef Mon Sep 17 00:00:00 2001 From: "Gu, Yonghao" Date: Tue, 24 Dec 2024 02:53:30 +0000 Subject: [PATCH 06/16] graph: dnnl: enable gen_index in dnnl backend --- src/graph/backend/dnnl/kernels/gen_index.cpp | 154 ++++++++++++++++++ src/graph/backend/dnnl/kernels/gen_index.hpp | 91 +++++++++++ src/graph/backend/dnnl/kernels/kernels.hpp | 3 +- .../dnnl/patterns/single_op_pattern.cpp | 14 +- 4 files changed, 260 insertions(+), 2 deletions(-) create mode 100644 src/graph/backend/dnnl/kernels/gen_index.cpp create mode 100644 src/graph/backend/dnnl/kernels/gen_index.hpp diff --git a/src/graph/backend/dnnl/kernels/gen_index.cpp b/src/graph/backend/dnnl/kernels/gen_index.cpp new file mode 100644 index 00000000000..6ca9d1cf67e --- /dev/null +++ b/src/graph/backend/dnnl/kernels/gen_index.cpp @@ -0,0 +1,154 @@ +/******************************************************************************* +* Copyright 2025 Intel Corporation +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ +#include "graph/backend/dnnl/kernels/gen_index.hpp" + +#include "graph/backend/dnnl/passes/compile_ops.hpp" +#include "graph/backend/dnnl/passes/constant_propagation.hpp" +#include "graph/backend/dnnl/passes/insert_ops.hpp" +#include "graph/backend/dnnl/passes/layout_propagation.hpp" +#include "graph/backend/dnnl/passes/lower.hpp" +#include "graph/backend/dnnl/passes/memory_planning.hpp" +#include "graph/backend/dnnl/passes/transform.hpp" +#include "graph/backend/dnnl/passes/utils.hpp" + +#include "graph/backend/dnnl/op_executable.hpp" +namespace dnnl { +namespace impl { +namespace graph { +namespace dnnl_impl { + +status_t genindex_t::compile_impl(const dnnl_partition_impl_t *part, + const engine_t *g_engine, const std::vector &inputs, + const std::vector &outputs) { + p_engine_ = make_dnnl_engine(*g_engine); + g_alloc_ + = reinterpret_cast(g_engine->get_allocator()); + + subgraph_ = std::make_shared(part->get_ops(), p_engine_, + part->get_fpmath_mode(), part->get_use_blocked_layout(), true); + BACKEND_DNNL_CHECK(set_given_inputs_outputs(subgraph_, inputs, outputs)); + + subgraph_visualizer_t vis(part->id(), [this](const value_t *val) { + return this->memory_planner_.get_memory_info(val); + }); + pass_pipeline_t pipeline(vis); + + BACKEND_DNNL_ADD_PASS(pipeline, lower_down); + + pipeline.reset_visualize_arg(true, false); + + BACKEND_DNNL_ADD_PASS(pipeline, layout_propagation); + + // constant propagation + if (enabled_constant_cache()) { + BACKEND_DNNL_ADD_PASS(pipeline, constant_propagation); + } + + // bind the memory for each op + auto memory_plan = [&](std::shared_ptr &sg) { + return memory_planner_.run(sg); + }; + pipeline.reset_visualize_arg(true, true); + BACKEND_DNNL_ADD_PASS(pipeline, memory_plan); + BACKEND_DNNL_ADD_PASS(pipeline, compile_ops); + + // Run the added passes + BACKEND_DNNL_CHECK(pipeline.run(subgraph_)); + + // fill information for outputs logical tensors + for (size_t i = 0; i < outputs.size(); i++) { + auto &out = const_cast(outputs[i]); + out = subgraph_->outs_[i]; + } + + // generate a hash key for exec_args_mgr + resource_ctor_ = [this]() { + return this->memory_planner_.get_exec_args_set().clone(); + }; + + return status::success; +} + +void genindex_t::prepare_args_set(const execution_args_set_t *res, + const std::vector &inputs, + const std::vector &outputs, const scratchpad_t &scratchpad) { + // update the data of partition in/outputs args + for (const auto &mem_idx : res->get_mems_use_external_inputs()) { + mem_idx.first.set_data_handle(inputs[mem_idx.second].get_data_handle()); + } + for (const auto &mem_idx : res->get_mems_use_external_outputs()) { + mem_idx.first.set_data_handle( + outputs[mem_idx.second].get_data_handle()); + } + + grantor_t var_grantor = memory_planner_.internal_temporary_grantor( + scratchpad.get_buffer()); + + for (auto &mem_offkey : res->get_mems_use_internal_temporary()) { + mem_offkey.first.set_data_handle(var_grantor.get(mem_offkey.second)); + } +} + +status_t genindex_t::execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs) { + dnnl::stream p_stream = make_dnnl_stream(p_engine_, *g_stream); + + // each thread's own local resource + thread_local_cache_t res_cache; + execution_args_set_t *res = res_cache.get_or_add( + reinterpret_cast(this), resource_ctor_); + + temporary_scratchpad_t scratchpad( + memory_planner_.total_internal_temporary_size(), p_engine_, + *g_alloc_); + assertm(scratchpad.size() + >= memory_planner_.total_internal_temporary_size(), + "no enough scratchpad memory"); + prepare_args_set(res, inputs, outputs, scratchpad); + + constant_cache_t::cached_t c_buffer; + + for (size_t i = 0; i < subgraph_->execs_.size(); i++) { + if (subgraph_->is_constant_[i]) continue; + subgraph_->execs_[i]->execute(p_stream, res->get_exec_args()[i]); + } + + return status::success; +} +#ifdef DNNL_WITH_SYCL +status_t genindex_t::sycl_execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs, + const std::vector<::sycl::event> &sycl_deps, + ::sycl::event *sycl_event) { + if (p_engine_.get_kind() == engine::kind::gpu) return status::unimplemented; + return execute_impl(g_stream, inputs, outputs); +} +#endif +#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_OCL +status_t genindex_t::ocl_execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs, + const std::vector &ocl_deps, cl_event *ocl_event) { + // TODO: add support + return status::unimplemented; +} +#endif +} // namespace dnnl_impl +} // namespace graph +} // namespace impl +} // namespace dnnl diff --git a/src/graph/backend/dnnl/kernels/gen_index.hpp b/src/graph/backend/dnnl/kernels/gen_index.hpp new file mode 100644 index 00000000000..78751c206b0 --- /dev/null +++ b/src/graph/backend/dnnl/kernels/gen_index.hpp @@ -0,0 +1,91 @@ +/******************************************************************************* +* Copyright 2025 Intel Corporation +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*******************************************************************************/ +#ifndef GRAPH_BACKEND_DNNL_KERNELS_GEN_INDEX_HPP +#define GRAPH_BACKEND_DNNL_KERNELS_GEN_INDEX_HPP +#include +#include +#include +#include +#include + +#include "common/dnnl_thread.hpp" + +#include "graph/backend/dnnl/dnnl_constant_tensor_cache.hpp" +#include "graph/backend/dnnl/dnnl_partition_impl.hpp" +#include "graph/backend/dnnl/kernels/kernel_base.hpp" +#include "graph/backend/dnnl/passes/memory_planning.hpp" +#include "graph/backend/dnnl/scratchpad.hpp" +#include "graph/backend/dnnl/subgraph.hpp" +#include "graph/backend/dnnl/thread_local_cache.hpp" + +namespace dnnl { +namespace impl { +namespace graph { +namespace dnnl_impl { +using ltw = logical_tensor_wrapper_t; +struct genindex_t : public kernel_base_t { +private: + allocator_t *g_alloc_ = nullptr; + + std::shared_ptr subgraph_; + memory_planner_t memory_planner_; + + std::function()> resource_ctor_; + +public: + genindex_t() { + thread_local_cache_t res_cache; + res_cache.retain(); + } + ~genindex_t() override { + thread_local_cache_t res_cache; + res_cache.remove_if_exist(reinterpret_cast(this)); + res_cache.release(); + } + void prepare_args_set(const execution_args_set_t *res, + const std::vector &inputs, + const std::vector &outputs, + const scratchpad_t &scratchpad); + status_t compile_impl(const dnnl_partition_impl_t *part, + const engine_t *g_engine, + const std::vector &inputs, + const std::vector &outputs) override; + status_t execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs) override; +#ifdef DNNL_WITH_SYCL + status_t sycl_execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs, + const std::vector<::sycl::event> &sycl_deps, + ::sycl::event *sycl_event) override; +#endif +#if DNNL_GPU_RUNTIME == DNNL_RUNTIME_OCL + status_t ocl_execute_impl(const stream_t *g_stream, + const std::vector &inputs, + const std::vector &outputs, + const std::vector &ocl_deps, + cl_event *ocl_event) override; +#endif + DEF_KERNEL_METHOD_STR(genindex_t) + DNNL_DISALLOW_COPY_AND_ASSIGN(genindex_t) +}; + +} // namespace dnnl_impl +} // namespace graph +} // namespace impl +} // namespace dnnl +#endif diff --git a/src/graph/backend/dnnl/kernels/kernels.hpp b/src/graph/backend/dnnl/kernels/kernels.hpp index 5da6fe99392..a17d2502461 100644 --- a/src/graph/backend/dnnl/kernels/kernels.hpp +++ b/src/graph/backend/dnnl/kernels/kernels.hpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -24,6 +24,7 @@ #include "graph/backend/dnnl/kernels/conv_transpose.hpp" #include "graph/backend/dnnl/kernels/dummy.hpp" #include "graph/backend/dnnl/kernels/eltwise.hpp" +#include "graph/backend/dnnl/kernels/gen_index.hpp" #include "graph/backend/dnnl/kernels/group_norm.hpp" #include "graph/backend/dnnl/kernels/large_partition.hpp" #include "graph/backend/dnnl/kernels/layer_norm.hpp" diff --git a/src/graph/backend/dnnl/patterns/single_op_pattern.cpp b/src/graph/backend/dnnl/patterns/single_op_pattern.cpp index 73a22f6c45c..67f75dc4ef2 100644 --- a/src/graph/backend/dnnl/patterns/single_op_pattern.cpp +++ b/src/graph/backend/dnnl/patterns/single_op_pattern.cpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2020-2024 Intel Corporation + * Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -424,6 +424,18 @@ DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, reduce_pass) return std::make_shared(); }); +// GenIndex currently is CPU only +DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, gen_index_pass) + .set_priority(DEFAULT_P) + .set_engine_kind(engine_kind::cpu) + .set_kind(partition_kind_t::misc_post_ops) + .set_attr("FCreatePattern", + [](const std::shared_ptr &pgraph) -> void { + pgraph->append_op(graph::op_kind::GenIndex); + }) + .set_attr("FCreateKernel", + []() -> kernel_ptr { return std::make_shared(); }); + #undef DNNL_BACKEND_SINGLE_OP_TRANSFORM #undef DEFAULT_P From 524a84723421fa635ab6b0c22280995f5313a994 Mon Sep 17 00:00:00 2001 From: "Gu, Yonghao" Date: Sun, 29 Dec 2024 22:32:32 -0800 Subject: [PATCH 07/16] graph: dnnl: enable greater_equal in dnnl backend --- src/graph/backend/dnnl/passes/lower.cpp | 7 +++++++ src/graph/backend/dnnl/passes/utils.cpp | 5 +++-- src/graph/backend/dnnl/patterns/single_op_pattern.cpp | 1 + 3 files changed, 11 insertions(+), 2 deletions(-) diff --git a/src/graph/backend/dnnl/passes/lower.cpp b/src/graph/backend/dnnl/passes/lower.cpp index d1143cbbc0b..18589444b90 100644 --- a/src/graph/backend/dnnl/passes/lower.cpp +++ b/src/graph/backend/dnnl/passes/lower.cpp @@ -91,6 +91,12 @@ static status_t binary_handler( new_op->merge_attributes(op->get_attributes()); rewriter.replace_op(op, new_op); insert_empty_scratchpad(new_op); + if (op->get_kind() == graph::op_kind::GreaterEqual) { + auto out_vals = op->get_output_values(); + auto dst = out_vals[0]; + // GreaterEqual output's datatype is boolean. we treated it as u8 + dst->set_data_type(dnnl::impl::data_type::u8); + } return status::success; } @@ -816,6 +822,7 @@ static const std::unordered_map handler_table { ITEM(Divide, binary_handler), ITEM(Minimum, binary_handler), ITEM(Maximum, binary_handler), + ITEM(GreaterEqual, binary_handler), // eltwise fwd ITEM(Abs, eltwise_fwd_handler), ITEM(Clamp, eltwise_fwd_handler), diff --git a/src/graph/backend/dnnl/passes/utils.cpp b/src/graph/backend/dnnl/passes/utils.cpp index 571f9ea603a..98e7093d668 100644 --- a/src/graph/backend/dnnl/passes/utils.cpp +++ b/src/graph/backend/dnnl/passes/utils.cpp @@ -1,5 +1,5 @@ /******************************************************************************* - * Copyright 2021-2024 Intel Corporation + * Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -249,7 +249,8 @@ const std::map &get_binary_alg_map() { {graph::op_kind::Minimum, dnnl::algorithm::binary_min}, {graph::op_kind::Maximum, dnnl::algorithm::binary_max}, {graph::op_kind::Subtract, dnnl::algorithm::binary_sub}, - {graph::op_kind::BiasAdd, dnnl::algorithm::binary_add}}; + {graph::op_kind::BiasAdd, dnnl::algorithm::binary_add}, + {graph::op_kind::GreaterEqual, dnnl::algorithm::binary_ge}}; return binary_alg_map; } diff --git a/src/graph/backend/dnnl/patterns/single_op_pattern.cpp b/src/graph/backend/dnnl/patterns/single_op_pattern.cpp index 67f75dc4ef2..4c74a3cf2b8 100644 --- a/src/graph/backend/dnnl/patterns/single_op_pattern.cpp +++ b/src/graph/backend/dnnl/patterns/single_op_pattern.cpp @@ -291,6 +291,7 @@ DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, convtranspose_weights_bwd_pass) }); #endif +DNNL_BACKEND_SINGLE_OP_TRANSFORM(greater_equal_pass, GreaterEqual, binary_t) DNNL_BACKEND_SINGLE_OP_TRANSFORM(matmul_pass, MatMul, float_matmul) DNNL_BACKEND_SINGLE_OP_TRANSFORM(max_pool_pass, MaxPool, float_pooling_fwd) DNNL_BACKEND_SINGLE_OP_TRANSFORM(prelu_pass, PReLU, float_prelu_fwd) From b481d782db48e8ce132c9846f296ad44aafa57ac Mon Sep 17 00:00:00 2001 From: "Gu, Yonghao" Date: Sun, 29 Dec 2024 22:35:07 -0800 Subject: [PATCH 08/16] benchdnn: graph: support genindex and greaterequal op in benchdnn graph --- tests/benchdnn/graph/custom_driver.cpp | 71 ++++++++++++++++++++++++ tests/benchdnn/graph/custom_driver.hpp | 4 ++ tests/benchdnn/graph/flex_rewrite.cpp | 18 +++++- tests/benchdnn/graph/setting_handler.cpp | 7 ++- tests/benchdnn/graph/utils.cpp | 8 ++- 5 files changed, 105 insertions(+), 3 deletions(-) diff --git a/tests/benchdnn/graph/custom_driver.cpp b/tests/benchdnn/graph/custom_driver.cpp index d82b7c44864..2a267163bcb 100644 --- a/tests/benchdnn/graph/custom_driver.cpp +++ b/tests/benchdnn/graph/custom_driver.cpp @@ -32,6 +32,69 @@ extern "C" dnnl_status_t dnnl_memory_desc_create_with_string_tag( namespace custom { +namespace genindex { +// GENINDEX OP +// DNNL_ARG_SRC: src +// DNNL_ARG_DST: dst + +std::vector exec_args = { + DNNL_ARG_SRC, + DNNL_ARG_DST, +}; + +int init_ref_memory_args(dnn_mem_map_t &ref_mem_map, dnn_mem_map_t &mem_map, + const prb_t *prb, res_t *res) { + + const auto &ref_engine = get_cpu_engine(); + + for (auto &entry : mem_map) { + const int exec_arg = entry.first; + auto &mem = entry.second; + + ref_mem_map.emplace( + exec_arg, dnn_mem_t(mem.md_, dnnl_f32, tag::abx, ref_engine)); + auto &ref_mem = ref_mem_map[exec_arg]; + + switch (exec_arg) { + case DNNL_ARG_SRC: + // For GenIndex op, the input value doesn't affect the output + // value, it doesn't matter what value we fill in. + SAFE(::custom::fill_mem(mem, ref_mem, 0, 0), WARN); + break; + default: break; + } + } + return OK; +} + +int execute(const prb_t *prb, const args_t &args, res_t *res) { + dnn_mem_t &dst = const_cast(args.find(DNNL_ARG_DST)); + size_t axis = prb->axis; + auto ndims = dst.ndims(); + auto dims = dst.dims(); + auto strides = dst.strides(); + + benchdnn_parallel_nd(dst.nelems(), [&](int64_t index) { + // This function resembles dnn_mem_t::get_idx but has a format + axis + // peculiarity which can't be covered in get_idx function without + // sacrificing performance, and the current code as is. + size_t offdst = 0, result = 0; + for (int i = 0; i < ndims; i++) { + // calculate the idx on each dimension + int idx = index % dims[i]; + if ((size_t)i == axis) result = idx; + index /= dims[i]; + + // accumulate offset for each arg + offdst += strides[i] * idx; + if (index == 0) break; + } + dst.set_elem(offdst, result); + }); + return OK; +} +} // namespace genindex + namespace select { // SELECT OP // DNNL_ARG_WEIGHTS: cond @@ -208,6 +271,7 @@ dnnl_status_t init_pd(init_pd_args_t &init_pd_args) { std::vector supported_exec_args(const prb_t *prb) { std::vector exec_args; switch (prb->alg) { + case GENINDEX: return ::custom::genindex::exec_args; case SELECT: return ::custom::select::exec_args; case TRANSPOSE: return ::custom::transpose::exec_args; case RESHAPE: return ::custom::reshape::exec_args; @@ -219,6 +283,7 @@ std::vector supported_exec_args(const prb_t *prb) { void setup_cmp(compare::compare_t &cmp, const prb_t *prb, data_kind_t kind, const args_t &ref_args) { switch (prb->alg) { + case GENINDEX: case SELECT: case TRANSPOSE: case RESHAPE: cmp.set_zero_trust_percent(100.f); break; @@ -277,6 +342,11 @@ int init_ref_memory_args(dnn_mem_map_t &ref_mem_map, dnn_mem_map_t &mem_map, if (has_bench_mode_modifier(mode_modifier_t::no_ref_memory)) return OK; switch (prb->alg) { + case GENINDEX: + SAFE(::custom::genindex::init_ref_memory_args( + ref_mem_map, mem_map, prb, res), + WARN); + break; case SELECT: SAFE(::custom::select::init_ref_memory_args( ref_mem_map, mem_map, prb, res), @@ -304,6 +374,7 @@ void skip_unimplemented_prb(const prb_t *prb, res_t *res) {} int execute(const prb_t *prb, const args_t &args, res_t *res) { int ret = FAILED; switch (prb->alg) { + case GENINDEX: ret = ::custom::genindex::execute(prb, args, res); break; case SELECT: ret = ::custom::select::execute(prb, args, res); break; case TRANSPOSE: ret = ::custom::transpose::execute(prb, args, res); diff --git a/tests/benchdnn/graph/custom_driver.hpp b/tests/benchdnn/graph/custom_driver.hpp index 3ccf903e190..3566b3217f0 100644 --- a/tests/benchdnn/graph/custom_driver.hpp +++ b/tests/benchdnn/graph/custom_driver.hpp @@ -28,6 +28,7 @@ namespace custom { enum alg_t { + GENINDEX, SELECT, TRANSPOSE, RESHAPE, @@ -42,6 +43,7 @@ struct settings_t { ::std::unordered_map arg_mds_; ::std::vector order; + int64_t axis; alg_t alg; // A stub to be compliant with `base_settings_t`. @@ -51,6 +53,7 @@ struct settings_t { struct prb_t { prb_t(const settings_t &s) : arg_mds_(s.arg_mds_), alg(s.alg) { switch (alg) { + case GENINDEX: axis = s.axis; break; case TRANSPOSE: order = s.order; break; default: break; } @@ -58,6 +61,7 @@ struct prb_t { ::std::unordered_map arg_mds_; ::std::vector order; + int64_t axis; attr_t attr; alg_t alg; }; diff --git a/tests/benchdnn/graph/flex_rewrite.cpp b/tests/benchdnn/graph/flex_rewrite.cpp index fec7da56b09..ac586a91025 100644 --- a/tests/benchdnn/graph/flex_rewrite.cpp +++ b/tests/benchdnn/graph/flex_rewrite.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2022-2024 Intel Corporation +* Copyright 2022-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -236,6 +236,8 @@ void flex_rewrite::infer_output_shape( case dnnl::graph::op::kind::Exp: case dnnl::graph::op::kind::GELU: case dnnl::graph::op::kind::GELUBackward: + case dnnl::graph::op::kind::GenIndex: + case dnnl::graph::op::kind::GreaterEqual: case dnnl::graph::op::kind::HardSigmoid: case dnnl::graph::op::kind::HardSigmoidBackward: case dnnl::graph::op::kind::HardSwish: @@ -1136,6 +1138,18 @@ void flex_rewrite::dt_rewrite(deserialized_graph &dgraph) { return aop.kind_ == k; })) { dt_rewrite_norm(aop, str_dt); + } else if (aop.kind_ == "GenIndex") { + // GenIndex: only rewrite src dtype + aop.in_lts_[0].data_type_ = str_dt; + } else if (aop.kind_ == "GreaterEqual") { + // GreaterEqual: only rewrite src dtype when it's floating-point + if (std::any_of(fp_dts.begin(), fp_dts.end(), + [&aop](const dnnl_data_type_t &fp_dt) { + return aop.in_lts_[0].data_type_ == dt2str(fp_dt); + })) { + aop.in_lts_[0].data_type_ = str_dt; + aop.in_lts_[1].data_type_ = str_dt; + } } else { for (auto < : aop.in_lts_) { lt.data_type_ = str_dt; @@ -1211,6 +1225,8 @@ void flex_rewrite::update_output_info( case dnnl::graph::op::kind::Exp: case dnnl::graph::op::kind::GELU: case dnnl::graph::op::kind::GELUBackward: + case dnnl::graph::op::kind::GenIndex: + case dnnl::graph::op::kind::GreaterEqual: case dnnl::graph::op::kind::GroupNorm: case dnnl::graph::op::kind::HardSigmoid: case dnnl::graph::op::kind::HardSigmoidBackward: diff --git a/tests/benchdnn/graph/setting_handler.cpp b/tests/benchdnn/graph/setting_handler.cpp index 697b904e8a4..30f651e1ad9 100644 --- a/tests/benchdnn/graph/setting_handler.cpp +++ b/tests/benchdnn/graph/setting_handler.cpp @@ -153,6 +153,10 @@ ::custom::settings_t get_setting( ::custom::settings_t op_setting; auto opkind = opstr2kind(base_op_ref.kind_); switch (opkind) { + case ::graph::op::kind::GenIndex: + op_setting.alg = ::custom::alg_t::GENINDEX; + base_op_ref.get_attr_s64(op_setting.axis, "axis"); + break; case ::graph::op::kind::Select: op_setting.alg = ::custom::alg_t::SELECT; break; @@ -287,7 +291,8 @@ bool get_binary_alg(const deserialized_op &base_op_ref, ::binary::alg_t &alg) { {"Maximum", ::binary::alg_t::MAX}, {"Minimum", ::binary::alg_t::MIN}, {"Multiply", ::binary::alg_t::MUL}, - {"Subtract", ::binary::alg_t::SUB}}; + {"Subtract", ::binary::alg_t::SUB}, + {"GreaterEqual", ::binary::alg_t::GE}}; const auto &op_kind = base_op_ref.kind_; if (map_kind_to_alg.find(op_kind) == map_kind_to_alg.end()) return false; diff --git a/tests/benchdnn/graph/utils.cpp b/tests/benchdnn/graph/utils.cpp index 75a99010083..58247d8c36f 100644 --- a/tests/benchdnn/graph/utils.cpp +++ b/tests/benchdnn/graph/utils.cpp @@ -291,6 +291,8 @@ dnnl::graph::op::kind opstr2kind(const std::string &kind) { {"Exp", dnnl::graph::op::kind::Exp}, {"GELU", dnnl::graph::op::kind::GELU}, {"GELUBackward", dnnl::graph::op::kind::GELUBackward}, + {"GenIndex", dnnl::graph::op::kind::GenIndex}, + {"GreaterEqual", dnnl::graph::op::kind::GreaterEqual}, {"GroupNorm", dnnl::graph::op::kind::GroupNorm}, {"HardSigmoid", dnnl::graph::op::kind::HardSigmoid}, {"HardSigmoidBackward", dnnl::graph::op::kind::HardSigmoidBackward}, @@ -484,6 +486,9 @@ dnnl_driver_t opkind2driver(const dnnl::graph::op::kind &kind) { {dnnl::graph::op::kind::GELU, dnnl_driver_t::eltwise}, {dnnl::graph::op::kind::GELUBackward, dnnl_driver_t::eltwise}, + {dnnl::graph::op::kind::GenIndex, dnnl_driver_t::custom}, + {dnnl::graph::op::kind::GreaterEqual, + dnnl_driver_t::binary}, {dnnl::graph::op::kind::GroupNorm, dnnl_driver_t::gnorm}, {dnnl::graph::op::kind::HardSigmoid, dnnl_driver_t::eltwise}, @@ -839,7 +844,8 @@ int get_prim_arg_name_from_graph_op_input_offset( case dnnl::graph::op::kind::Maximum: case dnnl::graph::op::kind::Minimum: case dnnl::graph::op::kind::Multiply: - case dnnl::graph::op::kind::Subtract: { + case dnnl::graph::op::kind::Subtract: + case dnnl::graph::op::kind::GreaterEqual: { if (input_offset == 0) return DNNL_ARG_SRC_0; else if (input_offset == 1) From 127224657208d2f402b08d6fb5c46cfb6770b5a2 Mon Sep 17 00:00:00 2001 From: "Gu, Yonghao" Date: Wed, 1 Jan 2025 20:09:46 -0800 Subject: [PATCH 09/16] benchdnn: graph: inputs: add gen_index and greater_equal cases --- .../inputs/graph/op/f32/genindex.json | 65 ++++++++++++++ .../inputs/graph/op/f32/greaterequal.json | 84 +++++++++++++++++++ .../benchdnn/inputs/graph/op/harness_bf16_all | 6 ++ .../benchdnn/inputs/graph/op/harness_bf16_ci | 2 + .../benchdnn/inputs/graph/op/harness_f16_all | 6 ++ tests/benchdnn/inputs/graph/op/harness_f16_ci | 2 + .../benchdnn/inputs/graph/op/harness_f32_all | 6 ++ tests/benchdnn/inputs/graph/op/harness_f32_ci | 2 + 8 files changed, 173 insertions(+) create mode 100644 tests/benchdnn/inputs/graph/op/f32/genindex.json create mode 100644 tests/benchdnn/inputs/graph/op/f32/greaterequal.json diff --git a/tests/benchdnn/inputs/graph/op/f32/genindex.json b/tests/benchdnn/inputs/graph/op/f32/genindex.json new file mode 100644 index 00000000000..c41f2568cd5 --- /dev/null +++ b/tests/benchdnn/inputs/graph/op/f32/genindex.json @@ -0,0 +1,65 @@ +{ + "version": "3.7.0", + "engine_kind": "cpu", + "fpmath_mode": "strict", + "fpmath_mode_apply_to_int": "false", + "input_ports": [ + 0 + ], + "output_ports": [ + 1 + ], + "graph": [ + { + "id": 0, + "name": "genindex", + "kind": "GenIndex", + "attrs": { + "axis": { + "type": "s64", + "value": 0 + } + }, + "inputs": [ + { + "id": 0, + "dtype": "f32", + "shape": [ + 1, + 16, + 32, + 32 + ], + "stride": [ + 16384, + 1024, + 32, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 1, + "dtype": "s32", + "shape": [ + 1, + 16, + 32, + 32 + ], + "stride": [ + 16384, + 1024, + 32, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + } + ] +} diff --git a/tests/benchdnn/inputs/graph/op/f32/greaterequal.json b/tests/benchdnn/inputs/graph/op/f32/greaterequal.json new file mode 100644 index 00000000000..61e77474abe --- /dev/null +++ b/tests/benchdnn/inputs/graph/op/f32/greaterequal.json @@ -0,0 +1,84 @@ +{ + "version": "3.7.0", + "engine_kind": "cpu", + "fpmath_mode": "strict", + "fpmath_mode_apply_to_int": "false", + "input_ports": [ + 0, + 1 + ], + "output_ports": [ + 2 + ], + "graph": [ + { + "id": 0, + "name": "greaterequal", + "kind": "GreaterEqual", + "attrs": { + "auto_broadcast": { + "type": "string", + "value": "numpy" + } + }, + "inputs": [ + { + "id": 0, + "dtype": "f32", + "shape": [ + 1, + 16, + 32, + 32 + ], + "stride": [ + 16384, + 1024, + 32, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 1, + "dtype": "f32", + "shape": [ + 1, + 16, + 32, + 32 + ], + "stride": [ + 16384, + 1024, + 32, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 2, + "dtype": "boolean", + "shape": [ + 1, + 16, + 32, + 32 + ], + "stride": [ + 16384, + 1024, + 32, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + } + ] +} diff --git a/tests/benchdnn/inputs/graph/op/harness_bf16_all b/tests/benchdnn/inputs/graph/op/harness_bf16_all index d26c102a82f..7e4e8abc9aa 100644 --- a/tests/benchdnn/inputs/graph/op/harness_bf16_all +++ b/tests/benchdnn/inputs/graph/op/harness_bf16_all @@ -146,6 +146,12 @@ --reset --dt=bf16 --case=op/f32/gnorm.json --reset --dt=bf16 --case=op/f32/static_reshape.json --reset --dt=bf16 --case=op/f32/static_transpose.json +--reset --dt=bf16 --case=op/f32/genindex.json +--reset --dt=bf16 --op-attrs=0:axis:2 --case=op/f32/genindex.json +--reset --dt=bf16 --op-attrs=0:axis:2 --in-shapes=0:1x16x32x32*acbd+1:1x16x32x32*acbd --case=op/f32/genindex.json +--reset --dt=bf16 --case=op/f32/greaterequal.json +--reset --dt=bf16 --in-shapes=1:1x1x1x1 --case=op/f32/greaterequal.json +--reset --dt=bf16 --in-shapes=1:1 --case=op/f32/greaterequal.json # concat --reset --dt=bf16 --in-shapes=0:1x4096x14x14+1:1x4096x14x14 --case=op/f32/concat.json diff --git a/tests/benchdnn/inputs/graph/op/harness_bf16_ci b/tests/benchdnn/inputs/graph/op/harness_bf16_ci index 5480ddacce8..be5f6e792fe 100644 --- a/tests/benchdnn/inputs/graph/op/harness_bf16_ci +++ b/tests/benchdnn/inputs/graph/op/harness_bf16_ci @@ -26,3 +26,5 @@ --reset --dt=bf16 --case=op/f32/static_reshape.json --reset --dt=bf16 --case=op/f32/static_transpose.json --reset --case=op/bf16/dynamicdq_s4.json +--reset --dt=bf16 --case=op/f32/genindex.json +--reset --dt=bf16 --case=op/f32/greaterequal.json diff --git a/tests/benchdnn/inputs/graph/op/harness_f16_all b/tests/benchdnn/inputs/graph/op/harness_f16_all index 0ce472c5167..ee77a8943d8 100644 --- a/tests/benchdnn/inputs/graph/op/harness_f16_all +++ b/tests/benchdnn/inputs/graph/op/harness_f16_all @@ -146,6 +146,12 @@ --reset --dt=f16 --case=op/f32/gnorm.json --reset --dt=f16 --case=op/f32/static_reshape.json --reset --dt=f16 --case=op/f32/static_transpose.json +--reset --dt=f16 --case=op/f32/genindex.json +--reset --dt=f16 --op-attrs=0:axis:2 --case=op/f32/genindex.json +--reset --dt=f16 --op-attrs=0:axis:2 --in-shapes=0:1x16x32x32*acbd+1:1x16x32x32*acbd --case=op/f32/genindex.json +--reset --dt=f16 --case=op/f32/greaterequal.json +--reset --dt=f16 --in-shapes=1:1x1x1x1 --case=op/f32/greaterequal.json +--reset --dt=f16 --in-shapes=1:1 --case=op/f32/greaterequal.json # concat --reset --dt=f16 --in-shapes=0:1x4096x14x14+1:1x4096x14x14 --case=op/f32/concat.json diff --git a/tests/benchdnn/inputs/graph/op/harness_f16_ci b/tests/benchdnn/inputs/graph/op/harness_f16_ci index 62c2a05f3ff..8f046c03975 100644 --- a/tests/benchdnn/inputs/graph/op/harness_f16_ci +++ b/tests/benchdnn/inputs/graph/op/harness_f16_ci @@ -26,3 +26,5 @@ --reset --dt=f16 --case=op/f32/static_reshape.json --reset --dt=f16 --case=op/f32/static_transpose.json --reset --case=op/f16/dynamicdq_s4.json +--reset --dt=f16 --case=op/f32/genindex.json +--reset --dt=f16 --case=op/f32/greaterequal.json diff --git a/tests/benchdnn/inputs/graph/op/harness_f32_all b/tests/benchdnn/inputs/graph/op/harness_f32_all index bc471a40a88..ff8781a57e1 100644 --- a/tests/benchdnn/inputs/graph/op/harness_f32_all +++ b/tests/benchdnn/inputs/graph/op/harness_f32_all @@ -952,3 +952,9 @@ --reset --case=op/f32/gnorm.json --reset --case=op/f32/static_reshape.json --reset --case=op/f32/static_transpose.json +--reset --case=op/f32/genindex.json +--reset --op-attrs=0:axis:2 --case=op/f32/genindex.json +--reset --op-attrs=0:axis:2 --in-shapes=0:1x16x32x32*acbd+1:1x16x32x32*acbd --case=op/f32/genindex.json +--reset --case=op/f32/greaterequal.json +--reset --in-shapes=1:1x1x1x1 --case=op/f32/greaterequal.json +--reset --in-shapes=1:1 --case=op/f32/greaterequal.json diff --git a/tests/benchdnn/inputs/graph/op/harness_f32_ci b/tests/benchdnn/inputs/graph/op/harness_f32_ci index 6585457760d..addf165bb15 100644 --- a/tests/benchdnn/inputs/graph/op/harness_f32_ci +++ b/tests/benchdnn/inputs/graph/op/harness_f32_ci @@ -58,3 +58,5 @@ --reset --case=op/f32/select.json --reset --case=op/f32/static_reshape.json --reset --case=op/f32/static_transpose.json +--reset --case=op/f32/genindex.json +--reset --case=op/f32/greaterequal.json From 69181fc4aa76ef00741361eff70f2adfc59f60b6 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Thu, 26 Dec 2024 17:23:45 -0800 Subject: [PATCH 10/16] graph: utils: pm: support multi-consumers input for repetition --- src/graph/utils/pm/nested_matcher.cpp | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) diff --git a/src/graph/utils/pm/nested_matcher.cpp b/src/graph/utils/pm/nested_matcher.cpp index a6948b4d003..01caf9179bc 100644 --- a/src/graph/utils/pm/nested_matcher.cpp +++ b/src/graph/utils/pm/nested_matcher.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2021-2024 Intel Corporation +* Copyright 2021-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -364,7 +364,20 @@ bool node_outputs_matcher_t::match_op_consumers() { auto node_consumer = current_node_output_.second[k]; pb_node_t *out_node = node_consumer->first; // check if the out_node has been matched by previous out_ops - if (node_oport_matched_cons.count(k)) continue; + if (node_oport_matched_cons.count(k)) { + if (out_node->get_node_kind() == pb_node_kind::PB_NODE_KIND_OP) + continue; + if (out_node->get_node_kind() + == pb_node_kind::PB_NODE_KIND_ALTERNATION) + continue; + // For repetition case, check if multi consumers exist + repetition_t *rep_node = dynamic_cast(out_node); + if (rep_node->get_body()->get_inner_consumer(0) == nullptr) + continue; + if (rep_node->get_body()->get_inner_consumer(0)->size() == 1) + continue; + } + binding_t out_bind(BIND_IN, out_op, op_consumer.get_offset(), out_node, node_consumer->second); From df15f6b489e352e9a09218f6f0c8963f2d0bed27 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Sun, 29 Dec 2024 22:36:17 -0800 Subject: [PATCH 11/16] graph: backend: dnnl: add pattern for sdp with implicit causal mask --- src/graph/backend/dnnl/patterns/sdp.cpp | 34 ++++++++++++++++++++++- src/graph/backend/dnnl/patterns/utils.hpp | 28 +++++++++++++++++++ 2 files changed, 61 insertions(+), 1 deletion(-) diff --git a/src/graph/backend/dnnl/patterns/sdp.cpp b/src/graph/backend/dnnl/patterns/sdp.cpp index 297d817e8c0..8fd49c2b725 100644 --- a/src/graph/backend/dnnl/patterns/sdp.cpp +++ b/src/graph/backend/dnnl/patterns/sdp.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2023-2024 Intel Corporation +* Copyright 2023-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -146,6 +146,38 @@ DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, float_sdp_fusion) return std::make_shared>(); }); +DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, float_sdp_implicit_mask_fusion) + .set_priority(21.0f) + .set_engine_kind(engine_kind::cpu) + .set_kind(partition_kind_t::sdp) + .set_attr("FCreatePattern", + [](const std::shared_ptr &pgraph) -> void { + auto matmul_qk = pgraph->append_op(graph::op_kind::MatMul); + + std::shared_ptr scale_graph; + scale_graph = std::make_shared(); + auto scale = scale_graph->append_alternation( + {graph::op_kind::Divide, graph::op_kind::Multiply}); + scale_graph->create_input_port(0, scale, 0); + scale_graph->create_output_port(0, scale, 0); + auto optional_scale = pgraph->append_optional( + scale_graph, {in_edge(0, matmul_qk, 0)}); + // TODO: merge implicit causal mask and explicit mask + // into one pattern + auto optional_mask + = optional_causal_mask(pgraph, optional_scale); + + auto softmax = pgraph->append_op(graph::op_kind::SoftMax, + {in_edge(0, optional_mask, 0)}); + auto matmul_v = pgraph->append_op( + graph::op_kind::MatMul, {in_edge(0, softmax, 0)}); + // Optional transpose + reshape/reorder + optional_transpose_reshape(pgraph, matmul_v, 0); + }) + .set_attr("FCreateKernel", []() -> kernel_ptr { + return std::make_shared(); + }); + DNNL_BACKEND_REGISTER_PATTERN_MATCHER_PASS(dnnl, float_gqa_fusion) .set_priority(21.1f) .set_kind(partition_kind_t::sdp) diff --git a/src/graph/backend/dnnl/patterns/utils.hpp b/src/graph/backend/dnnl/patterns/utils.hpp index b5aaccb1ef6..73e58680db3 100644 --- a/src/graph/backend/dnnl/patterns/utils.hpp +++ b/src/graph/backend/dnnl/patterns/utils.hpp @@ -400,6 +400,34 @@ inline graph::utils::pm::repetition_t *optional_select( return pselect; } +// Implicit Causal Mask +inline graph::utils::pm::repetition_t *optional_causal_mask( + const std::shared_ptr &pgraph, + graph::utils::pm::pb_node_t *scaled_output) { + auto popt_graph = std::make_shared(); + + graph::utils::pm::pb_op_t *gen_index_row + = popt_graph->append_op(graph::op_kind::GenIndex); + graph::utils::pm::pb_op_t *gen_index_col + = popt_graph->append_op(graph::op_kind::GenIndex); + graph::utils::pm::pb_op_t *greater_equal + = popt_graph->append_op(graph::op_kind::GreaterEqual, + graph::utils::pm::in_edges_t {in_edge(0, gen_index_row, 0), + {in_edge(1, gen_index_col, 0)}}); + graph::utils::pm::pb_op_t *select = popt_graph->append_op( + graph::op_kind::Select, + graph::utils::pm::in_edges_t {in_edge(0, greater_equal, 0)}); + + popt_graph->create_input_port(0, gen_index_row, 0); + popt_graph->create_input_port(0, gen_index_col, 0); + popt_graph->create_input_port(0, select, 1); + popt_graph->create_input_port(1, select, 2); + popt_graph->create_output_port(0, select, 0); + auto pmask = pgraph->append_optional(popt_graph, + graph::utils::pm::in_edges_t {in_edge(0, scaled_output, 0)}); + return pmask; +} + // Optional (transpose + reorder/staticReshape) inline graph::utils::pm::repetition_t *optional_transpose_reshape( const std::shared_ptr &pgraph, From a8c89aa44f1d05b5f1ff3f1ce1696c3bb1182122 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Wed, 1 Jan 2025 22:35:35 -0800 Subject: [PATCH 12/16] benchdnn: graph: inputs: add a sdpa implicit causal mask case --- .../graph/complex_fusion/harness_mha_all | 1 + .../graph/complex_fusion/harness_mha_ci | 1 + ...a-plain-implicit-causal-mask-fp32-bs1.json | 530 ++++++++++++++++++ 3 files changed, 532 insertions(+) create mode 100644 tests/benchdnn/inputs/graph/complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json diff --git a/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_all b/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_all index 05399f458fd..13d8e7ccd6d 100644 --- a/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_all +++ b/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_all @@ -14,6 +14,7 @@ --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/GQA-fp16.json --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-wo-mask-f16.json --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-scale-by-mul-f16.json +--reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json # int8 graphs --reset --case=complex_fusion/mha/MHA-GPT-inf-int8-bs1.json diff --git a/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_ci b/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_ci index 8b86b687abc..d593483363e 100644 --- a/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_ci +++ b/tests/benchdnn/inputs/graph/complex_fusion/harness_mha_ci @@ -12,6 +12,7 @@ --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/GQA-fp16.json --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-wo-mask-f16.json --reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-scale-by-mul-f16.json +--reset --dt=f32,bf16,f16 --case=complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json # int8 graphs --reset --case=complex_fusion/mha/MHA-GPT-inf-int8-bs1.json diff --git a/tests/benchdnn/inputs/graph/complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json b/tests/benchdnn/inputs/graph/complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json new file mode 100644 index 00000000000..b16217b31ca --- /dev/null +++ b/tests/benchdnn/inputs/graph/complex_fusion/mha/sdpa-plain-implicit-causal-mask-fp32-bs1.json @@ -0,0 +1,530 @@ +{ + "version": "3.7.0", + "engine_kind": "cpu", + "fpmath_mode": "strict", + "fpmath_mode_apply_to_int": "false", + "input_ports": [ + 0, + 1, + 3, + 8, + 11 + ], + "output_ports": [ + 12 + ], + "graph": [ + { + "id": 0, + "name": "matmul_qk", + "kind": "MatMul", + "attrs": { + "transpose_a": { + "type": "bool", + "value": 0 + }, + "transpose_b": { + "type": "bool", + "value": 1 + } + }, + "inputs": [ + { + "id": 0, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 64 + ], + "stride": [ + 393216, + 24576, + 64, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 1, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 64 + ], + "stride": [ + 393216, + 24576, + 64, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 2, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 1, + "name": "scale_mul", + "kind": "Multiply", + "attrs": { + "auto_broadcast": { + "type": "string", + "value": "numpy" + } + }, + "inputs": [ + { + "id": 2, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 3, + "dtype": "f32", + "shape": [ + 1 + ], + "stride": [ + 1 + ], + "layout_type": "strided", + "property_type": "constant" + } + ], + "outputs": [ + { + "id": 4, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 2, + "name": "genindex_row", + "kind": "GenIndex", + "attrs": { + "axis": { + "type": "s64", + "value": 2 + } + }, + "inputs": [ + { + "id": 4, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 5, + "dtype": "s32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 3, + "name": "genindex_col", + "kind": "GenIndex", + "attrs": { + "axis": { + "type": "s64", + "value": 3 + } + }, + "inputs": [ + { + "id": 4, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 6, + "dtype": "s32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 4, + "name": "mask_greater_equal", + "kind": "GreaterEqual", + "attrs": { + "auto_broadcast": { + "type": "string", + "value": "numpy" + } + }, + "inputs": [ + { + "id": 5, + "dtype": "s32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 6, + "dtype": "s32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 7, + "dtype": "boolean", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 5, + "name": "Select", + "kind": "Select", + "attrs": { + "auto_broadcast": { + "type": "string", + "value": "numpy" + } + }, + "inputs": [ + { + "id": 7, + "dtype": "boolean", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 4, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 8, + "dtype": "f32", + "shape": [ + 1 + ], + "stride": [ + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 9, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 6, + "name": "softmax", + "kind": "SoftMax", + "attrs": { + "axis": { + "type": "s64", + "value": -1 + } + }, + "inputs": [ + { + "id": 9, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 10, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + }, + { + "id": 7, + "name": "matmul_v", + "kind": "MatMul", + "attrs": { + "transpose_a": { + "type": "bool", + "value": 0 + }, + "transpose_b": { + "type": "bool", + "value": 0 + } + }, + "inputs": [ + { + "id": 10, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 384 + ], + "stride": [ + 2359296, + 147456, + 384, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + }, + { + "id": 11, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 64 + ], + "stride": [ + 393216, + 24576, + 64, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ], + "outputs": [ + { + "id": 12, + "dtype": "f32", + "shape": [ + 1, + 16, + 384, + 64 + ], + "stride": [ + 393216, + 24576, + 64, + 1 + ], + "layout_type": "strided", + "property_type": "undef" + } + ] + } + ] +} \ No newline at end of file From c9dab3b7946df40d0debc8c77eabf7ec7e760dbc Mon Sep 17 00:00:00 2001 From: "Gu, Yonghao" Date: Thu, 2 Jan 2025 09:16:36 +0000 Subject: [PATCH 13/16] gtests: graph: api: add gtest for GenIndex and GreaterEqual --- tests/gtests/graph/api/test_cpp_api_op.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tests/gtests/graph/api/test_cpp_api_op.cpp b/tests/gtests/graph/api/test_cpp_api_op.cpp index b0f05894ed9..a8b80217012 100644 --- a/tests/gtests/graph/api/test_cpp_api_op.cpp +++ b/tests/gtests/graph/api/test_cpp_api_op.cpp @@ -1,5 +1,5 @@ /******************************************************************************* -* Copyright 2020-2024 Intel Corporation +* Copyright 2020-2025 Intel Corporation * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -114,6 +114,8 @@ TEST(APIOp, CreateAllOps) { op::kind::Select, op::kind::Pow, op::kind::GroupNorm, + op::kind::GenIndex, + op::kind::GreaterEqual, }; // clang-format on From f30023bd445edafc90e4b211cb84f455e73215c9 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Fri, 3 Jan 2025 23:31:15 -0800 Subject: [PATCH 14/16] doc: graph: add document for GenIndex and GreaterEqual --- doc/graph/operations/GenIndex.md | 39 ++++++++++++++++++++++ doc/graph/operations/GreaterEqual.md | 49 ++++++++++++++++++++++++++++ 2 files changed, 88 insertions(+) create mode 100644 doc/graph/operations/GenIndex.md create mode 100644 doc/graph/operations/GreaterEqual.md diff --git a/doc/graph/operations/GenIndex.md b/doc/graph/operations/GenIndex.md new file mode 100644 index 00000000000..ff3306633dc --- /dev/null +++ b/doc/graph/operations/GenIndex.md @@ -0,0 +1,39 @@ +GenIndex{#dev_guide_op_genindex} +================================ + +## General + +The GenIndex operation creates an index tensor along a specified axis of +an input tensor. The resulting index tensor has the same shape as the +input tensor, with each element representing the index along the +specified axis. + +## Operation Attributes + +| Attribute Name | Description | Value Type | Supported Values | Required or Optional | +|:------------------------------------------|:----------------------------------------------------------------|:-----------|:-----------------------------------------------------------|:---------------------| +| [axis] (@ref dnnl::graph::op::attr::axis) | Specifies the dimension along which index values are generated. | s64 | An s64 value in the range of [-r, r-1] where r = rank(src) | Required | + +## Execution Arguments + +### Input + +| Index | Argument Name | Required or Optional | +|:------|:--------------|:---------------------| +| 0 | `src` | Required | + +### Output + +| Index | Argument Name | Required or Optional | +|:------|:--------------|:---------------------| +| 0 | `dst` | Required | + +## Supported Data Types + +The GenIndex operation supports the following data type combinations. + +| Src | Dst | +|:-------|:-------| +| f32 | s32 | +| bf16 | s32 | +| f16 | s32 | diff --git a/doc/graph/operations/GreaterEqual.md b/doc/graph/operations/GreaterEqual.md new file mode 100644 index 00000000000..912b066b45c --- /dev/null +++ b/doc/graph/operations/GreaterEqual.md @@ -0,0 +1,49 @@ +GreaterEqual{#dev_guide_op_greaterequal} +======================================== + +## General + +The GreaterEqual operation performs an element-wise greater-than-or-equal +comparison between two given tensors. This operation applies +the multi-directional broadcast rules to ensure compatibility between +the tensors of different shapes. + +\f[ dst = \begin{cases} true & \text{if}\ src_0 \ge src_1 \\ + false & \text{if}\ src_0 < src_1 \end{cases} \f] + +## Operation Attributes + +| Attribute Name | Description | Value Type | Supported Values | Required or Optional | +|:-------------------------------------------------------------|:-----------------------------------------------------------|:-----------|:-------------------------|:---------------------| +| [auto_broadcast](@ref dnnl::graph::op::attr::auto_broadcast) | Specifies rules used for auto-broadcasting of src tensors. | string | `none`,`numpy` (default) | Optional | + +## Execution Arguments + +### Input + +| Index | Argument Name | Required or Optional | +|:------|:--------------|:---------------------| +| 0 | `src_0` | Required | +| 1 | `src_1` | Required | + +@note Both src shapes should match and no auto-broadcasting is allowed if +the `auto_broadcast` attribute is `none`. `src_0` and `src_1` shapes can be +different and auto-broadcasting is allowed if the `auto_broadcast` attribute +is `numpy`. Broadcasting is performed according to the `auto_broadcast` value. + +### Output + +| Index | Argument Name | Required or Optional | +|:------|:--------------|:---------------------| +| 0 | `dst` | Required | + +## Supported Data Types + +The GreaterEqual operation supports the following data type combinations. + +| Src_0 / Src_1 | Dst | +|:--------------|:---------| +| f32 | boolean | +| bf16 | boolean | +| f16 | boolean | +| s32 | boolean | From 6ee19e03f2cb7c66c1d91634151dd037b62e6ae6 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Sun, 5 Jan 2025 21:17:23 -0800 Subject: [PATCH 15/16] doc: graph: update sdpa document to include implicit causal mask --- .../fusion_patterns/images/sdpa-mask-3.png | Bin 0 -> 182858 bytes doc/graph/fusion_patterns/images/sdpa.png | Bin 29134 -> 182345 bytes doc/graph/fusion_patterns/sdpa.md | 39 ++++++++++++------ 3 files changed, 27 insertions(+), 12 deletions(-) create mode 100644 doc/graph/fusion_patterns/images/sdpa-mask-3.png diff --git a/doc/graph/fusion_patterns/images/sdpa-mask-3.png b/doc/graph/fusion_patterns/images/sdpa-mask-3.png new file mode 100644 index 0000000000000000000000000000000000000000..339a9589122d1dc3870fff8f4cea5f64d351421e GIT binary patch literal 182858 zcmeEP2UrtL6GlH&6i~$8v0w*DLIQ{=y@OPxs2CtXAP`6bh>8`lcd&pJ6)TDe*c*2J z5xW6Y>=k>*_U|RP$&ovXl#n1kJ_?sh_Od(k&dfWrvvZvstyEidZlR!{plV}n?xLX3 zn5Cf5XqZw{a3%KR>yhBUMqC#wQ-$4^yS`9p9Q&DP;lT@?5Wr+H6f`Mj!k;wBIzenM zPm^M&SH#a_}VeS=UP%JdfF5kxZK~)+RedU zlWYclW-&t;;6E#dUkDp{i&+qd9SW{kk|-o?@X>-XUFiXI4iobc$TLPUI9w(>O!zg( zG;k}pio6gGo}MebOhPYl{pc(PdP#VbFPr1f;0P~)uSM3RSZL~)fOjK*DQ4)DfM3Gs z;Bf>~6%ek)^$!O#fega*7hdEsm=hvxJfoeN5tgpeQFZ|yQ4=LU4t{9p1S<|bJjj9V z&tQR%@SlJ_5tW8Mrhkm^iats3NMxNr4jhLh92U9Y%A5$l3%o6o=`Z@q!W(&PHjBp$ zhw~u3=*JEVWB8#9i+&=V!$DS4@D0&R0c;lf?a&LzjF!3sx!{VMQ{?JsravzT{T5_B z;jqYcYX&ng2#!pm83zhgUD2mN21f2f1(GCGe&jjB z>EbDZ--yBCF`&dF-=_Sf6EC0@0{{Y#69YaLzFh~Mc2r%+y6ErGV&Nmfdx9{6r;px8 zN98?G^wG#eS4JklJ*Y^P|M&>jN=P2za)WV^p)t!XR4cNMFZjYV7LagLe`XZ;F%bC? z2m(-Q7I?|hoEZg_9^L~kU~a2XVWI0#$z&jBcwwQLlPtkBqe37d6lDsJ5WNJ#LT51p z5dilC`hKQ_QUytSAO@A9z9dRclw?$dpRR z8zMi*Qp&cxQE>9+K<9>l$0igv6MugoSgs5fVl==Hk-orB2oI`lp#zH_#MAPpb3*!S zk}0VC3A_dHtBwwxtfMmkEIkgqsHPU??NMU@L-(!~jrD?5PaK%XGI z-e9Z{2b-!_3Gk2)V@1WScY|*t>6_@mr7Taq;XSf#D?q3w;hVTHgi%u&=E9Xn2UgjI z0i5`7B!nx=PT<@c;TzpJ^e}EXn=9u|jd0d#*CWw~-bQ&+2EAh0mNGO(#CX8M3B*_r zMwq_|5~U)ZB8$%DGW~${$3!>cPzZ5Gac+Sq+7ZMm5^idxVi$&gsra{i*Pk$Q!9*rR zwjEyMFjzo^qDlwbMA1rR9zpydPHZNib42kI6bvXlG}J=Dfw=5QjvoVk69NH~+r zQ8ZT8kUTmkkio0MkQ6{7fsjlTB$oe}=-q?EkX$IPQXlxCYq2^sMT?v&Z8dZK(5GEPK zX_Gp0!>X!?!UraZ8!NPBsG3x$PAFej&(w(u!xyQl6ID{3BrT}?>Jp01b!*P59>q&v z|4_ZEQ2$VjuV?zF2OT-7>K_f}a^a};NB{6Sx@y`o>BeNz*FRLRD%3xe(5`3tr!T1q zQgJ7740Qd`KYUuVn)+9DTUMq|SkGkGK+=cPLy4){i-$-$l6FhaK$|QPhRLrJR~cR! zWHn$Z7nczj3V_=}e95$;tczbHCrVI%h7b;=Q!wDIXkJBMkdUEsV2v2^yVyVoc|pik zAmwC&*QCc(RA@5lr^_Z0yHKF^)t7Xu72Y!ceBkMRB2f05cfiHsSXAo%p~_rFVw!AHe?yq-jXnP@@&?(+rNE#W*=<~!GQ>WoB` zd=SREg}-B_7&6tYHvf#6VndKrBk^i>enSED0Hg zKk_GkM5MKp=g%8+MYKNI_8i7Y!Pl^0rT zXm)Epu}F~wWcWKAQf`B2GM+R?y2l#y8bu{j&~1c* zDSxnuuZxC8RBB<%$VgszB#4lN-U5gc8)6rc*C>v%@xG%0bQYl$>WVo(84yH4L|(?e zRv4{|cMxM^e8H=5({kw-;8pP+Sg#llFDp~i$Wej56l=dAhCN#k-Cru+?973s{zUwV zz%ds;hyNY7-&_Zv|mhbAnY#@0fJ#%Bz7?e82B0Z>Y}6& z#_oysDwk$xz}`X8hE<^f$C)>Rjmu@NndH*ov4#W#1Bf*Q7LNk)2+Nd+-UuBisZu3` zD~XR(aa(qPKaJt*$kkccTHJ9gm4%0k3Kt#<0>;9S5cxr-47>1l!}^<;J~hFp)kTB z*kmT*YzUH>RWN8W zr60IJQc_rJUOzCBR+`=6~<5$+=a>34#1tbs_>oVxj8`D#|aL!6VvJl{s7NG z=nm9wsFo+8*!k5+xLU)HcCHNkh+9WG}RUh$_ozuLpUOh(MExPM(pb z>>xWGf2x~{v#$x6pyw;j8>vpb4TVAx?S{v|jV`2lupS4Q zM#F8{kZQ7oa5?J%>4?yc&^=+WeA&_1jUhteBoGQL0qGlHGSODw--Ah?=x*SsPu6mw z#%Q6rVF@y`EOjp@#fA|K=&fSq0$!62V-JZSRDeYxd+ibS&>X`N$)Jh>H79V;DhKGo zS+C{&6p~X0F`(j9c4?si_sisd2SQvxr5_WF;9y>Av|AOrG?l6LLIO(I%Y^WYgUA{Q zuIroHL!pIQC`o_RJ~2m?9J0z<537q{L-g`78NIt|ZS)951fPfy_NA6ODTD0?`FETg zfkDJYY2{BVgkGXu51^5-P>WD$qtZ|}17iwhj*}x?YJ63YpyjHE&pIuz?YYe05=Y7+h+&^cV$)m$1IBqzS0|_=kiK zfDobt>=kd&aA;peyV)c+MvK}cVwPP#G#s4hA<*WU^KZo^GvX>4PoOPTZxm4Z3?5?p z2@f0*I0YC+9?RVOGXm(~gcg*~hXFM5xGcIagJnwh3kei3IndsS-yobj0&7ywyk?4l zD1nUv9SsJ9Fm$bXzZpNMsA(J=RBk(^iZy8*rH11h}CL zbCejcRPZWzR?G`HfDy}?4Auq`7}gpL&S7q%9M@e`){XVup<@S)kW_tlLfL@6yZjw|ix z75Xsfa$^yOkX!|H8p5d6dmjcL^jL8iba05MMMNaMW)I7)^?eJq>zha(UMk`lWg0rl z6K;%Ifk%@{u*mM>Wq3!VzHP2DJr_SY=ue ze#u&GOA47&`1EkODrOagj|U%8R?BDzVF3z>lpV+_0zlat-pdR(Wzfm1mRd_nDbzWr zV>PeUl6!{ zpTN>WJ5y9F`UE|W2tN>wxUoSO;eFK^*g(rE(#w7{%Rh{2Nsp%baw*|K@SBw9%~mIF zsdkk=jt?*0T!nCw1d~~o9F+Dk2#eaFyM;P1qCNH)w^Y568j1u)1Z0odB$NOd3RekPz*^L2%woY91goy_nsmZQ5=^VXut>5* zP6`gEL)`elO^n_YeBURRrax2uR3ty3~mg5c}wc)K% zAgGn&twHF;qCb2a48EmWkq!1o-lrY|jyqW{G<80Dc+J#)}?LAXq!8HoR&P6nzHQfQP>!4rs-6auf@ zy6-R)!lZ!Zx@$&xT7t_7l|`7BCNPK!_9UwiXav<-@?o$^kPJXutpkik&;;()5x2>( z5MMPF^Wycal}35uQ!!nFW)v!>T2VZ!K5SRIkPH#&+U>AGNFd>JEdN%Ejq>EF2H{9c zR0j;bMhHHb07CYvM~P<#&D%(MTT>O-R@bGtya?=h5>*YHu*6g)n#hzVBo!M9G+3y<=ZXozVN{OI;lXQ_ zwkMTgErkpwQbvtyEdz^PAX6Y&xLNbXFZgmeh^5u0&*As1)$TI*S~$376?1CB2ZWC) zSG8~?G=V2&_f!h~WvT}Z()S>%r}4N-&f3>_n6ri$1I?(M>bs~~ufFIzAv=arHJD7G zi9h~e8gc69@4H)DTACZsXn>8vxehdQk|i(-!tW z)nuH2P8%xf6kRe1kp;a+F?Jv71J>@-(A%I+NVj&!foc*(axkHrvutB*&HQOLzR{2_ zmNMQ~i+Hat9Z`LHlu(rj*imAeQx6k?_0s4=%n$Fio(bi%{q4W9p@x`M_X{ z89R*2qlfVt07evk3U*ZE4qZcD20Vuj>b#I)Tq}333uHH)eLzix zyyyeTptN{OK_SRGrH7P4+4g36q+%IYnAN^rky_Cz4^fC#pTaD76e zaKTmWZ}G+u!6aWgiGdV0R+MWNK4#q$!y@g%#8xoe!qUPpJqtCW?rCATLj>ajbA2Kj zmLIyqWjMc(^OtltWMKVkxf>ijcfcuJi{=v^uIr&2GBHNSE6Q!cNYrr+b|8P1YQ+O9oFuxpnc#s3zp8*ao(D8@#C`BJNAFvsmo-XRCH#CF8&m%NGRXdab zNPFcxRAjl#7%XI=pdg@(B^HFM^KrU?EH2km&n?Eo+ASnJ($v{eUsnqq0|QbHbe=G*4g(U{ z$qj=^Bv_k+HA)c6>cl7!E;pW*5p=EvLgxz-C^QxiOgZKVRUimFW;hGHWNFTfV)`@S zJ>UZ7Hk?tbPDZMd!ctDQyR`7&vfZ^SEEKrta9q7*hDhvX2;ilFMj z5?06xMX46Wq^=LMs3xXRp%Yu4n8*jFpp-U7OkCLkJO}!Ox_|{)wIQ_1z=C|NH_Am| z*1N%FV24VSAdq1AFaP?D-x1@&U~~K#oU$x|E04}&`nfQ;LG*CoF<6%6G~nC~;TzpJ z^e}EXn=1`xL(Z)TsZy>fbXR5D+raO(!PyVl?;$EdViSZRWF73TJYpwZU{R|OuVCU4 z=w4vg9A2w*>bNqGAP$@)woHhaJBWZ(18tbufdNs?=ZWl7j|c?=rGYS( zx0))#s6rJfE`}EcY?b4bYeYp*+5($vC7dxqyjvMDmS7`gN|PkJ;EJ}_)t(xOdWE6N zD>_9H4k@=f1pTtQRidhH<&z2JLrJ6xl~2GQ$i4DGs|n2*O&VrQQsv16qsp%e;z{Rf z+OevuMO771_`o${%TSuFLUpPaUZEbuU#Y4Sj(w@Iopp8X5)zz0R!v=!Tc5CA=!ZU( zOsVRJJ}jSrqw33H$sJ(;K*A?6$SC)2H2;^;bDFk zn~Bmz-bk|n!E548o5=PE(XP#^t5Q|m#LB84S*Jo{BUlT$S5qAcZ;t}cw1LV}3py&< z*MtHI4mf6&6%?idYPXo21m_5aHqo3Rna0XU1bj?zzf4w>6Z^6FEG$f`BD6b5#fo=X z(IO(*NFf*Uf)-a$|6~GgAnAsK3Q*9}7QPOy`LS7S4!|4iMG8vjWA38`BAk1-J&oF=Z*|BvY0djL;`HW4n}2A zGo@P8PArBspMt8f?RtN=wV^W9U4>StUZj_Fp#gxtA}m9KVpsDGMMq1*A**bI z>Y9&(m7SRMk|o`b2S8R(@F@mqvh1KE1Ygi7>MwJ-wTqgPQX;A-z`!O%QrU)0GCt!` z$>Toj%3{i@Owc4usT|7$8W&@b0*4V5#fzmqow9;JQG1>m4+1M2&{f7Qm+cA+ane_M zR+4If9~I)A{2NLlVqRDgD(!$~WF#*<5`<8qJqry{jW9+Q{VSoHa3diVN+Z+>s*bh} z_7A$o@@KE{oV7V4A`+~0E^~*H197bGLt8SQ z-fdWOqV15lHHTYUFltV?9a09@oQPEeei?>LBmbQbhErpaf>*>t)h3}X@BWvFAAp=M zi~=Bh)I3EYSx8RN5j$gWRQY8o{*lcV@-Fv~6o&p4q|KlJ*1Xu`z`Z{0pCHT9VJ=nJ zdhxG>+Y(|gs*RDJ*w-(#$0uzXJX<2q)<9D5pvNkIo-Irr)*QPiUC&m`WrBpR;1|!* zy~WGmGAj#_l_)VG3)H5bxe%aQLK{Z9QiM}U$3<>EX9_-((}1Ow!BR5pBP)kkTJxEj zDmZ^t$KT3gqP#;LNwS=d49ZWXQ1yz7nFxlJU%j5Aj2@wz$W=B z<7@~OfI##~xGsy58y0I`@JYHjeM5l~em(gWdwe=T-fVp3E5?yLoFzS!3A&n6%&i%$ zCofg5(lNXTGi#gIb4xXXta!Ur!!xJwxg!GQn`68P6Z2$2sWAv|sh zLsZI%2;#YEC=OCX!Uhp6TmH%hK}ivW(Tx0wAQ4sGfP%8WLzN;K+`^2V{E46pY%FL% z!78Ajt|SV|uL$CECJlyI9`@JJJMC2<-IvFsaSJs4}ia72Dl5YH-09R-CgON4UY zK+05>lrezYE974WWq_k@a0aX343cytyOPK+4$1&Q-5`TCA%mhIS?+>js7wv)w7hi= z6E!%Ids|Fxpr_@ud!1xYWU>s|;#xkGTjI!V-vAO#bmq4Bq-M_mFLPVRNV8D; z+s-8*)CDOg!529RlJO0va)Zesk!Yx66iVj>33^Z>@d-#E=zxO*{n^o&OZa&dhB6eh zs-4QM<)dX2Z9=v32#>Unit#jXA#2%ap@V}kkzn1hp%b!(5=V6tejti>*z>yy@2l~F z@ky9^ctR)PR0^iNENPKkLn~Yrx8$aOGL%S+!>5lf7q%(}GS{#Y5L@Vod_=_(1@1wG zs{E&~tnA^FAPon717xXC7po~iI7>+&K%(s;vMoXQ9EPxEVZ-U)z(~~w_)wv&Ku;h3 zjsYJCG_ox^_}oK%JQ2J$WPk<(^>SoDJ`50kCxv|ZAkr!*uMH_8=}6`=37|wi#0Z~5 zYB+pc)U%%wAov6JE(n^72|E3olXNS)m!Vcfk5mZ+BSCSLZ8Jv1+Y7=|UJP>f05~+^ zFPp>3+-9({C7|>%<%IX4QKoR$lW;l+?-PU~6j)w3B;{gvV4sR~c-7)|1$gSqn-jfdk0)XUS3JS1*Y&G}Y3}s_RNs zwNt1n<(}?REgeH?kqUK8Ac^v>V+JtMBUK$U!0DKTTFbA3p*~aH8ov6cl=VW{NTRSv zs@jGxIITC@hEIZ^j+A@^aI$IwQEIP36DZ(r_-fF@wBkgX`@06PEvU30>j-}mlWaiIBG(VS!5u3|L!BQe zz_p+>k|)NAnohzo<=1mWjdHo3qxhuhhQp@?hn@lL79Dk6O|`Ai+M!laDAhtblvus% zhqJ095u{P55^iH{sJDbEroiS&)mCpvK-cck)sp>5>)p}S#LNml6g~{8)4@m+1$LY7j4! z1;Q=l=vDDStfC{YJ-|yka$$+9;EiDx8CJt!(EwhPPK6AWCO&KdN@MLCS%OhTn>OU$ z&XS4p@ltN$BJ7}{^Ut1&Rt7W0&Js2^NsGTQcf|FV_KDd2X1jEX&LdbdX;yTdcQMr~2mDq2Frs{-pQC@*X z{ivT3oq8|Q7c$^61^sbqPf1A>NE(Vpl^9cg)lzoK5JtiJhkvLgZEf&q5;jdP+B25J z_K&)wV;~Q3RcWlr*?lNg{}#1t2d=;+WaOh9uZ!Txj1J4)|MN1p;2Sa$ZQ!+aCOIq02W z7+4rXU~?)^5YXq=0K=f}}3Igq2OmSjdoU)_8dp4uVABB!HH3 zpr$J9OC1rG6}q@2K>?1ulQa_90T7MWx_c{$`Udiu9UwLY_^PZ1Cjw}&FjTA!oW>=a zAS%n96G2#MP**$utFR}lPJmk8@t2INLNL&#kPXCtP(XqHM9pVwas5^)dI>l<8fjH z)$V|16}ZEmGEvAVED)7%RZub60^ z1hdE3jIt=f0V!aqmOBJ^O!TwSa$&`c_VXRh)gtM*TbU8)2-bP5!L=17mT^q0bgEve zyXJ^YAk!gF1upb(J(n2ADD%-FT)MxmlcTw#M=-&PLli?&{@@#m%baKtk!++m3VA_` zD>x6s1l&N<4F|eq3f_jk4zBr$>YKzDDd?JG?i0dh!!hzAgdzNa7QX9i`W_nenT)s z)ELK!!C{6ncnl7DQTW}GpcpuI^hSR5QW`yNO)g$s5^HS-Otm<&7NUmKEYvhXt1TFzOPo< z?hJ;H3&i>xnCs9BA-L2)-HEq6;motxJNo-gP12(72*t0e(GX+^q0faz1w?!v8FGBs za7h$S)JbejI29A}d$_n#GIA5s$g{1=t?-Qsy-c^^c5V3&{WUxT%T?2kbyNSHx106M6zCnnLi=Y6W}@_FM@tAOn4`**94&-#;qq;Cz==KPfMp| zG(n0^DlJNCR2$GJA4dnns2`&(5NNSMJ`DWpB*rHc0@N$S2!Yg%yKevqV7aB>-+wOU zR~ynPI=f7sY&@ZK;u$C$(J0GXbN^{uO@UqVp%j94-xQQfg7{bSi8@_L+Kr5!as!9_ zcM^3lsB>j=9QAzJ6a45l9IAOhfLV}>wNNs}sXU$^K;Vdm68RwDxr4dUhnrc#>46Mu z1~agsM7|VHpl3fSf)b#t)K?#_#}DyJ9}lROIq2p}@@nCtK=!KH>go1Ev_ zs!Lfw=NdI6L@6~HERiSF5xheq?sia*Vy^10xTFz>5+KXr72$fzK<@0{PCpOrN18#+tFp?lHkHjgVK*&Hs ztYOy^R{HQi3kpjaY*0nbi4lt=u5g&ZuaLQIPlO^^Z<1eiSRhY0u>##LKW{FSHVFyC zc&5;Iz!LxPnshcp;pAKrdTX9{!(~Jyhy3?=H&donbZE2=D~L;WGMT6sYU9F+gY8D^P~E+=w9X&ofsb)$vQB=+&+L;N!|?v0u8*I za0aTxG1eeEBdz-OYR15oX-y%6rR;<{g6G}pWvxSXS6nh=Bk+%9uB?G@y=8!R`=`qq zsws%jd4#EB0y&hqIE_#a8|b&cq2CCnuz`O2cj-4mfsg_Ets$oZqL{M(N^m2P5QeUU z+@pHEVc%-J^t`HezkxnOX#uDk&;ydevNbXPDY1Kq5UugY%r!(_R`|>WB#7bq6L`E3 zKC!%Vofec6t~dUqGe|s1ZMRrMu1BuuR-@l#N5UflfFzimhYiI@*Nl83B%9LNrEPCLR#36&OKBuHqlI z%w!c=M-Qbz{88&hBEdrq137SbfiRPGkPEJY?yyK5HarC{Fp}xd;Fh8icx*O{#|$rh z(T^P##_)sf%!N-+=Wy83rEd#hv(N_e=%`4Ucd1WDF1Y&9S@^4?nf|=Y@t#Y9A93sIkbO~^l=RJp2MP@Worbu{$BBP86lSJ7W ztUCzX9>AzUsy)F8mm4f9vaT2_EL1D9jxQ+8!R6zCAONLiftM`JnNdLL;XRUF@oP|6 z=sHv~8R-X)+n`D_Cs_hVoVdHEnW#%OB|z{8-X}yak}mkY@*PGI=h%`S$^>7XVs6b~ zMKO3xKRWmbbpC3!3B4{1El_+Wls81>ky)2*dBZ1u9O&E-@YsanX5#M;1k07dLX1Wk z`EmJ>G--$lt^U9gJ}=|S4&XV^CkW>e(_Z)=HbZ)qfdxoZ$TlqCa|R7whh%lE0puLZ zvjXz5+(ZOD&ew3|(SeP1VQ_=!;nMIkD%l7_vTIyD@-s*}Fl%3)h{1j1vW+NsUK?jR z>I1z5+Q6_QIeusyOzh$dj47TFAdV3++&3ngfu03+&EYjE)l9?K35q5h9>!8NUo(vh z^WY=EA*)<7Eu4CanF)tONA^F!4-7}P=x`B1t;}*^7Z720BsTa^U_=0C3??IhOI<6U z52xqvK3_ne2`8d5Us`Z`=me3Nmj52~ZUOO8(Ve(%dxY9Vls}OOLH-q7sULeWB%Ekb zgA%BdwT~;fA>|8KbfH7=cX|&outMAtP0!zbzMSDoji}ekxz{9 zUy`0LngR-JGWep_RTdC+DE-g`i^ben`t^;>cuJsRfrYh7wK|-i} zf=jhm~JH4*eqK`f>iCpcOrjjLQfN1+j~T_!2mv*dL@w*wKIv^JCGu zT&5o=J;NNS?H9?3VjwlMNHv_eBu02-F)Y~C2d{C#Eb>s|4*6XiBWY9Nubz-KEaxEP zs_Rize9+43UL}YNbS)}mA<2S7kNhih>3wTd7@|R?s)cwArWsNAs5yEOdH}VWH8p@) zO!V=oFpFBRvk%A#A4nFn6hsLta-_B`%o$VRDIAm*2(}8K6}Fg?KOSQOavcMIrO_yi z9R}(lOaquK7A(#|FIB$YH#&&PW4MOX1(g-ipvM>Zz|y!{xr|6utriAL$f{_s4#q8p zWIwiY0~V4&5UBeybP`dyj+Euku0j6(wq<~X{OLmLB7gEnRMsJ7`5VsaEBTh=>)@;LV${48-ZYM)x znFg5<2Fw!wN>oEl05$>WkRFr)2v+cSxih6m9E&VJMka6=Q4yV1I?^kPBddE;D)00Z z4bZ%>;06jFMEcsZl0>6I!wE-~UqeZhw3KorGcuAF9?5G6XZ|*vA%QcOg;)iZfhJCa z$R7~nlXntZ`XsEh7Tb& zPZGE4p2~x1JBg-CXZImU zl{KXNI$K0(iw2>O0pS1-aB6nb1vy5^}k9JIqkocu=^_*_YYfv>EA zC&*BUgcGYD5Hg{($bes0^C^l3125kCk6qx(_!vS~k6_!sNSsh>OC=eA(ymgKbhVFm z1cSE&py+YY)@Ge;m4-br>Wd7FJF=s zP2Qb8-EYmjxr^D07R^f?7n`?3d+6l;CtYTq8*i+j2>#mlUNXh9=%rtiW^Fnr0y6jhpO6@2)-Fy4Yzg^FNDo^ww%QJ7(D?sULVW zc&J`q?FTDcFLiFS!?Eje4WoS{8OMtoEe-rKVoyjOrNn9RCzW|~7hcaCsoS)%hC*YF z`;!}gj?8i!ZW z(v#KqO!(9M8=5q7l$Ss6e`yOiUFUW1)1;wKK7B8YoT5gX*;~(W`NU|4 zx8JW^c$BF3#c;^hiI4dquKQ-IP$sl}-cs;^3W{Dz#rLukX-7X!d@=p)hc@bmqkh=A z+vNY;y7-FoyS?ojDJ7e{Kakc+bIWJ@*o-5)il094o8m=!ddMd=JzfPtVJtKH5KT}iK{Mt7`!toC1dh~JoXxw&yPaRYVOz-?~AM-_&a_;CqEaRn_Dk`J*PMMHp@BH&-Ff;Z&B+K0QX?FV-caM@2ba+-&&P z?s_^?w~1g(g~q9EJih*~jpmQ&XDi(k-NzjK)V~`k)>ykUc;}|^MfVnS9{H;OHu|t* z=ixEmkL&ei-WZ(JNHC^CBfHk$ZkJ@WOaA?Nlv!x(=XQDBHjD+&JF^*hUX25kl4~3I ztN5+Xj+?qE5G?=U;47PVJ0x^9G&LOlYtD!Lv5rqHEN)zx_M^|s7Hc-WfBsxMVc9xI z12fAj8=U>N1@5t5bjzt7A}X0Z->u*8vu8^5uY%M=C5JAp=$6v0AamZRp}iI8kH++% z{BVl0A9E*ly+et^%GM*gg?xWz6*?wsP4<*;!#`U)q?|rIVcz_O8|TbXa%{0}cZcau zp9btm+qreej?M!IrER9K9Ow8b)&BNQ!%SslMs_rx{N&`uvtAnck5W zZB^zg#Ayv?X(hh7oU*R<-noWn$9%nXnsj@?eqO+9e@2ruk9Qh#7k*Ayx1zDlqdA8? z-B%xJ*KXxiyK4HaXZ7T_XRJ8NRXZH?T_eN@-Na@i? zsY|oMX4>}o5f9Eg$L={8G4IJ#?(oR3D(Xou2lU?5@7x0Cn=KSuDBD>tSg`SM`yu;s z!gZR0m0#Ir0{8Wf4@2}~z1}?^{bc5tMW(l3nDKK?|9E#LbLDl)j|B(6ADgW;mD<0s z>D}X8!_RM;cq&uwtk@a#$)h<#wd;b{&X+8j{@1usw7SRVZ@o1&ty&EK-;b2%6Qk9y zOueO+t<$k+<=xjSw(+<1dpdmoj>G5tUVro{Zaj1A+tn5Z27!KEbQgGdXSFEoo}2e) zUBHNgiXD1+P^~+5P&zPY+U-JnW%ha(&(YkM_s$f#r-fg=p_)7}PRslT<#NigtLwJC z4T}vj`Jd-A#%zzmZx7$@_ObiEeyaV^^mf6+Ca;||Y?OzoL2t|J?rqpN4))V+g{;oW z=3Re==^cG{{zyqv_M@0fEw637Pj%|mwy|2aQ&8)x|8@BME$GT(t5!!!;*$qY4O{52 zaP3&T1q(Nvj{Bw-IPc-8eopRV&Db$%Bdmzwmn%rHymRd=PaE|dGj*%_lAu=yeoIrz>CL$?WQ z6K1DxkBccfu{q|&QA*(56<=4SX=lZyZ&yiom~X;ob=FtuqS)2({`Vt4?B*}-=KXPr4Bf4gII?-p)8fT?ak=^L&lDSx^e|#I6{O9wR!yls) z7DYe$nYD%zr=^seA82*+-f@2m*8YUSQbSNr>boMH@wnxtrTf~u zjQQ~5@CMymH?OQ{r_B7*j~$eKVCwH;%a=*_W=DmNaqvm=X+P@l=*MXr7TkG|`F5si zW=zaDkI+@_&2s|U$9w{qv|+^w*ROZRP5E@C)t~*vYDd$uTV9Opqdl(G_IMFXyUgF*eYi=2-R7m2{>usI6Z#2Q zNHy*0lh_A#@7=q5fM34zz*hTJ_urnG-o&F3C2zx^iR*8^eUd*?d++P-R2mfwf6;h* z>y8Z10fy%?!i#r#`&#xWx;1dpUP|zp?>?Dp=P$pM_){ghkL$7xHnG+@iH7u~w?~pb zv2SJbZY|34Tyy;HN|K)^)#ge<;GKkP{Z<-mI_{*wvTS2D6#(A@wFwh1@0z!%S<5lU zKk583TvM{mOp&gf-Xb8-;vU8S6XWBQh)e$qiyGHIpn3A-5Z7gW=Wfb8vx@rEYyGeR zJ}ylc{#w(#OQE`IpOMqatYuehIGWA*>D#~lAp3q$R(WhX%Gvm~I*Bwf=%~ zWSR-zR5Shq7&t)f+qw1Y3=clK-m1f}tUdIloqic^9I50zt^d6d>4ld@{|>%(dj8a* zw{Atyy!byu=f)3Yspz*EXtQotFBYeHB0G7@2Tk+9o0D*gc z6h87cO$F5{Ae&eyc6KDG{4)crhR|I-p9$(bx zw{gy+k-rlPmv0Q6*S0sNG&kYjSY~PfEF-0P`Z>3w zM?P&2RgUi7BlxL)@g2Q5E#psbE-V@`jQe}}YUb+L;;ufY^3q8jYwzBgd9+7yV}%#? z9$uH7a`&4*nUSi(3hdC$+I#&)F0Z-Kq#xOlPxAG8cTkDTwJvzBtZ1w;rtul87S9jP zeq)f)|M=4d;f3LojKI9Kdw2NvSjWrW$A*$N?#x=8p0>++%iOzz`k%aB9NpJ4kiXev z+tG;K9~6%NY~OjJKc9bdJ7=T+ueE~*Z%CN(tl5#YX35L@a9!f_*7K4g9*tE>EEY}sfK6s|^V)})PVqejiz1ce2s{9v!S@~wFK)1YP}U2ktU z;s3m!+xtVp`T@FHzYDvpc((6Quf+B%_nsROW8HRc>cG>gIjOIno%?jQCo6Tf!hgWy zbh5es=Fd^fF(osfJ9~CL(^hftkUjwx#nb1zwCM5kQ4()oNvkbICX@EPIMC6ik9y?$ z33JUiI^6ltchQ!UGr4BWmW)tyg`9v*S7U}>-GNXM{Zc!YF}LqwV{oayr2K?F;#;_>yUl3ol#H( zv-8{;SHcFl9SUw6$p1E6W7Hm1VAJ9>Shwaiv-HzytVTOF!I;m|X5~+6s&;teIjswf z1GWio_l>_)phgRsrJ56vHetN(tzN)yo!=}j_i4^8PLFJk?f4E|zQyEIPxYPpX*tI$ z;Lh6xZ~anzLY0!Iuud^M<+N`}y%C`dJVc+(J&{Kl{~$Ugef6&spU3}i;p83ws$SoZ zHQ~>om^S~aJ}v$IneI7F>5sfkhTdF3W?f_@ecIgU`$wz9ye|p&T!$Y{FnZ-0JlZqf z=>7QX%MPAi7dL)uhqzp?t25F)2X*hbFiXjp|LtKT`poqAX&P&TcebN#*=2KTY4>iL zQ}vJhG}v(MqitTv`#-!Ag^6oZw|eeVqvbjUUsH>7XtN;b^hOXUO`bk0-Gu*b)Dp)m zL&H;jjo6vlDXT0mds#v>Puap$#?Xx%Au3S^l*I%9scFgE0nKZ!MS}=Wub&-mqUr zUDy2@Q$S-BZqI7BWw?!@r`Hjwxd(n+>B61oBMpH{QRtLuhwlbS^nl-=$4F?YuRlY)2}ld#pe%> z)3OiW_Ou1^K{}!j`mH^5Nw#r*CtaJ282}jL`|asBw%f}gep|1*@j~0@qz<~bq|5B| zSqoI!jRA^vwgU@(Cquq_%7?V8#``}niSNs}%a7XT@@exxHt(ICM&BEE% zyf;Ug94MBagC9fpe6hT8B{rXDJM33vfrZ)&W&hj% zYt_;cMC{-R!%IKmRBtfKHSY{o%IscEhB!D3{J1vc6fhX~b-^qZ1y4&i7{h<+m}H!v z@7cC>>y!IkoC-OLNy*8~uWx>3hTq+}qDyjN(>{3_Jr6(4e(rZH7);=eWgmQ6<~$9j zPkQ;LQ$oisCUcv=pFF5ZmX_6q9f#s~yi9tr`^=@mdNjS3a~S!*-B#P~@3VJb$obs| zBUHy+APg;i>_S%`oEMlYxeEcy7<*|uUlw#x1{xd zjB=-Nntv`%WNzQvq>GC8@>1b__&|?#N7(AqONaNY{kQCM;=}$^1COSqiagpv2lk^K zPco-Ii+Xui*^_o`@%+f(9!8@~J^yQV{lI|hX_Wr$decK9Ud?{%y7)#PTl)oD5AZww zv@&g?w{KqhpN=k# zg|+ZQtnXl_!|Ca9r_XovUOU#%)6-;7`(f6;enwliZZW<#IDL_6x5mr&r#$4Sr}zdH z9Ql5OI>6vYU#ELd5Sda|$a$(xKQ+YuSgU#Mf(yTQ===NB_*v~X=XNaq)}v`?z@kwT z5Rq3h<&Wvbx~98V^UT-jq(j+lvNxw|vOn#(dVIv;unVt7HZ|f|B2PCQJY5?aZEV+) zr$4TA^elWc`&h))`Ac`6IdH6a^40kf30IacQvm4XV9X!WpEcR8X_kqxQKLJfER6;o zPCOmTC~i;scupDElSZh_c2OMKb6OWY!|0g%XLBzvOdCf5SiV#U%4)$gjr%FS3(|e2 z-D+6p+=F{>h6Je5hG(eJfMm8X4OM#{r|@G&_hI@S7Id{8|KrQp%x{T$-G&<4+l2r@ zAD$w*Z(91nC{^%8T}smXwxla7dwXjN@n$0p$jg|+O*L4%RbxHc07L>-uf0m)|Bm$u zUj6dz*Sp3;!|z=lfA?XmT6o(=thqx6uIbkIKRP#Ra@M+TaSIpsC$r8Yrsqb}M*fnY z(aAh_% zLdn|$d0V_Hcz@ix;Yy1Arm1*`TSjR&qVMB=TF}D~3~TVO4*MUO)#>1=pvSvrwOx5< zU)uh&hdLv6^7Ws`s$>05lpOn(J#6T^7v7(4mN4j_BOjdZlwkNK;Dfp4_q5}pSar6` zfLHx)UwOym#=2*YZ$?+1HTu3JL z8%C5%qerh^y?no>!z{*;mcO69Zs~ax#D33E|6|vll(F6~{?}t{CpUgV zWa1;$zO3JUT(!3szvZi+vAQ)F8Sw*lMD5kB9X zHJ)b|-`t+oKQ!8T24^^k#o8bLkotPWe`+*8QziPNM=Ge(f;{>{YPaGQ9e?juPYhiA z*n8$Rr!86|zp$r%wQ>!cYx{MoU+hiy6_E#?Ps`erFgVLQW!LWGJ4Zh8*iyW<&-mZj z{!emJj}#^A_E1XxbEXBI@-##x=j^#o;_sl`C1hOgLayOq#w?eD$9g*weO6zH+r`Vf z{aTy1gne>g-gX<;m6trX1@a#QVnw^hs|TU$ohDvhl3h44U1j6sVfK$vTBTTH!7i;Q zlfO*;G*o-&yM5z&vd&Ik2;6afp&zM~WpYA+&K>0cq*M0gH$?$Wdv6`~0(eIMdpSM* zeMCGI0zm!mrz+5|tA-wq2cHlgPxzu%TyDqA*Au}zO+?}YygLBA-YFvRPfJr5sWq?( zW5ML!jvu~jE}*Ev*H?}Y4e(LVI&xg?vbAFGMPJuiPm48a)qTyY)7$Be=8E)FWTjbv zcW>*}?NYA>a2K1h+zXwum-$%*o!i>O`FD2S_9NTobnfifQ;yv%}KHu#HEW?2#{!B{x^;q4` zXU??Gf!i6I^2Zlv{J!DXF(-9JTOR|ZEi-u>Z}=DiWVJ`o;TL z<6UD49Y?=-Fpx25PoGG?eUIEPz6`Kbq`TV9+tU?HWRJMq2RZRS#v^x*D-8R3GAC9c z?N~~|;=DIs8@bta!F;437T%b>U9q??OQrp~O!~`Z&9`$8k6@a;_D>tgYS!h@l|3yo z`fH}l9ruPiZCS6R8;%#2JprCEeYz;xIl8JGjx5T9Mv<|M(%z6GJSf+$Kb&q``WC_P+dcJZRdEd{R}YC zT@N3)t3y&s{_eaY@b>1b8kYj5j4|zGqC@|&dFa)CqXW&U-;90^S-HPxYq!slWAnS( z-0tvc(}$rWr`$UoJOW&h5j(~}=fRIhd`IbqRkqxseEjGiEXf3}CyEKe;bU`lvw ze%3{|;I?xgHdz|%v0~v824w{KrPYZ+ESJ9a%fP@>zwp)FkN5qR7$;bEr)dQuPjYg} z^p6W8shNy3rUid4bV>UB{bQh>%k59YGeaH)mjwUHNG~|ri`@jINVdlIU~R~~^f4#+ z)Gk{!+P?GqKV8ZQnfaCBx@+s|lsVgu7|fe!uIt=;?&<8j>>rmiVsg9nnb^u!V^zA% zYVE0oi7QW!zMi?+<)(KE;I&yxE^0Xs;PxBwtUZt1UK@$zTMDhu&l|_WAF5pJ7EOui zmA@zb%Mi9}@0DW)03m#x0$9fFF)CL#3~fAUkjsl!45n_stW{yQ149q{4KMud<+jZ6 zxSdw&_6u8Hb~YPuT}cr{YdUG#LzVilitlvM9;y@CoS)(cq92Frvj&Q@C3}ADT0h3@ z)YjZBz;|u&i95>i_qsk)KKC26{ovl|>etD!q3zxjnEtt({Pf3R7e4>^;lKdr3+wi5 zvl_K%Waq@yMhBg4jY@Eed_j5{^Jdm-GRwtqSF4y`cB#*t|6Dq||JQ<}c~9H5ZkM)w z&JNWUWB?qUTpe?VO!2dvG0wE-23nZy#%0n zEb85{8^0g6x~Fd2(MBPH5>LxE|BkvH*3RBFoEmYU?+5}g*m&7=i+A&jjQO3TW*uNE z?OUD-U}y~@#>d*VHReC|PF|*KW3ff+>@L#nlE*gRmyv^tpV{;ID}pp)S|rey$2l$= zWOE|k;b_sy*2{gyYjnscn0Ia3v$v1c2j09|Vn1%%$!Y1cZ0(Cad^@veh|f_kt2>{1 z6&@Q4kUUTr?hXk40g@{2e2?Y!+xo<0%uBzZ>2@r1N8jXa1N=Sx)M!V~$KTeT@_zQY z++wde$t1+rMnP{$N)+O^k_Hr>`F%Ni*mX|RA1&WCQ|~_Y!`3_NRH&y?@2D@da|}7$ zvDcLq+ti&j8r}R5r5QZ4F+lW-t2p1zuvG%nCYjvH?ECV#k@$JOZ@-%#YUJ^Ld+qzA5&+`lSrwz(G#%!6nylGFC<+YagI!t-rEYSaq zI(R~7p(B_-hR=`P&^Kz2(zgxFFO514(_hI>jC1yWy^x#O?Aw$5SFZFvb#}L7qO#_z zxBm1-TRs_%xF2f7|D?HkIeGH$&DR%nc?&p2k2OO_og6Zy+k=;U-G`~Y$9P@`LCs|~ zTC#E>K1tD~D8oE73otB*EVn%VZQ+AmkwfMr-!lKTznQ=p-*+%QZMXH7GmPT*BfG0x ze7NgARvQqQz?R+&0X$^8oF%;I@jK1G4_D6K1rikuzE045AcCFRs0kz{Id@jtli{o<2`98^ z?>=%|;&Y2sW6X|zUH!>=yVI-pvH$zDQJO$3dttC)7f8KL zo3>F|Exi*mP$RJe0S4M)@0HFCT*kL)wKyp70})) z+wSJCcz)hf`Ni{<$w41K9LZ)S4QtbkPSG%a-dK_42tp})1>jnv5ae(CdOS!W2HtqM z%ys)8#-rxmDNPldjGZ=Zrr?A7f)CyoOn-GWb$qAR4l5(~|9t*?|qtjkU?uEt?%JsFElz1+Ir`fSwZ{HdOpmZZ98;R^j4kF zL(zCz>+xUr{4`^B&VAu&uBVx}v=3{@sMKpI*K|w9F7{1I8q>m~bDMA=+T*y7yY_mU z+TMQMQmcM}Z=RUknv}b=ZLhcI%oCpOF>=r8+osnO58c*%l)8+ZW4Py4Te{yqH}0b` zV9euLaVtzt3{c#0w0q|xLy{j38#k~09NXj+;1DkDJa^mWE+&OP-u8X*ev0GjeeIq$ zqQ4r}qChFR*MEK69AM8GwrR2B<-sdg2LxYQ64z>x-@1l`Za5UlRr@@(*bXio#qqDTQxexaCn)^p1S>o)I>sGsOJdH2UZ?i+@k|M5So$=)e5Mt>V|wfFB?b6hkZeC!>f^@!On z$1!TiuR|+#XgU8`dyx`!?K{op%Fv?Q?g2L*|LJ@3qsgFY53lDRzV)Tk_#T~#VzRQE zUr+>PK;L$0oJuW>UaZ+UKHamIO2;^d=7R^?urFlH?LM#V5nbx`|3&Xs9TKb5u|MmW z%1Y~;-=o%#&!?>K!F#HlareZOw9osaqEhWw?D8d_Cr=Fhr86TmpJH>xaed67_D6Ci zTEZ&{I^sU$9&C@)Gyk`e%cYYrJz_hTBV<2DAR zddk^x_IDiF+aIRPoVH$DDS5chuusnm$Snue_%4IVWfI4df@#6Slmc zyj*?mo0sAgP;yOCt#8gQdPapc2|`a5pE9R8X-;ac|2k?p-Ne`(g1HtJ8y zrKe{lwMZsg@F!kzJoG6ua@oT?FHPH+%L|Scg*muG{Z(3`9vcjc|+yVdFY|+UooU!fl#>}Mj=6#ODl=Ny+)G6jEt=Qh{ zNW#)rOVXDyp3fPUcaiuK>etSR(#OpDszs{`ED$Lp%t$b?4si)7nvs{jY`jX`d z?3E?E+iI$(Uiz}6Fz@c-^sGwe&98Mr%3o3pSW+I_s1-ZuN>)M`q1(h*>)0 zxJuHRjUQr*vQvu(I{U0}(1=*|r}wzDXTL@&Cy|Fg0F>5+#hF>JA7vjYda$7IMeMyV zGfOPnGB=(GJ7ba)_E4-VS^(ZPO*|o!bo8 zaGvZwrT3KM&%PRtGuJSj9aYlHX2rA*+^eZMYYWycExb078uCr=GPPgk(=RU*3Mc+= zt?u-Uef!=$34X)5wk7(@mLrqCPVqYM#(ivo)h3PbjEoth22WUW`s_s8k#;A)&QV>d z_+wrFf!h6g4vbIDO&&RR`7(#(S@&`R3QnuM^tc&nq|$l*mO~Nq(v3hVr!7*5s2FvV zG|sUXT8GH{e;?l+Pxbr$Z#$B$i0qY}k-a%alDv&@%tXdHM%iSPO(G*j2%*B;>R1_3 zl68!5>?nI>@8f&F=<|7h-{0Tk@#}FO{o(bxU-xxi_qFcpdS3U?ot^J8A5~NW2$m1k z&vG0i((xB5c+PShac?IXK!}q75D~v0#FW<81qsgtUJA^*pgZ6& zdU=0;vVWQHx#_G**&96Ds}@AR>pz8K%2T><6aTA3M|EO2dzfHah8{jKM03)`>Btkb zb-H}*a~6K_;stR0#U!wF@~|g*H1`bll&+Y(LxopdF_{i6bJWMXXZ_)xbIQdog+o9H zf*48`F_-L2A7-1rMyHhQloL?JsVWmx*IFo}9jcJC12bUR^UMOs;$3+T?=`vx+pLoD z1MY|_BHAEKW^=Z2-OZYLm2ZT*hJDWAjQ#P(YXx^ZmwRK)h}=&!E83X*|2ErotE?zB z`(%Bd4y{6EF=LjAV@?0kpZT%GF#JgujgQcfbISZA7-(8@(%Jj?1&E&<`i2_aUA;%u4e3x7#~C{rx_NJfT>-ykAvjI8Aoj813owA~DsxlNbWap+5Hx zE@tdxa>6hDat7_a7uwSToYtc(t@P$Ha;e_srR(LY4*@3W9w^1i$~U;uNfqqH3fOK+ z3=a;Bao)~o&iBO?{=JY&`nsL9d}L#dt51EdpB>qaQ~R_G0MMNP=q&NHo<^}w2b=Lr zR5e_WSBJ3h=VnEqe&s##8@SJhxk(ZoDf(}Bxvb|Z73r;VmWpe#k1VJk768?OXDodIH&*=1NA9GfuBm5!e z`%A8Z8Qq$Zk?#5r_E^*C|2TrEE5G5S_=h}iNu0Qw#(Xx3=yJJH9pvaIh>>>v!;!fKB&1BO7_wMOE58*Yrpit%VnU7GW>(a zG+ZOORU8S$OeoPLYAG+a+sOm~4joheJ6FayVPl%R6hWe zG<(!7fdI$@0T`C`|MYcL%9i{lcO|XEJ~;rTWw0Y(3)>`<;8VL(g4p)Kodq{uApd)6 zSY8mDsTZOE76`4-%LE922tvnYy7%sy0Hho_vLIJNC#7U@FceX`fMHYw>@`gS2acB> zz0O9|QpMU#qw!9NoYtEXTN7Xk33Uoc2%!z&iVt-F_0M~O97z0W65EQoN3DgRjbRU} z*<3QOA%>-=)0y)p$)rtN4Qg{Em3Kuz z?&bPFK-pxfw-&QUWa$cab13I!c5B>(X!l7op9_F2wX|<7u($-TaJ3A|ZchaB7hceg zqvpTU`?ggiwuloPKks5Y_Zn2P0J4#n*3>fbdTd9Ujqe#|u8Xpn_nV=eB?Zc?EMX^j zLsLHIx>?dhGLb$%64aqgtR~QnyDz@89IZ7**@ML7Cx%LtI5JrFK3{Yfx%Te__83-! z-$Lj|)vC}#W&11mpv^~y;PH+_zlYc}ZSKYCb=qjP^y(SC%es$A=gj@Pabe5{F8&))?w^unz@-aD}T~j?pkMMIo0tD{-w3*TUK(J-P+p8yqEQORN*|jDBWT& z?*L1?68H@Z$6|&28%=NBWWDZ?MV0tjj*9HjFc1|9Z7eVAq84I_R_K^G>Jw1T>gbUF z7nN5s6aIw4?AcXb^FG}XS{87?noNDW!DHW5-fA(*i+q;wZ8ULFj7(u`_^_dP;2
W7WywuELuN*y>Qo{7%;F47E zCE*ziH%Ag$BZa(aGBrO?*T~8$^P8d7uO@MK#BC*T&^8WiUgkvBpWL6nNgGbg8NE!O zzJPd5Ykn8t4-D7E)MHBc6Cp2oHJOcpS|^JO~8O-gG|r@33`e)_|+o1Vyq40&#NMyd+Vsevg`3|A7=zD-QIWalvfgu?xyX(*iNE=y3=-8HiyFpjV5}=JiM6jkXtYpE>4tIqQiYQ5 z9HLA{06pq11dAIkvELo@VL|b^s|uCooV50K+De z4h-AydU)xuM8m($5-i=-QMmVNSQo7=%#Rv8L$t+`T!5UL9^~AB6?M?B!*ln5ZcC z9SuvFI?(zwQXil!Y_a=ps^Z=Wjj4as4;I^Jq%tr>nl9-8#Iox@-qzi_bq7hr!~0p@ zE2iYtZDTJAP*Ll*jD3Eo8`Rz-FN1{ucqyp8L)PfOT|46ba;~nfSv_UVwC$I~UY5M{ z&7Hae(y?u4>DaIlu)uh(Tlbe?ivNi|b0}j(Mk)Vv4u$#ev`Yp;aHL;g3(s&oc=%eh z&;ej&vgVclu3BmrKfaQGQCK+%Wxj`De!4gS9kF?zWnuyFC?v#=th55DW~gXyNA;%$#^A>bCCpqg7|pz~W;`nlFY-mb;J!OLiGcmAL2LoLTxiCIH1K&o4$waw1Y8(pJ-LEb8GzDc3B)({?=Ol&9(hWmyfzZ)E?5a zS04kA7T#)iQQLu*b>JIaBTvsXM+1tp#?E$-%At|GYmIh`8a-IcC(^G-iF8>iyCx;s zLnwW`_+8>K{SH!*hK`g9eB)$w^N!%^!`X0d2`a~z)gnDW(E&R%o08sKsw~fmbjvhM z>X^^`uJ#o(Ttrq*Zkk!WIoSSXa^Qu}z+S3*lO9GJBt;jgLK-K2HJtARzKxB_2Q7j72!a zcU}@7o6^PW2kjzw_`5W}?S}*m-}a7*?_N>WNY(plpU{XRKR(viJvqb2&OVYBIXC>ukIi_j)1yo>6v&qLBwYL9Vg*9d;s$oIem-z0mrv;3y zv8?xS%H{78zdJ+g>0ZmNy?TI-FrV2Mo=foOasEdU^0S^kTjx$cqS9nflvm+mNhIpE zE1AR&Ga6oAM?Jl%C-T)yie<1xS2t6k;b&vkSk)G(ziQ;HWiK}f1JmwpztR0m+`ry$ zvLg8U{V&&>xFLVkL<}X{YF8t1$ut8hKD$@)pT47eSizcf3$CB?roXrKivLRILl9Ao z3wq(xapmvru~ZqQGO!*FKU1$9j>2VNNx{dBHp}^Ib4&^YkmDqvHz=9T?x`$r-XS4^ zWc1Q-1M|DNZmXE&{C(UL{TFHq8vKxR>lR!G51iSk z(cNJvl`43R^1*uXPd6Ouc7X_9CqTvmo&!`*e(;ubQ4sEPcUQWVFLEe4V0*2rMo|TH zCarWzIB8V6ajHe4z=`0>vVuqT_-+~{TYlIgJ+L^L$QUII6V^t~ zWw2RqjYcx$Pm$DiA5+^n%+jQ794hzJRw-{$*z}!PTz}7sd$qPU;I%Ee>f^Vy5bqXv zGEut6QG4*UrJAJnXWiODIb!V%$MCV!PqvLK92fQ5rU(`j+5y~~@SAJe;nzuwu+3C?at8lj9hnSVahzrFHpzZkT!efAx^B$=(NHkJx>m8W~Z17Cs&-;S&; zTv#NZ{Ne9rI?4Hu6;A8cIMvP6daM*9e~XEci7HpY7(6oey8OqlLB9us(S3V$Y)%1~ zETZd|(E8|=trSHr#T2Px|CqjjA2gAznv+cRs@Sb3EkMKH?HGBj{N{;#%VQaT8w=cw zK#!ukL3`?Z-U%7U`S}hmV|pzpa!DcLwzV65jdbz+6qi1~dU8k%SYK|(@M=J?t4g>% zY=uvq1U<@|QkxLAWNH@f)?wJ*uCu0#DBck6B`p|NQOfbd^C zZ(ccj;ef}N?-L##`4V(5z_)CRtjaPWrbGDvvBXWE2tBzuG-_8V%d?Du4!!;ySGXdS z5qSAoDUXkCUJRYFQd$2rGGnABC*vCM!&|W(KYhl`4a@_#q%kacL)%MV6ou8XMy>v< zO|FxcKAWgx(Q4m3E&B`fz`(%NRA|xHo{s6g$L))-D45)Mz# zL;rLq#T z?{58`8clsiX)PyNRMT2w^~m71H3OU4HG0qm3R6)LK^8@MehTpvc|w*^$@4Pt<=86v z^1!tZ4aNN%KlzR`)yXUE+ePF|<%225_kM8x)y_&VqXdGs6+7Gl;&0>Njn)V^wqx3R zLzQ`;auqt;S>b4Eq_I`7nR)Wyd00vHb642>EqbCO+63~^ADbB0Is&8GTPnB)Ze)Ax zxxC`8=Dry3+fmn91SF8{?=@z2gVoceV;g6x4sCL-;rbGlbdtK?@w+W? zG}=ctGS8LcqyqP>E%+CwQ^;)lrS{%u`R3)#|9(9bd;(vWL?JX?K>PDx({ar>W!&-h zFRK>O`W#)x0R!SQw+S^Ej1b>hiadE2aQLQ?UJ<>dp}b-a*QEt*LyRs{-3^{)SA21s zb$rh?^L7`QHOUJJAQQ=MLk#$Z%DH9a9))wqbUGMp#hc%cl5*)6Fi!;Of$$nu)@t^J ziMF&f*>SLW3L!bV_Cr zK2aL?{8_-x$K^p%h93&dq*_K`XYueWz|*~NjaP%Mbl^I@hkl4yw;yg4!bDInk_9sr zrOxqu(E|*o!FKSnHZz~D5h*k%7fJ_2ooo-o-;nMG)dk?teZ7Iguf4B=eruaBqmL&| zaQ@M5R$qgxrGm*9?{Z{fAR<8$A`GDRj#P@QaW}p-ys~Bu+UnTgn5ug8@~oL*G0(zE z3NG!Xv+vfpoT1cK>4ZnGce*@Pu)L{ItFPiNld}vOnF5X6!0ChyopTOX!)v~Hu9)4P zjp^Be$I~sVIsNoQZGDTVY!1`2s|@5o@ElH(YH3q?=aV|m8~E3TFH(+<%f7nmSYI-M zMP&t{mA@bTs9BWuQBH_M1Ua5(=sI7gm1#bVzFT^B^HBQE~&G{E} z%;c#kZq}b@`0XE&^*dxR1?{JJ$a#kHLT(w4ND|37*gg<@A_6d>7@#~$F>auEY4vWlhoU8GnVRGg_PESxVYYs}0D{Pcy3;=h+(5P<{A zlJ@)(Mk*Q%&o3`4vptFaIa4L7sEv5+@%iuo2;%~#x;Rq)5FYnbSKAcJBf_y=?$Nj` z=eaeLNaFpp1n>a-=m8xCp~&OC`c^7wT<)R~3gt8q-ChG393jRfkZyWG9XpE2cBy;A zjd77NR<>y?>|lzBulX3TR#AR~k*t;o5;h+?_qtLR#p*0{(3aKAj=s`7?FG8$!5bR; zxsY8RuNNqi4S=Obd%W->q0qjeQ&YTdbR0I1#irq_YevbcNBXS!;P9YJOFZt*gTo|< zDpZ?L&4b^u+N2G~5e>sT`_V_D=v(TRmJdB|bVelE&vZP$obNV9;}y^i$_|Ou;~u;e zPxa@eRJUNa2Qm~Alb&UIOveS?BfocddoA_ZgzkK1-(%W3J_BSR0UWM{xci-v1QK^_ zdI(*3aoOw8Mdx>K7kznHL`8G^n(Zd+NE?;SN`tdZt&Fh34K7vV+%axLCag61d<*g{U4RwxIJNZCYk*Tb-M+7YYRS!8zT^xYJR_iUEg;@eC3`4~YNpH~{y6{dHT{n&P4B>dnmjq* zdUiudwmfhnt5Ln7n&8h&p&Am94Z|bY0n$@p-8m+A2bPZDdg^!fG3qdi*d`J1qIWcX z>uE&Do-GT;&uaH(-vob_3x!vljUNPwgNmTItIbebUl9@5fAAXuonxS6icXga>+1;e+=ovY6)BqeSlGrTsN74DH zV@mByWED)f*V!~YL)s(`=6`i=*bpw(&$h{GOz3GOFA56wZRJXD`ze_!exxF}4y}sp zKdXYlvCX0155GJb)y|&Ls64InH#-q}fqH#NsVxg+`al{jRpm-^ekUmXQOdva$YtVW znmE_TcdvekUWhAaLA?5n##yc2OW6H}0uor@(%#pc(7=$8aWY44q{lhR4+Lb5)H+9( z!n#%kKg<1$EH&f1$oS`c_=!l8R?3-1(pE5W*N&5k{grW&{V>R%x+7l=vV?Ux(R>L! zQ)eSO0BtK@(dQ39Td+^*?ZZh|sx%NHYL0Z3gR$i7$am_Yk5qy2D%Ci9-vY1^blQx% zPE$*bhMy+{85Y3Io|q6#L=LKIUnN5CeOopQ53py0L}%o>fNii~_8-*)t%Y?7woPXe zVGelV1gBk@2U(kZY1_)v3qROx%m`k?((U;|Lu(#^1keTB*3&Z)9?=oRI!i3ThyQD% zfT&#NXZj9Ze8X3IW@%NL4uGf(KBdvb0xuGgaw@;^4Fi5AX{fjdq!{xXh?j3nny`{z|vRLBJls{6KE*~ zo})5X8i?Xq77f}Wd$499W~TG!5jm!WXXDv>PRX%h0KH7ZJbyvAh(K@7f+=2d!5>2o zZGq@w7j$b3aH&~=_D>{al86_~w7csfn=lqiU<6)^EXIt2|zfx)h^GP9=_ z=Rj++kq0yWBLiSscV@taQO@}NvsHkFVB^6;{!wl~#TZDg&7!018vd{aSU1SH{2y}m Zg0Wl$z4dXV2YBG04*aG@iJHx${{gaqQx*UK literal 0 HcmV?d00001 diff --git a/doc/graph/fusion_patterns/images/sdpa.png b/doc/graph/fusion_patterns/images/sdpa.png index 87f4443bf4944b2d5ee6fe1a28c6bc8e3b456af5..07add3d2afec9d3d4f53b900b6db64906aadee15 100644 GIT binary patch literal 182345 zcmeEP2V4_L7q=dIehTW@&VuDU3sRC$6+ya)fQo>qhyembfj~$=x?Qk)7Ho)Jv4E)9 z0Snkr!3s98qX?pav{1g;WG7j&OOX;1#2>%AWOuW(JM-TEz4w1_WH*my|L?J1FTXWmApgK{Rbj&S~ zXY{5qz32>g-fNIm!KL6Ra>4^Vy%+B=2|eWHKy{^|hj^t zk^joP3t0??D~s*{`+;}Rf#L2>b3g|cy(5*$L`IYQhUg(DhAaAZ=mDfht6hN{u)y^x za@3mc$Z|p7f~>}C7CAndMt63BEt6C=cxT+G@M82$yj-Y2xOl(k9>NX_uODG=G8qhT zKg_6MPs5Ac7U|JkkDPbD?+qqjI~sF4Gbg15Y&A7^HP7ib-exM)MaDix z@@O00R980I<)&1YDI2^3B6VIYUns{+2HV|{hKv&l=$wxWokg?opgJHw@&Qr?j=8Yh zT+!1209^GLt_+^QQ%Lqy5)FvFt{02RSV+U1QBqPh&;ebe>%?$pVU9a-|Dne%(5FBe zMy^8zlEhPff}!5Ogyt1VTbl zqww(1iz_TtSGqF-;0{2a5HQvSn9iaDT+u;KxX~RQ5rOB5nIYAU4qlz2Kbhw0O+&5& zCwQ;gyz)k-)I;77`9X$KyycC8lQ&bU*Fx~vgyN>-=m-SMg64`C4e)_&5BvlVm*y5a zF!(Mkc}FU9;pmZM3MzkGZvlK&QlgTTl*WLe$AOoME=j#n@+b;nqNt-fUmJL-NQ5vc z}> zMNKuB3kwz%SmhZs;KX|%AzV#%0_WE7emIj!b@%dMcuBZZ!|S!?^+?pAw^5svL9bZ6 zr3{S`F&;3l12NWv=I*G2M5&0U=t}kSqB{WVkBM#sp%CJX;@kp$w8M>6gxu76#V$0* zYVmLFu0LVqf{9FsY&$&0q`3kW@~$3i6Gbcac?7|OOk>aiox_ixpkP4ZQAI5jY{-kj zW;)Q&A0ZHc=yIc4IDj%v;bt)gqR<$rTdUjItTwXM1HvrrQfWnCQDOmA-e8A3QzPr-nnqInf{ZbF92 zgf(KwcY%T6@q&=6K+4Iu=S0U;)M+wWr^_Z3yHKF^HJ`SrK#~k~*P4#~L$nx3pdQE& zQbApu5_|^_Bd8i`eYtgj@?P8$LZn!mr0fAQ3ZMuOo(IPq z_;s#=gRtTia~-dsayq6)f!8SSIL_jO+Et`|OxC%&Rd;8&bL$i_g|_IS`WHD874cRF z)&^&=k;2L*tG?xR4h*~?m|#-e@DR500E#%&Pcf*yMH3R$hz=3UpGlkiaq%0Azx>Y# z5c3oJ{OScEOG1X>kL1Z85os-A`Lk-1zedU*1%jRA$zRhGy@)H%%|TZ8?D;84G4kwN zXJVf$lw}8{@;qw|&2G~t7AeAj48OxBB{qmA?DS9=GdG(8M1<9ys zubU)7#!F~6Dmp0Ga+Ai~PKBjwY@};4d&caUmIh`Xp3|rJYRK24v|m7PAnY#@0fJ#% zgk~{2X*g)uE2E?k#_sXwDpzM{z}!LpgjJpa$C)>RiOV&Onefoyv4#W#1Bf*Q7LNk) z2+Nf4e+V5Yky0guBZ-exaZ`4HKaJt*$c-7;W}I;>l7WYV3I`qv0>;9S5cxs63_I|U zRTiIdM}wql44KeC*Oev`2m%rX;#Be%om8Q`XVWMwW?bMz9PF^bW1^Xy!2c~!GT%RJYj@QN@6ab5V!CF9PdRjp_{1WSCI7tYQX(sgA{8Eh>{w(5#9sWE%#(4 zY%-IuH@L~n1{gH4(hpo9DJ-l_uOIw$T&vR$G`)?%ccD4FLSQcd^Ivz@B;2awgg|nw zIuU^(G|;UD6(t?*l!Z+*DDD%gQEGa;!_~Iq8!SFJ8VIQXp3`meOg?5Q zB%;hdCZ<*5{sW!^&mCyq$hF2K0>%-5u?p5vZqXD!;b`M4Oh%jW)*YP8`=15xfBInQ zL6HwXW_W)w3Tocrhw%>db|3^{MxQ^QU%ZDOW}-0u4IF=T;88JXk@1JMmxwEcsDF*} zfPo2-s6QNVC;%|0D1iciC$8`UpfncnFfAehHI8_liHOlOn7Ir-C6C_F8z}k88Ug+|8zmzB;slv1fOq3|Apsq<6gm5@p0qKa)jnF-zx!N;)uoFXg!bu<$SOU`CNyl%B zy5lT5)z4DHOr0!0L&aAf%?%5YnKh| z=Mhl6wM+jicgeU=X#oII-+E@N;HsOiU{ACg`!!8*b4tDh5YID>1w}O%x zL6wXR(3XZT6j1*R9%B0mHyq(Q1sFyi%iKHCoTy-j7L?D20W|Wsu2g%Pt1i`Hp);4s zf%Zo5176<|Sd&)fHB&VB32YSTRADgKeDN~AQU0QFuu+NioJn}sS)T6h)Mi=v5KGZO zNlPejN+?ocAXe-wDG>KG}&6NrEa+`hL zLi742!kd@!c}B5@j@pD9V^-kSq(Ut6Z|sWlmm@-bQ58rli}i35`=ThAVa6L5bcifr z&R^LG#Kylcdk+C($OATEAq^VRHXi0NQb9z4z+?4}XzP>U6r?$YR@syn&ZxsW7Z@Kk zeMTL&03$c3EwIG~(}M6Nn{`=I$ehBbhr?AjtH66a_>kgSMneb-P)J1VKsFEnYM$_3 zW4I}XPF|zbT3AY<&Ose(daV|oBg`)zg-uG7TJx)XBuuTFxlmuI3?E4`^HhrN$jhkf zxibw+X&U4OmOh5=3ufx-s2g)51H3TTxSkP-L_xV7Sf;P4f*xoq)5rHCv?mQ)&Pa6z z#nGK@;bv{2Y-cxZft{NAB3JpwEJ{RpLp0*X23>^LHKt(=`DyIgQ>|Sc-Bk>!J}UNJ z6b~2pCbfC9jfqcc@LfA>%&a6oes{0s(MQzaCLLC_X9D9sg%3moSOPsgt z5=n%^i^>&dcnK>Q8#Bbs&P23;2;fXL6bXz7h##{_C;?&=t`f3<&8W|q!GbXeR$bvS z(S(sCm{x;fk#LEe2pmq~CL8Qd@aVo{MJIP48va02{pe)s-`9-R;W=I zzAjrW!7G3?hqp$7pjL~w2B8;={_tlo_!eoE10ks47eJEaDF`j#wUOYMLTG>-55^IY z3{fEAy*5cMR&N1W;L}rLC<~&cpm-}m$lEik^H6GJ9cOu6Uk24&b;>L^3)9)wGq@{B z;N#|+d8oGKBFNAI@^Ya%f|4Qbq!}(rgYU=2@lf_EepF?jMZR=XPhTH~h6Z~fn+n;R zx+z$KwYCXwNWeo8Ue_2NitvVbz{7d;jUg!-%|l_CIl>o!OiR-(K_O3q^XgCx;|m+c zo~c0_SIxH(PG*SIxgqQxF^bp(Rp8B%PNK2$KSq=v6an(-K@xs3yX^I)On{uqRxF zph{4!B^d^rggP`Lo<-0Aa--?tpJ$KiYvg&XKy7phno+2mYDMv^@vvR*LNY|8n|Fl` zB07yAsr-8_u^}QjlIPa}L$48n4<>*Rzv@xq-a+#=k_1Se+V0in%TrC!l2(N3h%#ZlM5#f8H`20IbmaY!T0Bp#CZkN?`rVbUE zIzRel2i9yhsalOTc<8B2$}?)NLrxV4?AzD<%Yo zQ8_k;2ana;p45l6gvD2Gi*QTJz+xB36bKh?Hhu96z8nr>X`|_L_@2$Wx(vP+4vtyf zoLY4%S)yv;NN562#O|pG)|Y7+z}h+lgDn

(Z&AhVD! zjpaHU4=(mZFikL_i%{oUW9p@r{=i_19>d*>MRjMj0ZbhLgPcbEJJfk0!?;%BTo=^f zwv6l2s?Y_cy|DZkA-bSF-gtVf-ngT>lF1(TnLwsmnDMBGXuZK~F90Gp_~*_Z^t#A+KuZ;h{WB zd7+=L(roXUi_B5SL9k^3-JPW_??fTXdpUYg5t+gxBDa5;R1X(Zh9eDZT%hEL@OI!@ z5=9*~A8=CwHD%OOZ`+gr-;dDvG;C7>Ano<@P?6!%qq!mj1qA_REU_TmnA<5kyLx%q zsLk}9HF@Sj54P@fGj(Nov<(bMnNnG%?55oUk8gEw2}~s=RRbNsPH^oJ#!fkL|JB=0 z;S!V8*DB|}WKAQ>8(%oeM6(w80XPgu;7V>7Od`SB9IR1-Sk@v&iEy~_v+%oeIZ5Lk2%S8&QupYBa}q`_;z0nBAMqb5ji{`y|d^%Ry`vfb5%2bb+`USXlY zL9b0%5NShTLA-?pwG!2^0of#nSXQ`I4~DQpR@IdSw6__>q@_2rXe6enKqs~~F_8>R zK`CvFm{>5JSf*649XulA81_#Dupq8BM72J!AQ|J0auJyEZgUwJZZsAXm`Nq*W0PR` zFZue7pAj>I#$YVw5%7lV}(m)u?+ej6mHDEm% zTnx_-*y_hAH;Ia%v;{WTO4wuEc(*=cEW}2Nl_m+Vg3F&?*L-TkUn>k%Uj8nMut|y4 zA?TO2tP(YJDxX*=A6Kc!N_8q9mp_nr<%3oenlV*Tm@#3MCljQFL8svqjG zd;+$rE`cR?gaH5vpTv+T{SX8Dgo;9)_KC|AO1is8;s;l#Dk!Tci-5KuQH1QGvP!Gb zKzx3vks8==Z?C>yo5V{of9gFS(3KTbQHB+U@J&y?k#J#)y0U@>8ZN?SB{yurC%*84 zESTf1??rQV1MqDiIHX$_f$hL%dea=y8I^);DQ|wIBf1ulI!LU8d3sD1=DWa5R4?-8 z&e=pJLs?-oSvA*h84{m`w780+fNBeTHNj>mw2HF~Rj>Sxrvt$KtcFFs+Kv?$zonR6W@Yr1}Oq zA@9Ylf6@ULkd!?@b{f>=!1Lgk1H+ZUv!u5@R2@K_E2*Fc_`qKnjQ(H*F+I_L>Ex}&2j z@?#$tI*VrELFI19=L5=J!LjPVS--4A^A?2(Id^23k;SALAQE5`b}%Y~nkmwvc49GX z`V>@?P1pOisSWj^?&=JAT19$E85#iSE5b4)D0WS6gp`FaoWHdsY_h>NsHOQhSlNk5 zFBwuDSO8>&i{Zov9U=Js`uw2gMNMHT;a3!3U=t##c*7YcAqU}F~wCTRE12b z1j_^(7h{kDhY@wfi`6}ynu0+7@;pr*1lBj8tB+eQ-W3>vq_60#q>2WYkit0 z8?23mMAO2aoB)NJZMGA}99N&+gf8oiG|UZjMLgT7!Da4Hav+Yid}vF?)4OeBPBa}7 zH|B6^3r5Wer$dV2m=m#Tz%Rp2TKY8~th&GcD-l(566%uffARSN$oax30K!MpQxw95 zq2{c(x%O`CGu=Fgar?Jtdi&1!qj2Yv5TVhYz15<)_L3K+FTxNYC zvK}QSWUkt@G8Y0=ODModSA=jX?6^p*=S;z8a@w%8K3GbI>&QwVmNtE+rUA}h!|}Ji zm?-H`M-nclBZKl&5mde6VkUxSB^U7URQU8ZA{@8p3~mOE1j8Mk2*;&7evK4u2({S; zcYuZ#{2>`Qp#?THppOsPAp+QGdV)weMk2yk$pyP6&j+-LK9FG;Xke3k^>H?M3P2$G zB-|>C!V?yoUhqlEIDJEb626|~iakCZAZa$f{uSd$n%t1;MhB~!QuHU&T)k;XWjSCO zct);yOB(S)B*A_H5u4sJ6+9+w#XA^Z0qZw_&ovb9pi9s~=Yb%J57|)huC1lu*=;Bd zqAnz}f(^ILj1iK-`idcifN_@$F@z6Zjc^w+8iEZ0+9dFofe<1eL_)aT6o#mj5D~<4 z(@-2lhJ7;hJp=1L1keS zlw1+S=S zoL(XMGAITdb(=HT0B4Y-B)lq#+h5Y%ll*c39z50WJ=D2B?^#!gGx%3=HkP9)wI ztF+P661sYw@S4bEF|x(YyeYTPmfQADBvt;N+k%~%ZJcK7PcdUt9erIETIexMW;nUc zK*~w*MNZsgeA}+vpmRthRn#$Zqq1DMYfvKb2}mF)fsF$l89tap_9M4Z5)J&1mlJcoscmU+NzuIhA7@)_wOdWuE`t5Ct+IQ z4xNNu$?fu*q(yiPjc`@m!jt~VP$Dr7pE^2R*s2)FT*FF0Y@s7^i@GHWT!RW#?N47{ z*~2G6+BW(I$Wo^+Hd26amXbh#_|rwiTY~U84BnE3ZM%O1Bh?(>qXK0GdiwA?27DmU zh_~qAa}TX?NATv50aX~N*CGRwVSw;6DJ07W5miBXb4U?MNjQ&5040(kM)(|3+vel^ zHTx;m*8=?2wf~w$^hotUFcK6;@it>bygfHOMHgKJet?9}gREHss&7+HcpVyL@@75p zf_1{{xS>exQ4JywxQPlFw+G4;R$kyS(d=ff4(q#7!%TrCnlR*R`Uz+1(Ak4c!opYGbcvZ}otaGBUnnicU^BV$52|NP95V)qNMAX z22AvbRL3-MIwqvnlB-~-&(yMpuk|Tqt57x)KP(cdw&4p-TaC8ilOU)gB^d#nxSBwe z+N;w9a=9BxH-RKoNLHX)sjP}$2)yb2Op=O_=~aex;gTQy;WJB(RLe#fv<581*MuIX z5tFU&XyL>#P*HW6?CD6NlQk&vm#s@d}*ZG)@kiftH_gT9vw=o-mSw~RT2u)C@Mm3V{@ptkSV6X z#giJX-jIN9-mR;J*C%atTUQe^EBI7!3(W)jEw`>F9uou&8{nFwJHH9}resLPVp2BH zUB#wsU{)HcYs;6(k?{HDM*5(1BB zAt)+PFd(eK?-C~{(6z5Mz;$;74mx*ch=VKD%Zu&+;yM?Wn=7!d$PTE2jaa$SpF4}` zI}0(r5@CrO;EiDh8CJt!(EuJ3O@)k#s$jDPD2>f;WC>d3Puh@pJ4+_Y$BTFo7v2gQ zO26$M+0B94BxPaO5JpBAjg}1i4t?6DgYAX~LJe7KUl6VV`@#KUQ{7_J5Uw?0&_&1+ zl3E2z2=O0AKv2vfGB%37b4|Ev(uk&trji~AhF5@ zv8w4TTSM$q!>wVXER>`NBN{@r^JGkL(u~FjH}8s|S;u&%)Njd!ddxRNQ*}JK06`jPf-rPyCG}JM zUGMq&LIyl0cYU1ZQ&JKIl7{?Nh1!%{wG^K+1S?qo@Ea{j+Z;TakWG__@S1)MfG{<=u;pKBiC7o#XNPGZ-O`n zy%P)r3u6duP6dh_gkX`Dun^0R_~6>_xDB5OaE(lm)P;wzvI!XrF_O(DFR#Ku5Febl z&{7N3G=zO=A;Pjw7nj5>z>##4MkqS~qS0n|Z-r4`Lo%}i#D)N0mDT1%01f7aip_!3 zxMUMVWr=elV6cvHKyfbQudQ|S8vwPW<1ZOkg`lB8A!`W!p@0JYiKfri;?`TG)cypL zKt2#+GeX;dKWu>4reDnf%oyN>L>AUVhQ~xx5r@++bs*g$<=k=lRmT%Zu72S;cL6~{ zM4wD&Q{7nz5kcl}4%J4Va93gDO@-tCmw5Mh+@cY1Z8D4kN(OyG#?MA?^O~F0YsMu4 zC0BNcs5ivSV~AIgdCmN4cR;gpc?%n{C}iXnh)TMvEKF4h)4)yd8n=qVDq>$4%_w4G z%{IgpDdG|($|S^E;|b+9Kd(6%E?gCY7#Ztk-uj8|$O}--n3Hv9Ql>e0O;**N?QGy= zM6)pF1@D5*p&V(RAOXvAkmnY?;H_QTwofjol(yuyPiTGtj8q+K9}w)aLr~4snA=&+ zboXFuc(C>SrW(xDm}+n5?QX&o9wCy57oK8bQ)lob#4#)GQh1>~6MW)U4ZtJ9)dIiM zPlX)x##BZ)IM7s_oQNtTD5?|<3IGsZ#X~J=Nt73)FxcXW<*jszalawrt^aykW>pV8$ z*zyz0xJ-sxn-62T@E>sV-8NO!T`qm4*sE*3h%=FD>Eg%?ih+;^}zgo1d z>d9sx#Zkx!c`t6Akq)?kr0fB73#5A7KbT5W@~+ z+iQ_Pm@vhH4g{hW!5l&f1Zsq65+-2f9cT%f2VJmINdXWmfan2V5uOLBtwTNwW~*YC zzT?tC+?5cJ%MhXn`3+`S1i{aclfSZL9%`(xQ|t)4u%T#%u~1N)u(AS=i6)^V!~&6p zCAH~UAie>HIwahrBxGy|1p=S25jm%fo?LQ?-=z4&KY7?Fg3E(y#jZJ zI}M!27S5uF>c5W|QBkhIIYm}9sjI|I6uu-bsR?~cjHXHks^b$sEsp9C#xd|*4^UT7 zr>J9roqHE9kRuTVe<>V3taB~EG#ZocMq|;K=q0>&uXW!H6zj3oi8ZS>o&&*q4nl4N zMuH%&w;a~vtdjwk_9zg1031Hp|k?jjV5HGLvrNW!LY@-8du3+h1gON%|l7;vW3BL4@%g9PFaw8N7 zd`%d1ibaP1zb73?I2w(Fe-bPvlX(0WrvhKV1R^<{hH->gey~XrKTRq(7{36z4%n~- z9urNLr;YjJ0`@QCIw3ED5-phUPEfB2LO*PdMc~{1dNkkC#Trej4QP~)qk{>&Xb7}e zZGnM*i^TYZLV$XO7$Fe3akqCO0W3G<{`a>_`89`h^7k&2BpXjCop=TcM>L9T5>Eqc zJnvy^P24r<_=}aTrodJ5p%jAUz9}e|1o5xw6LreMv>O@SWPoOr{ z%9G+NZ|`C5=H;kllNXpmU8H5~7rv z7?#Kr>ImMU5v+F5iej$Dj<~Q9hZ2BsaKtXWA{=iq$esP3+-V~v5Yk%EcO?)&vCEPX z3Se8746V!HdA!)_Bb55UCd3H4uq`1*oIk?mm5^N0GzOT911^P?Fv1`%#wC-yS0PA< z5ekGDB*fZgJz=E}|Cw7@N@IX3YNoG1Bys5jiG}r6ll*GI0{NPCyY~Ew04$6I5{B_i zq3?ht{^7BXZ(Y{VNI`^DK`nA9cJv#e9JbMKze2wec3~U+#toJLGS~dq zTIn}Jfe-`wtu3cgvn&EJC1&3wvhq-gzW8(h^xaPAG_I-EXy- z4MKn;;kDxp#T>CRwryqNAh}iRtS%9yt8A={O+7S}JXE|HN=6peIx04vU<)QhV+3X5 zP9E~ko>b&09${)sR*{v|P#VM$R&bKwriRW;m@VT9Gg%2afbK$#q^L zE`aI_-zZ~{o1dM*mKnfh4`9?F(lx;ehZ_tkGOida%u_3}l07KQ!R6zCAONLy1*Z)4 z>E1x;;WgkuwX|Im3JX<9g-k})gRj1BC|HO}Ri9)CN`#14_tfLBQcZE<{)5+X*S6Lt zDTD9zUttt+jt!}9bnxmF{mC>}ZyJm4Kn1rz`>$CSp;v}Uo!aCLQF)}-#arI+i62v{ z*Fx~vgyN>-=m-SMg64`CjT-Xfk|AlL5EIqb155b4j0MAqWlCMd>qpG;!oRT@QmYRv zz!FOFh6Q}ipv~(Lu8!4!oMUZPKr)7#h@i*$8Wt=nu(2~}UM^G*QTQ44q$n+fcBE-V zeg;ViX6nVCVOePhX{{V03flSfyLIkxw z!-X9{f(io)j0j+l!PRYmOWiE@hu!mQ_vg}QypE{POQQwkd@-&4H|W*1@^^4BaQQMq zZ6eB_2!$ZO3$AdTw%Yq&(6{^n6E|EDa-#VQlt7(qeq6ydDOtFp3>|`B(|dtv2u(~B z#w~vL{M!8`3|E>w9fs^_H4Bhczq05dj$z~b#0dW-?D?X304|fpbFOM#N(tri>-1O9 z>v&4QWBmz{7@%(b?4(d8Rjk|^Rblvn-CL~e%2Nnhm0T4>(YxtXu;FUa zP(^G)TZX1P>Qtvz5u&R@{DlE0fnYbiI^kGwVGAm`x`cvp)9F$}%$Y>%ADTz1Q~z2; zD61jN-ik#3aA}Fwqks5hbqlzY60LuzGgYVlaT7w4P9>4}s~n2d0OGSYtw#g#bpVaj zz=ms6eZ4}7mtuZZCzN6hl$&b4kcdPUHeDJT$S&>Nf+pCkwi`44JIf8;#G_&kQ zb9MuR&q(~S(fDUtTrFwbM9YErq2|(E93LJ%|f*Eb3$-$=pPb# zM1xAz4DlFDGd4NB2r~xFnl;scSxoftRghMZ1YfJpJ|H`MATL=BqIeZKBAXUA9aG`S zYgE-4Ocg*YY%wK&EZQREI0pWzqmetq9n?eUVt1`Y57ockmuN@a>bP3JjEHx$76uE+ zDu1pHCdC2Cer)9iEF^;<(DG&IB%*R15zC)RoBRQgtKm;5hlWms5RyM-Xk8>v{)oyt zL@a++ZSq%x@ByK$pZvkVSMubq>C^v+^N!^IFDNbem=_>`EA3we%A>W%KZ0RL7hK61^#z8ftI^1Ng!oyBPph?QG6J-#D6#N(6F0q`z&;T9B z=rE)Nuqq6XiDn)pUz{M)s5F0eAA(d(O-in_MU=K^6MA0|eu=fvJW845lkV~qlYOeig4;MX;MilWWH3#R^K2e>{y29MiA$ULFxvwU$v zt*r`41C(}gDKRKa&5w2jjf#roWm&C)(v|N_pBqG>dd6N zp+81{0IE$u(Ft-pt4`FD?b;1(XRNO?bJ4uqaDUd^r-6CF{)XG!jJK_I8$+zf8S@RotiRy#!l^_op+gB>bv}DpZ^r5?E0^BkE9bOFA|xb7d-0r zZ>OEgrjPy`c-M63uxYzibw3v~?oQvJ`4ioecWRvHO!mJL%YLvsP?@tk&iT^>8=Kn{ zPR^~Y;$5Y=c|qjs-8xNbhx^-mtjnom3%dW^vBS`IxW7p<9FGn%?b-|f9VK^=;aECN zQ4PuS<_Ei`xnRP0G0;Mf7lYyG+Wl-l6kwkc>y@B+hxC z%$HRcX~majj(XC$TFVQTv`Y+C5fQjQ7iinScZMlNyf-&nVL?wOVkPg{d=`X z`;fp1ucnW!_Kx$NWVJy5ZtUMaq#`kOxI?uby?U>ie>pfN;_Elx_eLVR%7^h z=j?)$`a4WQLlRO}RIbe$f2la&)rh>V!xwB1SXpr*df}Ik<{AiRHfd>p9R# zm^O~B0d26K_deOa{gE!)`*Qx9oRvFh*Ixx0TP>u@JG*zBR-9wrAwR6h{N;^^m+4-~ zPxfB>WB{nMP#V8cLU$djE-;tD8xB7DIp9k*> zm@k_cWbeLsE%qhH$t+Gk-)^wM(n%*}l!g@@T2|;icGWkPvB@PB6U)obR=oRj@fVYP z=QC?%N0p?lKC?euFqjxMy`X>fRC>gCm4x)`k+kxu_X^i+E-_eo>PC^Ol6}g8A*oU^ zdH?C{QxN>NU59Fd@Nc4)A|>s0!R;sU0adFt^RJ9bsw^8Im1AmkT0L32p!AMjX<>B6 zuDS9u=W{1ZR|6l_u9I?J9^7zRXYbvvpF=9f8ZG{`1gnaR`v6t6I(1;0%+c=PQ!Lf_ zY%dHTT%8CGetX_YrNfeA!`j8zsT|hV!}jagwP1i&4EW#E&P%2Y2j`8S>DprZ?^q}} zuVT~Nt$q7j00|~I$;`we-CiefFl}w*!)Y>;Bfynq_off8rtJVUIo~_h1Qgfj4g3rnt~P@{u}hin>Ctv8nUTAc2ar15Mb zyQ9PNuOaCMGjB)a_|A)yvpl}VPwD2(F`EO!Zu>th$Tv^bv#xSD^u+c_!MTqtJzG`( z@GReceKhkNquXhpa?{&JeLvu425oJe`lb}^AKxd;^b8u=KO@~x?aBQeh0&Vj$>$c@ z$VYJ4RwX}wl=t2sJu}K~)l&Vr@-J^48y3~8eDAS}b`^(D-~Sr$$I~a$XW9?6&-{aZ z@!m_d%p~#-|4+)Qp9c1ocaLPa9=ufMXaD>}?oP|OrSr7Rd%u6LtTld$!>S%VhR6=S z@3-Q5VBxxjm3M#MEdRp1B$rGsR#sO1&@bZ2qetmaCjI>RbIX0FxzT=4(-b2VQ%g$J z*R5as@LKU7Qnw6mo%W6P$jpq_d%)SM^guTDq3sUWPkTNcUK!xO(8R>cqT~GtYjgJZ z=wrp1r~WYe(2e!svwHV+Yo}LLUc0t$Q;h2{y9;@;e~i5qt+#H~g$4JbO z_FiHB_(nwe$%X6U+y>ZgJb9cm&%#1)PE_1U-^y~|D9fHdr*|2o!Q4_-CbiPg@a7-+ zy?aLlEmG6Svfh6;!tF>m@{UV+vEy7^CY(HT(pP6hdAA;er_NtgI6Ce%MM+s*Q+3+t zsR4c#373NgFVIVTn)v#1uXkgwNgsP5oywT<>dJ$D(~K*W`pwtPd3oykizCtVKTLf2 z=($bi1;@~Je?OVI$cm6K1|Ma_L^li)%1y8q(Pqk%ys-(_-NJ*9l>b4B{1 z@{JVd!fuH{z8TpM%AOf~kDgL|=Iqyms#KY{OL5McdKs(6+OL%LRB`L~BiTznDt|y= zVe})P6=%zVm=?Ot={oY#3u&*=!HGe4pWS}-<6hCH^{>8ce`r2Hv*(W;A^S4o&;Kbs zFLr#zqa$G#+$~=htuH$Fm~?>gve(D36r*ka9H|UHms`u(3B68Uj6C4Z`I7v??AF?# z(xQ`32bKJ02JL)y zA+W%VXP`Fg_CmlvPqb8-h>Ue#mJ-B8#(K<`h@FeMWiUdYu|DM%Q}vyH(wCgA3O6 z%-d%!f68)AtQ`4<^X`~qU03~hetK~!vy@6=#9wr9tcY5w0$#jqs=rgoelKg4w@O+C z-chlGO7{gSx<4o7hX-+j2Hzaw`83}8k_`XF&qse%@racD^tW6{#o2MEzeNtr9=L0W zn{C|Tktc^wH&a?t@$!84^&8f|(v`LiFAMxqu9bfH@L{906Z6_@A0ZQEfZZI2%1$M3#Mp6xa7#ifNF9v&aI-IB^FWxKrZX??PE#q3k} zFHHHq>{wC9!d<&b(FJ`x_sM=5{!d87mJ`>%m0a5_KX`E3{OF)5S55wUA2IyZ%gfLG zzxPf%Y`bfM6)oakYy`7=QT7p2tfxBd1Z8Coyt^d_M`?i*E*=dWYYx62T z3U!9J=k9OC3#0+)^Xv9ZvY*@v>-M-nHY~z!S0M zDE`k**gRb_vYfuzJYF+?x$DbwgX5%g#~u7;y`OD+@!p=|h%?8)Z#SR!Il1L(Sbmk? z*Yz3SwETkxD-`v#w*8^E%d<~~ob{E(qq}NPO1Lwab7@}heq~ez@>r9ctV{m=diU8g zV!xN}-{<#5A1~;ppsJZWJ|;F(+dJDg>**O+)!hBl_vl_t&Z6fXk6!QpRa-^PP5pfI z{96w=;b*eXJfHf);J}sTT|X3b_WP-ACtGz>K@q5j>E)R4!Iuj=IzLUjmeA{*cdn!M zCdE`Q#W=0XMG2s@`fx&yKTzL)|HA6ZtrYi3#wG2`Ouwb4{oHvWRoZLQbYJ^mhZOll zmySKxk?hlcddpYzWH>5(2<+GXpw*F=5&6u1^S00S9DP)INqlDc%FhF>!#-A>p$5Dp z=PO+K{_45rYUivJEqij0?fd5? zK;!`rW`kRlSl>NUwto|Fm$bfZD(=*i6C7x3qUaOzBl1~Sdu_{C(^u^1VL1I`XiCqs zEMrTj@h^4$uumXY-26FPFXyuRFtu!FxhUP7%hcfM-5v)m4%$4Yh#jYTjzd+*Uy^Q~ zJQiT=re|Ysv6LQ--4Y^G$1Oqv&*w}#9vkx-^f%{3k``m_rM1BeqARUK$YC4X2 zUVQSxpE1$&y<7gvxWQ_0sMYV^HN!U!<;WNmc2;eGIMme7v}*z@{=jYJIMPYKCJtbC+M*rm%z!vsP};qG@jPoxk(* zs&uu`%Q;tEk>^fU4Szf!y$4x_V|Ljn$*PBag1u?jf+V*j%aDqh_XpW*r)zuc%RHNX z!1|$-cGbw8lmt&h&IZ{5hwK6i6zR88B0?P<$k&q2XP6wZ^X{JeI$y3V&0m}G&)2G; zJhd-RytG}E%}4ZEvnez;Xx7?ehqs*_7o)gfO!$>(t1C{4T8u^R4kcEWR_ecLHm8NX7GM zHz?DA$9$m6TQ9S#R&!UR`y^LQO?& zNA}0MA9sG5KfJ!)vNBVLQk9>zBJ};^J$=r6z0oo5$%PB|QTPWSo8WWC>6 zX)NecvT|ML*ZTdAty=RkG8G6 z=d~0%$-lld+%}-Y;iXf{{~HytI6XpML(_US;BCtFR`I0UA-wR>_M(pc%1k{I=;f8l4JeUwIW=M;W0%c;p+VbOapfBMIt4m;CTFVt zITvL%Fb3%Jo&AC!;nfFMiD05hT)X03PKC9*49lEwM-2Ty((cdOs+S)%{ zPyhUs(v?x_i__IUY`SQAIc?LeZw0sCUoJ8qm)Y+~fkR6Cut!d(NrTt27Jk%CeRNXp zQPQaS!*3@l{!HvG!}*wc;|vuEe+G{41duT#5XG|#sf^XLZz?W7Gx?VD{iM@D!}ea> zE<53DX20ap_wSSXbgMG^)q~*S6LxW~*Ozr#8_+9*_lyClt z-c#j&MjT3yc-YJCEG1td|F-|qkfBA7Ii+K5XGHlI_pDM5vMNqfUzgeE`0}(Hhg(MD zzFjrY2#%C`6MC|p?$PM7IP-gxd!L;-uW!MG$_lA{Qe9Rr=`yV7W7)B-TcrLi8QZaP zM)J1D-&0Iwd(HY*d|O{TC-Gju_tTunU0GTuMiiCKE8AOc_G+G*&C2m1S2ecu;(*{- zYc;L>qo2vWEU$t7-!!+$`U6~b{-2cA>Kzr^^*5Wk0|b2bjIaLs8$<2JCI)HT`|hd? zSV0%V$<%ot=oEvhDFZG4?rO??i^${>4I|I-30DTiZ%NPkQpE6QuT(diowLTKY)$0{ z_u%OG_32^eCDuvtF7H$0hh&EDH2<>O`_rVqwXBZ)uXn|%zmnp$=cm4mEi}4fnokbA znLa^=b1rU-dubO#&(M)OGeDP~JW*Sh+RJJw8DC{p-CYr6k-8_)CcRtx!DrXn6?OZu z_5nG_#^ZQ$$*msUOOn5o=9itKl^@f-x#PL?oSP}moTN9))wQqQH0@-(rs8zi51-=i zoBu8iD(k+9!Wr`VuJf0liIsaIdijPud_vj&+$Yz%pMB<-HQRd_URvih%;n3oK40De zigrWF|I9C#x3kX-9C&e>9x%ZEMYg9ZL&x)UlqyviYLmA&uJE+whu*nt>Cc;!;>Vu) zc*SIJ>aEA2OG4D`?S78_`5~xusP@G@3XApw)yrrf_#yA)sQpP6{zXYcD*J{0nbNaV z>a)R`z}0haYLEBbD;sQT+p$Yb-qPcppR@;$bOS!pk4wi!MeQt)QONpx+o{FJlGD7$ z_trVQTXklB`@jXxhDFA{2?zbcQ(b+Nep-aDf8_f0U742*Cu+{^b!9wUPq%w={j4ma zXyeyKq%f0{OD%Px6~fal`kCj3MS6a}x<2jm2bHW<^Wtvh4$Iot`^NsLjNYn!<^C#3 z_$Vi3`(59nLh1GK;=|JayEk-iNMMy$)eEf+pUr}|1${CZtnU)O$?W1E4v3KOQWsz*vktOd}kG?oWhI47pU%R7rCdR+={4cPt zeS)U!j{A>gIA``*`KP3F?NI4Xt{wW&m&%FlX>T|8TGsV3 zZEW`J;Xg_*znr&ki1W%0cv$_-|)CA5)#Q6=I$rzICtC_0Gur-CyP( zoMYT?sa5I_%Ed~m|Bb|9e6%T>BVJEU$X-X?$KQe!@B2Nr}i<*zPfRJM80x%+|6t+ zDrejcv++It-aa-h!F8fEU?~D>b<@21J%y^S4wo9Vk|baE}e!8+M7CfelKe~)x>CeY*& zEzx4tYP0$;))6?$2mV462wlO^uPYM?d|R*dc{P?FWvT9%685Hg;nkoIM??d z?WMahv$(AA#c}rijkAEOeRy`FR=-!tij~_HuNoBK;%{@_kUxsr`v75I0 ztS44iq-^@n)%qv$iPaXRnCOgAs#pOYd)mY=R%t^HKYreW{B(Fx>D1EG z(stjX-@GaAS&$}|`ErtT@c!eYi%R_iZrTl!dE#9pb+}^Ap#FcUj83Z96LwAN;FQwo zsZ-<1hE%CVpP9Ae&E26hF9!99u~WGAVo?UqbxWAaX`dKhb9C{feT-t!`(%@SPm2?o zANF>9bfNzOTdyJVIkPPPGPGZ*?Bb;DG9Y-2)%fI+iy@nf%~$1CrNxHcK}19)G?4T; zG53bl%0$R{9N|lhsQEG`kvEFPRNX5pZ0h{uC$KlO!i!}dUV=Woukaa zLG}rAm&X3Mf&>_O+k!!OfdnvL{oH^o%>B?Y2K42OUK~5kgcb1~4?67JOl4hp^KitR zRbQ4^|50I8c|U2*=arBD4YSA_nEd!fmh{8moC$v$zqxEZHACIA!o0|$S83JDQRe%e z%6om5%^Y$_+yCt}$H1xoC;=bo>j&M`8DWZRE*TdX|6wjl11? z=}reGW@2Ujm>h-O*A-`h$l~cgMH&{K?PN1&4F6;O%$Vr1eU}Csq?sQmN;d*lM{1vA z6`L+IrMP&^I9vHOff}EZ)`c$}wP9b_?(vM-IdQqKI)~nT%5>A(mF1T7^lV~KuL{}1 z==~GrBZD?y%9eA99hR5yJDDD9-})G5(M@uUO8p-V&wH_T9V=7HeE(sM$Csyg{*fG9 zyv6ry<$nfGlt<|mmuFA(V`hy$)YE!%QRv?C&AmA1Dkhz^99*JPH8wr@U7607z|Ei4 zR&-rx^Zizzq>mX_XMj-Y(P6WM4KJe(2AM26ae^^x(z_Q<79an4ElsgB4-Xh$>7KIx z`fgq65$BQ}GQ!HVFOMtf^fu()&?2_WuE`F@ISg4%vxKTbtr-*K_n*1G>*j>vN0c4c zoc!zhygln=I(Djx$?G3zUm6hnv0bpi$Evp$pUULA_ZTAomvyg@0E4|J#>ON>IQ7h;p!09Pw{}~{70KO~x{*E5`QIZm zUbvd(7?t=xvbQ_2(&GKASL00QsEm?6m@JpQ^p9>mWaUz?U)-_(i|@Gg1FcI}U)fu7 z&-&{flI^;W`v)7B+Roi{WX^`8nQ;dzhEB~f`Wc*>@aT*Ow}T?D&yBM+$T0l3o15q2 z7#ExLt%>$aKkQyux#-0?qsrZyUF083dH3?_^)eNmmx@o9v}dn4Fg{g2SYb^?=_vUh zUq_`aOC!0w(K&JRN%)qo3rD5cou!{*gcTj@I*~Da$|c#ubEGZ4ytwRN`R@%Ed0*vE zq!(r@u3MM6&#p}I%wLthLGSqPvt!m&ynVmN?qylk><1~%y{(kfO85Kcx$H4>yVp0V z*e<|hU&8WvL7@-sa^`$pV2_b4T`{R~&rzJ9}?<$d(%zQYCNIcgt64YRaflMY&~5 z#ul-2msf7xbHpGi_MbA-{9cuNZ{jGC&OK@P$o@e{y z6jlcBIPJVU;6sm@v4if7dciQs=@)Fi<9tZS{gR0jzh@WdUF@c~*5U9hdQz_d-;XbE zdq$7mnKXa#xPwUx+Qs;-*G@j);k)${a?tAw$M(b?h??|<7D=h1?+)F(C4Ix{mcd2cS!>RD{J3A7ket?ADqLmg-I>vmA8DO(BNSEAhuVh+O#1dU#dZIgsIr{1 z$DN;!F9A)sS6f=d(zL^RB%wo%(V5)a2tDKgUJO{7@dSsrc`d z15P(TC!SU^G#Z^Z@n!KvS`}qANWT9y%pzir3@5>_xPKtiw_}B06)D^0L_zWYR!M0-u5?UBW7FIqd$k5(Zjre|Bd_*~DD`(nDW2 z|HavFn4J@Msm%1=(SpM4A%T_i2E|18wqB8LUA!?lx$vWWCu7CTe!E?P$tj=o{Zqve zb?-k1F4C!*Qu$z2aLVQ!wMePxyUEXQBnG9kqjcnY&|0o{^fI|)kIO+nM$G=PX@k^*@?8^z zweq`nQCIlg{PVzC(+@`EjNLsy4ZeTs`tJV~WcM#vPn*@TiZR$>#gEA1(`$X-uk8uG zl?JDL40awMH)z0vlDvWy@x3*dn#QM8MP=pv2+pdw{i({dw?!T;Wcr?%OZO5#IKMAl z8+NBs);?o~*5}-m+}`UJ%}jqsi<2tKxKJK>)V(WzD62R>sh0|_^iAj-u;Fdg+pcAA zzm0fYl)8rIMCD|ebRH32zAQ0*{3Tmc|00dAbBp(MOuUk`{khuDkCD3rX6$&b_|^_2 z#0S|$9^5%Q{CQTTS%&=s6WcLcpRJ8LwRPOyi&rx#JBz~}@A``~-qUj2`+mcWHzU7==YBDZN`h}OQzIET9)MdT0hpxOto}ZIC>CO$$W#}bJD%?4I8@m zW5}l#Z~451UU*~wLaKrR=e&QE^MIoMo6_|44-AP33;i?rVPI znBc19WvpV(?A5FCe)PB#UA&=RQsw3ew~j9CUL1bsDu z{f-feCcE9|hD#OZ&NCWvc|&ySrm(y<%gX*s_qU2#_|Dg>(|*X*;M?a7mh`{T2W3wJGb%6(e=B>4EM=+!Cgq{^!c6UMCeVV6!w zarpX)Q5w-LCdD~@*@M2*!V--FBc7gdd@^wJ%$S5XPae?aWd80OTl%xtOX&}kj;ydB z$?U&=`ev<=9SB4@+4z!d-=vp?S-HEs^BfXq-#%d1Q!XGj?@V%)y!(I`b~i5!7!jv? zb!mE*{jSW9$$7L^Znyvv-c7A+XueeakOK8 z;*fx!H``6O_U&2`nH9A0`nDh4Vj?1U1ugQ4bnCEx?NPmUrEl70KJa+9RqfH&6s>O& zbLN+AWoh4I1s=#xX6KcTvf7@W%DLDl#o$nP&LI2Xb06MmIUDC3PyYNgIQd9f{`za$ zc|#7QMA7rZ%8y$Ybq=jKT3-5Q@htn>n|_c>t}V%9gwGi??I>`mVzM4Ml*ZiZ@hae^ z$Hp~@U1#UdBl~s@^EHbqtSl^ z8P{@PS(y(>VLTg=eLN zn1x(oP3!(4JASfl#Mb20*uX6(ONz5@oPT99p(;cyBfbAq-=8$PAn zNqPV5{?7Cb--{0>X`Np(wuH7%(e*FypF^tVj;qKDUemGcLvFh>p~v%wj*m_`>^7oM zdz(s0R(?QmTK~WW2FnuI8)P4FuF5S-dfMe(>ZG5GwJqZNT7USmZQ#$#VJl^C#}9p= zxjMUQZ}1`Kx1%{Xzo>oLs=i}NpYlUd5$|>9WUP7c*q!r3`ymJqr)!3ojSJ4|*nw(q zx1PQ~O2_)e9jT9Z0=0Txep*pMUDa(hIdRg}40%Pf34d*reRJeNPC@+MK~;-(aL&hN zXR2O(|7Lw~m5yiU>!}$-W?#vlny~%o*QyUo7LPm<-g{C1`GYts4w+t(b=tq4p4A>u zNijB2&N9a+_RN;_uw~~TnEr9559f5?Q?q4rgKZPmjqj9(zNx6!pj z9y5+Px^JD~lNP;wMS8>+V~{Hgyqr45c+FWWjl;_}S?$yRYm^HB3s?F8MeC4%V>?E# zjy-oqd)lj2u?1NvsXMhV1ihd~YDPR$$ugb8{5LwAkp}p}qetqD=@n6ry84r|9Cugi4KYW)UZF-n&Fw-@T#IH`nIm9od|=+pbXO zf$=%U#tT-(8u`vo+K+pFPRHIZ>d1JHeNh2Dml>o3g%7QdCiJFL=Gt&hp8rqblFc zoP9TpsxteuVZ-C4=Xz9W?1()2GAgTT#ft}i5h(|sw3A7U|E4&}?FcU&mp|cgkj#?MS9g=P*{-ahb?8Fyxb890=ghl| zwk(`mILgO-;J9f97rOM4%E*`!P&zC9e+SCurrMr8-NmwCQ&rs8e%iZ-soJ6%qgl5&>^55%rrd)ZjM^LoB~wgtWyOl*P*UpZA`I=7ZPnCKmL=7}Uz z_kHhYk_&F8tB@}zcY^(tMkWbsPleR=OUs^^Bj^xLt{U(1-m zw$|H2ZddDst##7oy1p<=2yyN*WF+(txAlAotQ9NMC!eEZY}8v7MlA6?s64s+j_GYq zoIEQB-Si=v6hyydvu4?6Fm{lUPzQmK8`@_*ytwO}R4MJ< z`%b{i`JF3GEwLxXvrr{IUKw60VP?B2Z#*Y_8`({CY)W@UM-y;a!$xIVU1dQxBb!6nT1w`ZJo%zw>kY*^^NcxLcbCPHh1p)Ah^j2rQmfwGQy?u~_p3DR`=3HOR7F)V^>5pOkPW z+i5W>dheH@p`ehk1SCuFIWP8+5OLsqRk}@o7HmwW)S=8@Af&$+tr7 z*hX0vX-f*YwB~@Pxj^%pdOq^`av}(crN^vp>2j|GIWg`54Y37h@k$4f*Q^s;p?I`++bLGg z(SZ)+HA}kZazvsc{3!S7X6UeYYLk}c@y(l^2D#a+g&G#&>B~xSRkYU~sEYVWro#Nj zfZgc4DV<)e_T1B$mG$lAys1#g@;qcWu(Zf;WWA*L;p}YQ?s`C5+fLtXk9^yX`^rar znjECn*z~jeuP=<7#8>)w!n@e~4 zT25SM5yKYN@}XqeA5Ix_?nD#@O_HiCZFi$Y4sX(%u+SV3f)y*>=0M?IK?3P{wo4!id z=b{_=HDV^8Nykvcel@4YmVYXZIZ`m|p{Z3AUbRwmw9KDtE@?LR(B4X|B%XA)r>;W0rI-~iVHJ!86yEnqGrsMp!JOYbI_kzI6 z*i+rSm1+U}of;w=a!x z{`ps0y=FLv(^S#9a6w-uyt3ngt3^ilM69YOHVW0VpIp||{#VE$q}j#q{>8WJr-$97 zcA=(#yXpDD0I<(>-AbISHK&PGOrt+%e04|1?LewxgT_v_RCIGmW}KdRctkP6 zAbpWG{8yqK>`p=z1XfyDl?@x0rFY+AgDA4Sx>HAC7kcZyw}`BilDq4?*2b+nHhRDE zjHQ5Tpn&tN<7wL00um4+K>(Prsj>{)J8)+1${DdXy4+T%y zHQHEA>d97z1n|XAuEpLk_&c)>g_F0B@cGWe1d74n z6U;&1tr$oTEaH^XbAEca2JhHaziZXbxLz6F0H*eQ__U>?M)p&r6MB)sPhg>0M%gV?)m*C-8Oh?6j2=#n^YF!o&Cb;Y=`F>L( zsN1Wb3ve0;zzsqNZjR;Z+tDeoMN*Df@D5ydxM196D3P0;W865?IV%$6^z->+4ws}v z;Ewq1m0NSRcP;8->8lIM?H6O-lkQ>K_%`|O`gyu^IHD^z?)e`)pm1_xJ^hIhRV&@Z z#YHGA@dhP08S?ftGGbj`gjG+BDfv0ho8y&VBbY$1Sqjrzp!$Ep0-8?qVwmO6y!e{t z_Kf8)t;SZW?vsz8-2|TIqMmo+dWYi@^Me^HnO6Um!QYjzRL(t4jHw6{5Lor+q_+}W(KoFw zR7%#iV_N*I+Fe+19cc`7c;dNY3@I4WyA~PcG7tN!dh?+;W_QK}A8s10uW+k7vAkU3B=hwX1Vn5y&0qX*(jZ6@@=U=YQYlp8W= zrC@^bDcSV(qEuLKGOl8PiH{T0Cr|u{G^#>VUlmQAnn%i>bwb7x0bO^>fXoD+j%Kug zjpEMeKAU84dl2d4wdR2P;Q@6Jch}rYRzlT6cB=nq0K>`{tb|sUmn4$h9yrY7nVc6Q z$gms!oKHM2Dn>FYALez^&I;9Zc`)8NI7gjhlnc>^DS_@6ga=_2-cFyUjl^Hu4v4e5k^=LS)DIOE>4oEY{Y-l)W~~L zukqztW1r7@gKKdTEd%zo{h*6?a^-4!Z3;Uqz9?WMuhD92%1Aa{^ELOI)gjudb6WS~ zAWQapOET`zTK=6jnb+qmpg3co`CniXD!1L1S|9Xd^Y$G$cs`4ra`Q(IT~`{`5XaBo z`5NScW9u)3Oo6y)ab3lB5inI+p{`kp7JrJ?k{|GNgxksmcLDjT(&w-gAPL0%!*@uA zs^x5|Tpd%z{n#-agH4P1D{7x*|HE=S%A_{mEj^y7PeQ6PA?8m1`~0y08&ws(pIiQWK0 z)5jK6diQQG4n-%DAu@#4#MFWZwSq+;q+k=*V>AnNEk+>>^nXmEK){kc6!SR&x;cwq zGIZDRPk8XUUd7AXlaZQR^H%uffH_cke|D*+2eVLb=Dl(lU~px^y@;T8EhGWqBP>81-{%(hhZIu0KAwVNZIX zYbGS~WK%odn~JrnC&|ge-WVKiZ2g9#*LAi9x{((~h7jsdXMI0boPfFgfF|b5E{V)` z&+>H4_Q=Na)#0$b)t$E;6(^xNPRI4wL4#T?{OD*K46?v*OxWfC_s4)}c1Zm(cC=xMwca<}?naN$KeIdYpWT_8%lXni)mR{s&l$uk)oK4K)z!WLM8M@+ZNUI41PGra0U}U*vU-|AtIW;K zHp>h_G4^Cs~jcAx_Wyz7gG8JCMj&cYE{%Yj(E~UM@EjW|B`-H1$+Ks4B$*i zq?pX!6W2)?Rcf_$ae(b|K|@1BtT7vV(fq*MB6@=~yi7iRl0pw3&XN!A&kIZ`0DbXI zI_zbmkbLZc^D?uIxcTLP@UHlGI38Qtte7EW#kI|Ft#Z(lXhTe{m54Hgzr$ zYO3ie_%Ab37HL*7TJDKvle!_{b_n38ew>=vV$ASD0Z&BATJ5C9PNKQ!fQd()--GS0 zCyI*OhhGEz>Q0BmRM?F{OvCcXpgDH<*sKC-_(1jjs{>{rqvH_U&f4h!Uu{l{{*1jG zoZa%-Qt`7NoDqbj`(Bn7SRa9MxaJWGA@$N1(s@6{UGUVdEcm+-7wj z(Sr4RNMTkJ;S)w3<}QT?^~M(-kNsZ0YRpkEyrGjg#gBVa1?EiN$3WR`$pn<`!eBxg z`KMgQkGRdN&IIXVwOIoWqrQHlqSw>^@%C4Z#qN?kxL3P7VZZaCaNOWv-n8hZW5Vv7l;c>_7SKom3ew5llL}EN(>-5m7DS&jrZn4miqKU6Mid>9V%<{foEgTmv>$j?=o5 zbD(Mk)kiU#rAJ%Xh%hhma3&z4$Z*qgRJI_tY8 ze+|mLz507GJgGb5cW`(`3Q|>H!DLnxw*cc=rV__L0xbb~RcW9DZ6^nt&a!-G>|FId zUf0goO-;UE-!%$uMMdcvz_hNUmYSA?3rqmk#KQM=ld7sMx~rjKIy2L-P`jdY&fK?( zZuh6CVM1m0TLfp_&ijOfrVl4_3k$Oy#4}ULdw7&M{#Q6KgF5v0qIfdf_v)Ane_26 zCKtp_^Fxue`*}Cd*z-7dgDJB|#8mHM?lb?S+lqJ+8It%Jz)k1d?yI7N-YFbdc5W=0 zel*JJ4Hg#HOh#8LjYp&>1_ZHH4G;UxmsmBez07a3F5I}W_AY2X4b3=7QDSAgB>CTQ z%+ur73x)>q)Kit@%Y*pP09&2B@0O#5C0`0-<}`iq}Q*F-O^1` zM`2#hU!%)>po9zyW94r3WqrJTO^%BS;JSYgeK^AW-h3Vi&4XR{=z8Gr^rYDL^VJjF zYn&?MQZWnaN)D&F3l|qqv5kb~T9)10^MDNIzX%}n5EbQ@N82)b@da4g@d@wob7Qi` z3$m8I1ziazS7c8>))vM;_j6bMg|5ds^xxNtSd-th&FKK?!F4{ z)C<7*>}VG9qhAQDrvhuCUnjuF7eRON$uoK%`#6vnpvkwjZvse~4u5jN6_cB%jyQ&h z#2Wl)>g@ATU_CE*`Sr}vX2{iZfy*O)-4UbI6$XCg7Z*x7id#?0f{(WypZ*gh#{*)j zP$}>2_oFeWrwH0)r}6w*U?4t3(?lrvkDf)MZ4chG@0fZBp?nX%31$6$>G;tfBzTL2 zfkkl~M2HMBy&0-wwx@v0ugP;V15aK7y?}7$6h0EtM3hN9LiE+4#Ecr7CHGsg>XYRJBuMCNg{o)T}-H(=4tAmA4U-B`XOR})b_ z1pPTc7Gepy69F9i?Mn#iXwPz(C4i_mOmC(klQWA0n9B{~BWENKkzE0oms7GU2g@;@ zxkSJP**M=BKOgMet4g>pNG6~rdJ%= z34%;1~Hef__wau=-{TJ>Xtl!#J;540|Z&^kFoJU&^Sl=F*) zK27mI`N`VBi9uIt;bA(l=?o?A7suPK%a5N!N?Dd@NSyH2YcXMv@t0f|EnRI+6IDbJ zmjZ_#aiuh{b;VirBZu}1St}EkF9CgB#A4u)nY=nSkM-g(#j%(v0cS=)FW7mVAd}+= zmm9LkVB%nJx~?~wFOdW}DI-g9?3c$ASAXwSNkwh8HJrhBFBvg6l%Ef?@I(48b4I(( z?&Io&1InCZvosOe7kto46N1)E1)XuPbwWz-R=}Y@=(L85{82ZKef^CQaIQfFPxom$ z=ZHXhL}%(N!f7r$>wdMibzUNM#vRIz$6Qk{h~NASk(LravEww;Kzq8a?Mj@U%R5&t zhrj8+|2^RF3jhhcA|)wRu1^|_iOAm{sa~^U#>`f9hGap`o`iNKT@f$Fo>*x3_~8kE;!zwl;D0=;{@B4G4@{-{qTGKh za{nAfWEKqKmuBe_pr1`kf*8r(e)+}#R_ zTb%QEpT9?UkNa{T?+eAKv3KoH)|zY1xxO8$rXq)pNr8!kgoOQ0URnbQ=?M(+e=WvS z#9yS2qb9^(WETxNNu(dcFL#iTsFB`DOK5#E-Zw)36yH4!u@Zbne~vQ;dLplmFHuZQ zyX35$E1zW2q$4h2_la1G{e7h=GR--yvP^U~H8rNZCGQ!xO{EttK5XYDD>xbYks2St@4}*HV2b5{?KMe=*U}iH6 zt~r3&i-E!aWzfm{)mZ=Abg_^bVKR17o^K+=)n%wpy=`mwEPK{}iqc`ddsaEN5WYL_ zX@nLmUdtl)A!t70TUz?NPp`d#1-!HjeV)(aONdzwL>u)|6*&zgrD&%+73#m2wO0C4 zkuJntv1tmtUHkd((_NSL>lYnFMNANt(!cMWoE(GbOJ3IU(Bw6CDr7SENFnsn%FV`B z_P`Kb>--Z%alJJ*I8H9b9AIlfbA zuJxcV?a9u0tuH~5k)dZ6dHMZvK!3a=jL@KG0l$`fcET5U{(Zy$wVJNMek{YZWtE#E z04k$C2rD&YxD}&egu?dK&R_>E6J@J0S%mlY*ffPL&P5);;e`)s$i8CE3HF}2i?^y} zT>Ac#6xx>gP=7ZSS>;fultM#Te=C!}5BJt(|M! zVivEmV#wzqIENp}c}w(r^Cvx2FF?k!VsgXsrR+R*e~{mya>>ZEeD=86lK~YE)C@S{CVA1H@WlQ{Y2BCW(QCx~O)NP~eQb9pewP0`yDx|LKNnHlR2C7jrps?Lgz40z@s1<*J5fx1 zXWe#q81~D%m}mj1YwQmwg-I7iBrH#`QnU^4 zXhd51=N%22KJx4JXzpNgUt^vm(~2N+Y}p~w6_mC#w8Cx*+bAann7gU3W?O|nCmH?v zR%j)+6eh>Dh|@p2Gm3`DShq@Sk!P}5pXh$m%t0~PbK(W_7CKcZBFSOMDT(+Wn)E=j z^uc}SAZ@+*qlBkYfPLO%fE$g+PqleU3b^w+b&rGmQaX92=wZ^(L>oYB`hL~^N^!GkWPO>_L#PfIxBq7#dpJY~EhWb_?CuS#lmBffpWit1Vp zsQJnT>r7BdHN!2qM5iFkB{W&G1VY~WYKq>i7FSBrKSR0GLlm4ASs18H%2 zof;Z|VyD2C95<7pmF>*@*MmApk@V{xvNWwzkE7O~n%Ul$F=2D7Miq;`f2XHITlX%Y zrKsWs9KS7?V9hvke_{|~%wtt!8f^62?>Vy-HMbrp%!A-nxM)ese$!`5-2PI}Qs0A| zM&wwgf5Yh@5(HEQx$k{iix3}^XydKQY#M0|%g@A@>sxXY0oc!tFJerPPAq?@X=Fvg7* z*1_ob5s%GtoWdchE1)7ZpP33?+`-v67vo7!$_vMmixy~hh~$SRBtRmKVLR0^PyH>F zE-zTVz^P&GjidP2+6LdbrFCj7DBzQQtwdj$`XXhKVfLMl{?R zw}g0C@-~iXGMBg*A!7v2tYD@*ySUfPmg%A~8+zRTmeCTA89m36!9=GPh<}VeSI#}< z+|(W=R%MdCeRCDp(DKLXCTqxHB^OOGu8kAf)HPZFl!yy`zxyjM1W}PcR9jg0Rey>2 zl|+C1IS02AEm{zSb~0EqCrKwI#bcJdzo*G{FL61+(;1wI@$TJ#{wlX^>6>vPQm7mq z%st?Be9|xE&nTzhc;a^7jOf=KXc_AF7mg{lH4 zOppA{g*@Tu1|yGEB!4i;i9M^<;MuqRUcFY-;~Jz^)ewHM4`J0^Ru&(V%_I4*S|$6^ z$F9p+T&6TtMV;822k2>Jj@WJSdZ*&LLb@b2adXz9TG?rqr#_#`sJtE5pT*vl0b-uK zUu(10)#K2*J_C{eCj+6rMwED5_^V?qGJmtLe!`#?qiD1IS4A-n6~RVqedH-?1;HUN zuv2;>0yWf|#uR4_VNq!@@87U&r5Yf*8L<*0U{uy=FYs9tC*9QXS3HYNdk4W3d9rQ_G>GJu>=E>wg z*R=rF9}6PxxsPPN>iw_$>Hq)M|6i~Fr!kQfw=bt!Bi$qR zK?MZ*clyuF`N+;Wg@^V)sIGy2HdPj{QJ0CL%Ne_RA}0DET84LO&tOf_YrfMGn_pe| zEHV^m#M}57Ag^-mXamV`3^w&__*zZ-#C^Rnz9{zfnqut4Tu0k@(!9r1jrnH}$&fSg zO{$J~0~e7IIrJ$^d?VC*Fdplf7@a>zJ_d)TSYlqagIY2uK2MXEjx`-pGa#c1+qmt1 z;ZG6AVT2_K#~_OhtO3U3ZBVPJW5$?8JNrtOHaMNghUkv`qmjCwAu3z9O`T_1gkKb1 zNJO|VX*2-CACIM5h6Y6bEaq7#E}?l2aBNJ=lQjs#zpy_MTO;dDV0<(!nm8&o=Um=C z#4g%Xk47I;HQ~RE9Rb@%gg%SOh%cLEvQE1+Efgp;Ro?j%=u=73a!ipJ2LUfyWD*wM z7++Q5iPt~Den!)sA1JZ88S!7_(!Y;y7%yuEj12oI9*Dvt)p3ju-QBZ`!(4S$isN`8 zUJi?jHB0*B2fTQ`Dtj9LZ&1X9n!FD~r4qD62tOd+NAS{?E*KcS{__Whq32T4s$^Is zLQimz(fq?9O;E#LJOa}~)R`AKaym_x&sC5QvPMOD>+oehm#W;&;fv8ehhFeqF)-O9 zc5C>b?_pp&?&1QhgI*;m=v`+%GvIm2V(K8q*}CioX0M@vL$aoeb~4 zJdzIa$cw7I5sq!Pz)B`+{WgaZ);z3XJGN+6He`<860w2F{f11csF2j&?mG9~B?Yx}@ziLs7)W<3Y;-6j zNXow_3mn^47c?Y~7I?ukF%bjVA>hn>>q7HqoFS)Xq4#(7;Sn*hPexg2Qe+1gPqtQI znMRQFhC)J=`1|){5wcS1dS>5^f`=^Syu5@s8Gt=0E~K1>t*&#{ahKe{m&k*_?@Lo{ z%FVC)9>=Kc7Zw09_dH3hKlbeNRdZGq63=?SJ&3gS3nb?EdvosMw@e+GBvy6jKzy-| zz_^wGp9q_xR>h_!YVVNzI|+?!ewO>SJi<#QZ%3bHH$_vU&EI5;W?0Wsj2c~~`87ec zRjaEE&-XOKhMF}sG)$F6V&c<$ZUQ-YtFN}?+D~=u+*i(mNEkek_XJ1K+B!*O9#`yB z{f)5}K;d!C0H|3Z{45%`RY;d#`XR`EK2vy7q;4e9VoAKIoeJ7NUPh1;m2d;nMI)b+ zC^iPxlm2Ps1lWaBeUP=XArec?q3vYVEqXtj58U7KzMa&2Pe3pcq`WLVjFJd71K@i; zQJdXBD_z0V>UiFSSOyj~CkfLeQc+yu(&lY&jbu6CIfe-%3=Or$CAmG^0srZZZufj` zr$QnCIGRJcaH27i**hl{Jf3H?sGn?nz4O$7tP5Cco6Dz5a}Wh^E_&uSjY1_ss6Mh8h|%Q8!=s?=F-uSAZ_FKRA_R>=b|6JPbHyYu*cn_ zpr(G|5qX$sty0f#yvqi*wLRID6=rE~T!Ht4FEAmlN3uSLjnQEBn5o^<@K%rOog(5T zaVGkV*@vBWr9`o}_7N&u8e|(m7Y4_HMO?driKL07s(8Q8qM59T%Nv8#fh6lV+;af= zn&4l)4WOITa;?T6IAUWY^oN}0>2E$!vs~QWcs!RH6~*k)Z8eB!SJ=$$chJq2PTuVynjFykJ>YHQegM%Of!bXsotz{%-@ zcHy}NvLn{x)C@)Hj?EJG9GQpOuQBn?UrMP*f7A57q90VHVkT_Xs~0>S)&CMc1N1er z!Rgt0WgPUK1TA5%>B{7^zSBJs+EHODkzxM$%Z?HxBuGJzVN=ok=`&Jhu!uy-r)dtQ z+o~*&kbZ~+4Sex)t7)FX*K&<=NW*7}*xr6he@z&?B^ttV!{RfqT<1B;;%K)sp4U55 zt{c9Q7xL^B3c~@)tEfPB&*)5U*v!0NglrH!_59wvUAlr&3Cp};gn`Ii!kRFTrH(8>K&6yu-2FC7qNRH%6*5BG zAbDy~fa@77zle~&O;k&Eaq-~pj< zH+8#?@)^ogh%xBHgMcKkJ&^OSKsqJ!TS(%(@{ru00IZ8+qP&_M_HHM zo(HsI_2Aq{UzQBN*5cr!TSxL6M~ag8tvH)UDJk--O^IBTN5<>dqu=w?+{ueV6U!)B zCr3Sp%am`|M!(ycbKJ7C(tcrlqaIyho7j)0R$4zN)e(!A0^2K{c4SZy)lNY8r~OJg ze9pB0bc^9jcaffHG-Cg6WBNL>R^A)GHx`lb81whL$5q15aF6u_ zGuOqUud$>2?dEGtse^yF286VUJ?ytV;udMwTV^GxJtO0`TV>s~T#x&n(`+^;|K*LK z5vki)(v}g~t6Lrnq+8e7h_GtASaD5cpBOybSNlS~nrRlqe!|ohokF8NE_l=V~q&M4J26jq&!L3KVm8Y%STk=@Idp2 zpzUc!rdymzxKNF@j-(Vg!B zVHLC!oDz$_L|^AX(s^$_lpav=6NPe|Z;nS~hy;GHBUfVh%k2@B8>w!vdj>ef0 zv{@b&UBM;dhpk&+#I__id$tz%+kJmrJ%YZat7gZR!RLE)BUd%A>qc8zu{t)dWiOcw zUXh@zOf=_p(S=ry#VKnNY$8<^{3Jm>QIoVgN2F?*crc%Gr6EbQ?6W_!H4y~n<+S$* z{pBil`XVC10D^IH%RydfldFltUjtw<|H4Nl_44diYtg^ZQ0J5w?7J^ zk&`X^RYR1v*hHK(LgF)CIj*f-s1L>0h#RJJO7*Quo3WlnS>0g!vkqN|_8H%8;M&Ow zI=ryJ<3hyogpaZ&`wQwZ5!*u|e=&UV7B*Y%NrY>8@VsLl(^Y;VK5&cL-kcf>PlFiz)3t(5fZW1H*YffQN&|XO(c@yo34-6nJ?_h20J1GZ)jhlb z8M2|e2IEhIPsCIcH(Zu1D?7)2!{6ru8$(i_IG#0#?PiVSim7GWdG6nGk=GYI&BFFUQ=PBb+mhHnFU7~lCLuPlUa%9Uk@%`2(!+AO;K z)`TnyUfWSObezQ>cfZ(vn%bI17a3EO_?kxMp+uIzIfH+G5&b|WU)aDhOwk(G)LVAG z8yFh!i>h&=CtYa3&3N_*Z=8#pM!>7TcK%xFvNnoMRYW^KkfzL^NPw~nF=Cmv}a6X;-aDmo}PRW4rdW@Lp||n^5Zj1sa6W+8?JL$o(=AAg}BEiq(UyU zUk`UV4ls-V-r4v_oV7vR=zvHLrR8%5jlbsh2(4M9;M|z|3*gtD){V1Cf`OC|tNtS- zMZQNsZYtj>s-Rd4wyX);Kk793oPk`w$5Z~izgYd4_NLxoD!lMLar(;jQw!U!HWMyd zx(gon2R<|dtFO8%9^@C4tkaFJU`ZnZH`6K@*nprnUw_Kw?NMaihi6h}OA=pv^W}ty z6$@G)`pgwn&SQbmH%SgmvDo};a?fUkU~G4=y_Zvt<_PQRB<62iFE=9LKUqATjPFC0^R@Jru-C3re*l8 z)6OLK-MJ({^``)}+kGAYfn-q$cyv{aZw$mstDk!-GTkIy_aRX3vv7D{nAKH^ZG^DL z{Dy8X9|PP<{)S5arr1pl%c5nLF0m4zQ8qP z{q^4t0<)qoegOZ8hOeN&>~3(n9V{Ur_%TZH@E(Y0__7*${hQA{0a5@@mgdh_dPX zJ1EWWUP`5XbVgbrI@I7hCl^5-Jaog+0xEK>LJPLtzy*`vD4PW`r+9{1()s}i%!UGL z^%{&(X6?RfOG#V0*~Cu-QEm}FOJb1@nKO} zQ?zFpXld5=3ww9VBIi|SC1tyZ#Z_=?M24M>E$RERzE-0eRYc1#;c2{$Mn@Uf?2Ir( zz_sGVe#%o}I8w9?!U?sEITb{{Q^jNt>QY4s+Yc^;S1m3mJCUU99r?G(B3fn7Cx`|O zqCa8dv6WShrE5X;(?n=8gNe%eOz28*otcsmE;sH#nJObJ*RPeZNr1)KccI$xy9(?_SbRa6SeZZ$XfAA5CL;u%$IGO4@P$#n>C>y~r$urm zlLxbuWtp$?;BU7oqpnB1C#DABJ5Piwuk#QTx@jflbMeIsx{XB^?cf10)v)>|T>4Vi zR|Ms}N|j;bj$C{mpF;g94THjGPNw{GyM21t!e)d=Bi4F6$xuWlJ5is8zru(kM5btf z|KsN9t2>cw!A1xV)y9xpya7rB#^~aXvxyg!B;z>{N&i|3_-!pp6d41paOx%Hw;@5_ zc)sKh-39obNhl)<{s*2h_g(&*O)k8DJZNBIm>~rRC;&{BgtIu5I-|6n(oWuQQVdAT?FZYbsT@|oRIQHse>7r{bWw!$lTus&=RT+Sw9X~3$F3KMH zPCme9{834odY+&nBMN_-<=rZAan^y$QTDAePONatLs>nzk1W~1T?>>cUvEo8bu)TA zUg)|MDKDB($c|hn&EG&tDW{R!jd`D_dq&dr20S30$4oMT?NxZI*%G0{ADk% zI3AGqp!vNvS-sZ`S$w`$xJTNeH-dKsV)`Pd1p7ag(q`{*Kk(TV-R3|p)b73m=>D)9 zp;bgr6Mm?hdVTt_?9cpsHe;7={P$C=dv<@hJY*P{7WP!id{0V!;kXHlTW?Ksk^yB@ z4r8$k1jHrS=zp2K@BLY)ZEyGQKs=6Q&VO?K7>Gju_drY?5k)K6JMF2%boH7mchYp> z+r5aVg}J;s$~9iwb361u-E4f_;K7#tt)|`W;81sEFX?v7QC+vI1w3zL!8c4dF6P?-ys5I6bq(+RD|!@ zttT4Y_8iocxiZ*V5uQr;Dr#9`DP@ga`ce^-=2o%q#Y+$n?hqNdlhvAOHir}X1@F`o zGI#a0`MQpj5eDy`=v{u_uz^dRmtga8MO>@cB`$1VfzhCeL|J1~AI_}pSdV5=lA2&M z+mHhj*^SL!k`?$66cMi&vXZ-4jH!JR%!9f9`H5$Jn1egILOe@(zVGdC&a^saa9DIK ze&JkkwToWx3}+BW0Q*KWV0}wICVb@OeB^pu=#9Hp$DzRxTJPRr4H8VP?Uw@P zNDW{2a;{~7LP{5^D21pYEi65gi9~$me>Z2x;YbM85nED_QTmBJTh`vR?!c{dS#6Hu zg+rF3Xh=AdJp@Qt=*FZl>7z zPfTgDYc^#y;wLZ?Gy16A^1%KFJ(>NejK zl^RTMG@dW7luPzI6HhUq-RrXH=OZ0h9bn~8|8!u{@x9AtQhQqS1A3lkb2=7~7;+HGArl>Ht zMUkb7^dO&QqeInO8l01pZQD~f_8ApL9tlK7U0f`rQssM&Ex!bc+}Z~9!-?mPXI2?Z#(Qo#Oh&yXDiIVX4(#a8|Iz-O6W~87 z=P#MSPS;-*Q97%#zgQ^WY^%O9W1H)9_Vs8!_t*Uri)r3qu{gNk(q0!;v`#V?>VwsQ zZP&iR%e1&7FeB76N>FH6fUH6_>z0uB!idKecE>d3up0fN-sJ86#cR03wknwXP4~E* zJf6P}s#r;`LTG_esGHW}*mr`piPEFc>j`^IO6x>~roKwfXOM2UTzTf_%#)9Z%vDZM zrtz79Wo?TI-ZqigX!En2T$(DKJ6x3Q#0w)yo!wR;l+%X5*<_~}&8TIj+~lKA2c+R; zPjZfiuvM2{3Ea2bNF05|m6BLNnR!;>Jm@XVCuuo%`~u?p@J2gjyp?a~sIOdM%p~+Y z4o$pQw^glf&^;B~3TMCT(vvWLjx5380Y>qmu0^h8ef)qoWnXjh{jG;emFlLY{g>7! zF5#C0_S76ZB0q=W3O5-W>^}gob#KeY4FUP4u1fp86WHvQdby(N;<#f4j|XfuR5&4< z_Mz6DPE^w8f912f+)lj{q*AoJxs3^;_|SdM=9^6e22Zd(PuXNJ&~_U zwVZwky(R0EWPxYGNfb^kpmHlt2`W^h$Pqd?lH2*9K0tpu2|u|06PJ|m%wjm{{t^HV zW4&-}VA|OND6@<%R%J}NWc9%F9blLul z%>?6F_sU93K&LIGXdIjqVG=Mp{!7e0+NSot_e=Umrv!)ubz~2h^0m*T3QyKapZxRH zHcO(Vbm{2UIi62?CZgOzaMqk!ilB5gKVX}|?#;}?d8+bvCPgj&q{wKUn_78s8KgwFhn)U9m-rK3F^}DFu zsYjozMuUMhkNR%ShGWS)svVIc@vb}w{K{RT2u?6q=CT>n4JFI0e%Uq#)%iWToB8ns z)regP^Zt85^N;xX?%HUodEAGCY?NU>-&xannM|7-Yt^Hb&~Ct0Q`o-dd~#n;S^2lw z+j2Scl~d0-Q4fBa!8vNkA6C0-IM}L!B8*uv#^Q)fKu3=$mjT#TI=iqG8z-)P+Qf5u zCXFar<9(v|_B?{QhAmCL-#+`h{G{GM0qf|Med&YD=h2^ks~E3rEFDC?ElAufa5^e@ z*&aP<#iCLd7x&cc>IWh%oF$TkiQ+?0#N55? zMM_g=Iwo@?=SdJAWPTl+ZEU+y>K-e@<#7pAgz$1;6W?Z)h+y%Neh+E z!>=rH@~2HReXW}d@&|Aup%=bjLB^S~nLc<>pwxTrO&f+ktX8KF1#f7bRh`g6e%m64 z_Gz>&W(PE)(%Hs{6F(NOI`M;MkNe$6sq(3=-2X$Khv&0gQB0zInHFH9NP*nTwR<^FQ>)L`axpR+l`k)AZS*!yX{Ma8P+Mycq^S2a-CGHb&q+HL?9Nq z9m#kfTzte_Hkc--t%FBQ%~+#lV<=)N@nQT8pV#rb$qz21w5-}v)}L6k<<4$A`2VcH z_UuLnMTsr+ISw!)nR7~?%70hdnR~qum<|$S)5ZqBv$0`*i=A*YvKYR1v>W;2$Xh<; zcmz`a_Ey=&kP(l?rBe<80XcK~btt!yLTqRswzV=S-wQQ6p6fP7sRrO_2OKb7* zLaOj=KKil`Xnobmzx$T*P#V0Dtvn^z z9q};iS164$NfLgDcAJDh!ruMY!g+Q|(<_lhysp2MOQ?!ax(xluPtK*qPfuv=)lhC& zx5^LXUw~de`^6z#d<()CUk`Wg*p$7*EjI}JIRY);jTM7FY7>*;oP%m`ehhdf+>Zj= zAZxJjnM!tYqK|RQrX~&a;nw3)gY?-On6skEo#DHq0p{pnw#a?Q^tZ-FeQI66MjsYU z4onxSBZM*Hf(m7*QiWfXWs;Xc7$@$-dl5F+m(Qyx(wFu&MfK*h=v+(Vj0T`Tx#lG7 zpDve6`b29N&IYmUlCO<_2p7kI2a-xH(#Wnmbv8a%nYT;ZWeDJa3reNl7xBD4UXYyJ zITruv(+?pbwDR~1o&D!Mk8YP+Jnyfout1_TWtZgnNp`}|B-)CFzQA7<=AjJTf2Pp0AH2m9V}t6c4{Ydm#2LfCxWN{;+p1iVgi{$N`W+=2;` z1r^_9;j@06YqfBxgs0Yhyb{*bBAQ$5gMWEvo-TSBAD-Cv25pkJZa<41-TjwK^ry1= zZzq~a5?)R%;B8^n4Q%2n)U0s7S3rQw4r`2xucL=%h38%Ps_B!zQ{S*3GPU=Dt}2F3 z3~;kJ5@sKcTrE(X8M53DLQ>p{M+ZenIlKbUft(tQc1p17PkrxmuA0j3r~Ejz=tj?1ae6$?AO$P**kF{M(ve2e;PZx^xkY>8YZ-i~oD-++fo9Zji#!ps7X z;?KgglY!*oDEH*H^dEaX@Eyh=t7C9uVCi{3J;OD+ zq%c#ZK*(hx=8a8GStD6+FaJM6%Q&>lZ;WxV_D=L} z0#+w|Ri0RgWS5&2OYk+zD`2x2hvmo4y}mJVmg-qiL3j}T9)5iIk1|dY=BkHN!OUL* z7iq}1);GOH3Wc~Lco{LB)1KFcfF!BB7BgMf+ao23onl5ea+#tK1?}Q|2qSsw7&M`A zpQcJRNu&u9A@#L9k(pei!RmQpGl}=qW#iNBoLPeJt9ja#6uw?7i3*p4f1BKI7L3_u zGsP_A9uw_UHgLD@2meQ-xoEp+jruj!pA4&i>~ZclRpusn38v^?qVIgS9m)2bs_@$r z^w``u5<-X^|9Dbe_It4L>RB|DVX6gwSN4_pz-Y$lIVCs5Fg$)%;L7$QSt4(W@Ih+o zMe6UBJePQ!{_@F#QEWkX$`D?3chiHds97GOPl?0IzL1-zmk&Rwp`2-?!wuTG?A&{; z*T7cry6$&al30$C8*h!`!QC$UuD3I6U&z6KC4h0y80dvjXyYs?x0h8|D&PG|V7Pb~ zv?9DH;z&$i1wC*Tf}e%cjciUE#}pclJB>Oe?U#FIJ^he?-QLlebh7`AS2qR>gcURp z$yGUivm?*2eEx8$Z^Ls}wkSgU-fQXz;ffSaXX*XSm(HFB-+b{}%^kP<&BeDZ)TbBR zH`)j=?rql=V+HPR;(q-oyDH}43Y|cOA@s1e>z%WDEk1PB!(1#iq1M4FMl^VuHX2@{p;p}v!^2L*)=>(jGyI_a6pVdf(h&_h?H*j z1npbq%gd@;Q7DbQAZI1>1kTtrD_01ZD9YTi$Q)>)8b;ZGscAI zSx>?!S7S)U`tk6gM(DkbDD2w~q={5iQmJIP_z+Mo8u)0vGU8Co75Mq$6H|i6WjY*> zaK5-b&qG*#2)-;HTD39NJzdkv@nb+@-;o^A0YZ&!jv@M)cAdBGWZLRjZnJ>;y8y?z2%(GP+60#R51Kd#s5 zVXYau*COjt=bpw)Vu-U2o`pMYru8g)i<16kEZ~CTG1-BOzlg5hcs7 zqPx*Jm%QW$D@mlV7=MK-R9xgirzP&Cl|#y{acd1wyG&WuOG2NRLr(L38DYy=o64!w z(O|C|`TqDqK{|<+G^$s4d6p`vKVt5AzCTskRU?mN4WgN3{&-r24K>g_9tUo5eK)dj z4a1QOK^nKyNV$i$uBJ^mx23~(vZm`)SnR}7qacH-Z{X3(+Qb)wL2Ic?q$f8~xoht( z3*K`+m(K}wveTCht$gzi3Qb5Aq*b&JVb?8oA*5y1YM-NfSZm&B({3NyzNq#J$gi*f z{k|T?4tr@6#ED%jX!l)dbeni4{ZYdz870W2pO!ai7GmT3@)WsqfXGLEH&?>g9a4Fh zJqc6)GO(JFyO9hXyh5w8H&0Xwj9U#Ld%0V(-pKswaWS4 zJh_eW7Vp-d`E%}{{{Buw3>Z96*RqW0x2_-bHkEvTm%x>#P`jG=CKR(|-a_j5?Cq2Z z9tz;WO~*)fBK$sLSX-@02_xdEzalxo%3*7vG2HP7(UDp61R{M%NNzZd_Q*~JUpWC=EwjRH4 zg&o=lWyRh5z*|u@j19RFCc$CaV=$F<-Aw<;x)^(Vq@sCZVrYIHjdL^N%}dmf@7h8{ zCeWxSe8H-kXe9TbUxXoM5m2;r`6@6zLhU4|FZSa_FP&Psxd-3weLU?(E@0RIXAP0X zqvm%SfwJ)(^G@Zore_ozuil+RY#U1Jd>|#j3E)A+U(L%#u&nsbAJZV1?yug=KXU%A zW(!-(a?*LlS#8x@QTv@$W2%LAa#{|DH^507Uoi#UaLLs&Ya%6TAlFRN2AGh1teqF5}fM9PJpRB@nq(tQ@DKs7-Z)ga?7PG+*y zENd#kneKY|f%ly9g0jJVT*BxfPm-Q|DEjgk~%4B_`A&3FW#F2BY|NcZ{g4Hy6w3SCCZ&r-`JAzV&k_D3NI6a$@I`^#%*moJYY!=mFcG;q8y~h67z?XD@?O zIa+9^nk`?|2jcr?aO8u;R^P}kklj9E$>ozby&g?C-wWD@(Vn(DD7HF-JjE4K;sM>n z$LP<_dmBHOMJQ<^j^ND6VbWe_URg1dlCZ{W@T_KvHSY}M2t4Y)*t)O!y3$8NQA)3R z0Sm{rKA4+Sje))@MkhzPgZoQhCDu#D;nVIY`qDRvUv|<v|oS1Js8ot9}=Y)n%ykwm{R`#=T2c^TSMco0xPmdpr(UOB$){Tl`diTH#9X#)z+L zzPSd89YFz))y~?mzg2jPSQPh06LaaP1QBocBh5P)KQNiY#i}?(u6sx3jF*odK4nZ0 z#IV<^P%Q>^+Cn7Gf`P7Sp?LwL8gE*BnTB63=6y!8oinSMl}31+NB1TItW&M4tRJT# z0t5&fJ)UijNepsoUm$si%;fIq+?dRrT|iSQ39XN%X?jhq5^Ny)>HTZdMJJ; zADR;2pqoS%-y-i~o^xTlM;qQ0R_pk2n8}KYNGx=XS>2sMp)9r9I*E)3;9n{{!t~Nr z|D*xH?Ky+>v-NyNs0LE(3x$TYzh6s6insY0T~AGOi9sjJBtmq+RXtR|H8+Aw;|r!j ztx{$Dx}NiJ+jD09o|_2QSQit4f!fh@2Nszj9zk4BmT$LE7!Mml(p!`VL{3Nk*lgGC z)1UGfxD+73b%ul*@m3r;nOcm420bVYS<*TedX%ai`%S=p8?eKcc(3LsuFxW<;JL70 zo&d-&4to_J6Dvl;7F&ipnCTrWQGRn?tQwqYE-Yv^SJ=B~yKRo>at`j;+r}v=r{2~N zET-zR+O*vs#^~bTxTZ=kLi@>Dc0^?PID@|n3Y)oO8F-MC5NL8A4Hpkc11H=r27=b= zeoSen7qXn}LR)tQtLUwb$^xAMA@>J0yqq{_Z7rtx+D0E7eYnOw`^C-mIW4mVR1e;(>;t1KO{>}9k-yaVK7}KSLY~2W@&g}ZnpB_8xcsEQ} z$e6TE*Bfw^s`AqQDhIq;MYP7`f!ur;O&N$7UN zVb(8M@So}8qy9zFgEch7<+{z_L0@qMV{or;9Sdr2TNp)jbf>uWp4erG;#`4W>!P_a zUIkKRY-gV{?@Wk6De~A~-09^f90a~S2v)CgeAoMehQUa8H2sht4#*XUxaH;s5NH!g z*X{Vax@v$ioBoi=0EKr8A1q#|g6`S#@S@${-GeSxOBE+vYqZ_O-LUjmF7l-HXnxn5 zX46Kaq~6WlepD*IJ2tP%q>_Z+iei9#CwtCzr~Km$l~VACBYl5~ZAs)`%p-@~k?m|P z6ep>vN3rNv-Uf-Ld^uD+V}u=^m+Mw##wUkV#l*npjQ}upUzUB8(*XbgB$Vp}xf0C2 zMZDxepp#??ij~7ubhP-8qB6U%dPA+MSS08zu+e^+bd-TAm5L}9u2qq zg8w)xEHCJWV{(}oG0DgeWpmw?7WZIajR)wdDZwM!cF$t{q*cnv>aHkPD9 zyw+C$FV%+th~`Yujq2XS9jjo&9X$i|h%)KVgYS=5{QAFpUZA<_;w#(6;};Qk$jqN` z#pp3ggAy_EFI{9!p9^AmzYB;a4Ppj%Q16>c&H|pr*eDA0je=-2u{&OUz~hm|NU_Aj zKVQM4+27I2BoB(-k(8%H-EMNFQ5dz7rP*g8#m@4LLw00LuX&j{0qxx8y`9TfG1*<| z`*mA>`;xXRdgq#R$uj!sA43fCgO6Q0r0D$wmz3koVj==rmaQ>VDf?SY{C~wGbSeAK z`>EPSg%(|bUw-jQO+sWXQ`sE}N%!b<33_7hhAR~X;@M0QOlOq@Q8s-e_YDAZqs|$f zD;m|Wn1?Oz0t5c@R1N%=flcIodo3?<9alp}YCQ=I9*lbAUx3=TACf*d)Zp)_=B}0# z+I+|nk_-B?;GRIVv*OM9WXHYe?$l_hTPRtkwWQT zy`AlEwh&$3$~IS^^#QoDiXrrk#XsMr3|W=CBs}R}hZw50@0s&H5e zE%z(48>@x=n^tea5OK{=L@#qb^7?%YuPs|LrNs zcE21AOY9J+LjofNNmc_5RCP2$*ucGe1-1>p|hRjEjbk%D8H)9%%u_U1Huj~V9b;YPU>O7XK=Z#4!$Zr z?`I{jIVfURJt~kzmNldwp>%9anukCKJ9n>=v(ULS@2%Y_50X;3Gk&z$Atw|+X~@w z`!Yjy{KMOpL(T7YPv7=~6wkvkC0*Y8kS;rcc8{jR6BMzFD{Hhng z(Nx$CXh!X*GDHjGpZhBOEj#-pJo*Y8*#q2Ze2O>=mBYnI#WA2*o{(%H{6={w3emqD z?_bwte)v%P!l9uUFyl@X!dL@B#VIpd)9Lc@xcy~s9RG@`*i)ETyqEv})e=2WUL>MA z_|h^Qu3C3S)Nr}pYR{VF1z~`?jUM<;-A2ZDxOVeHA!jhr*@3=wxtpeG{T(iCgdexN zckQdQ^}>FJvNXq_M$S6)<5*&OJmfsAWVyz-}GF>B^E4-yTNIXFqD1cz&uWlCw+=4Tp@lM zTtYm&1(Fim%pcR-S;YOp{m=|{6lTTeI|;d6bcM@V4Lo3`l^cYWS1elO=yf9c^B%zN zd-kqA_Wnr;$@hl{Y_2a~iZE_-JNp7N!xA+BQ0u$+{MYQweDk1`e$7M*p>UfblP%FP zG!N)oJv~l5y@GwdJXUQ@u?Z2YbH^YD;9WA)TtVBIzPLDMy&Nf;zS{f{uV`nUW^ek+ zg#HA;$TgoM_c>N$-!}h=M-cH4O}`iFMz(Fqb7D*ac~bf!(qXepx&xGxE!_6$B)D-r z3nwJ1!U~db{IN`X8PR<~yr(5%uPY8(;ZNd98M;C=2QiJM^UhRz6M%mfILebxn!eGD z`*(S<5RT#nF=^>VI*s%NXT(ud}eLtmMm-dIi|F-sB zVR%pq3}D)IYeo_2TV?U-+repF0aXj!%prR{)VT?E@H*j35N1k(^4^5feQG?ZjMmh^ zFxtwWOq3leg&v@p@Q3}aN^xP7Z0x!bi}u`SY*PYFzL zei?7{Jg-ZS-#uBuMk?s@iW(l7$57{~fn4?M>h71+o`Ni1IBUC(Y*cd(JO*lOntgx< zR)qvgcb%~wnv}Ja-o0Hl+Bf3Sbd$4;iMI^-%s57)6kfG2l(@Nu{g8b&{O@xh}#E|lY+k$uMBm-*quLh?? z?VV2X_^i0PN2fyHm`~|OWIe-m$H-XZ4Q@l78M7JW-UwWk(m?36)bNo&K-HzWRnu=} zvl758smgeq>+(%Zq=P&t^VdDg>sJX>!50WW9c4t_>+O&hP}n>gPWRT@VGP-$`S^Qr z4zcci@{75GL4V}xmxoVJQguqGed19nAW7)y>Zdl%k#VAG$R~X(trLBKbjaOgR4ee| z|FrehVNq@E+khYnDj=mGEiE7*E#2J=Jsu>bOQc)r?hxtj5|B_ta%dQblx~oiL1u>G zxAC0!_bi+mf)h4=@*{LU8y8g9qR-Dzthxq z%M$FUNkpT1z;L9Ai!H>$W@+wI-KaCoA@MjG!Pj~OuZek{*u5~TSc~P)0#;TomFFBK zud2jYxmkJbr z-Q@@?&}h?o)R__KXfv4_4@_%@H(2yldh~|M_=a-S>Nf~mYmu>&810~Wb=Ikri!R!z zCAf+7Ky)h#;P^Nv5}JyE;S_f>`ryL*d>X>3bgr5aGIFV};84NJs*5hlO12BCAg0j|V z<2XG)eVBmMc8K{gVv1vt@FH}P7k$IBq@*l`t>$Zv-@x|UkCr0uIByoK8y(N;ZVR8N z{!?zN0C`rAiAUR%nKm*HrSy6aV^_UFv0E|xnA_s1Jp=7nmsO$A0-n0+n`T5BhWP_1 z0O(QZl-BIml#<@O6SPGR^XiPe7tI#Clx$EbXcLc)&U-6arZcJ6*~NPggDK3C)oT<> z3CTC{;z$WZAbpYSqYN=SQ_03R8B2Gk5KHP|wB$WtW42%3Cq+ zRMP0nz5Uu0%`8^=qoD!Pi($t_b^Lob+Foe6&~Zr-7yn}Lo>B=g)D`E`fBdkdcj7pg zdWInyQ{A~gB^{#1HQ4v^TvJznh#~2G=hCDi^zH#==99E1pxcU2p8j?oFn+OF!87?Y zM%iVssQaTziAObl`0hZJP+$ttc21&QULo3(E7m|XtRR^`$)z3n%B@4iM6ba?z6jhE zLFy~`a;>Ny02C5Z!2?u^y$V37r)i-a+sQ)m0O;7}Wb^J)cxUW2C_K9BvD(qeYNf;^ zHu`DV$d>b;M*kh=E~!MV5_LK(lw$O~d9u&8 z21^8tK#lz}Nu+O|F~`g%M~zQx;F1&2zZ@I!D6+1fe_HLu4_C=a-q&5k|HdWPRyL=R zCj#OR*h_1+z4%^PF(wl4sV8FX8*s`Aa8vWEum-4lW@4@ekQ>|7tt{C_c`RyFup@Lg z0NMBm*XL89ypKv7b5})_B`u{U*1qFBXwD6}r(eFTlzCyd`K~D-sUKa}b&x6KE-+Y% zpn^fUn+801SMmf@49XK=`JTg~kBnN$j;TRbzQ^07t>1RQWsVO9;UNrW+68Jb2Wxc7 zQdvJpHK!!W*k8Sd)6A=opf05`fv5~X8SJh`t4WpH?qn`d+)t1Cy~Vw>|LVuw-&k;` zqs|?p-!?e8N$TY8i2JJLu%O_w(h%>`d4@*s8XzDQ^IrggoPwf-XAObK9u2*&nqa=$ z+2p=rH%AvkQV4DIpFUtTSxUGBI3}RaMs8WhTW6lo0f`ICYM(3brZ@ZKsZrqsi5F!h zj*WfT<4RPC4G(V^=U=MkaexK+CRl}Vee#Q^pU6>wV!Ww*)&wi>G$I=PYQ5(iynQ~obQt-;?1MwHf*m>0p*!L1b~}3 ziNlbo{*^AGEe>6ghBx$;x4fbOEAYJIj5?Mo_iyT~ov#<$c9GglOwmnO8SXb9B(;x< z|8SswQZ848d?B3Ge^0UbX51SxDj0~DM&`@Yy%n+PJeu@g96pYBtT zvE!k~rbx-~djO&^aXH+#>$N^k7hS)Vz5^SOFv7}kYE!3qMe*D>Z#P?_P$nQ*=-f}n z=*IZ@<5sytYPgSzzvNMQ$9Ua*VHtEoI(~zx^bzo$&iK@Z$_77$|K_d`DBN410itL% zTEdO;Mh(b=_?qK}Gkr}29ZYM{A?V#0pN81rPwU?%x7QFgUo+dI&r8xF~Uk<#v&(Zk^Srxv0saMg9|LSqM~qlenV?k@z+y@goHjr z0nFEii@Yl%n>nuhO0IYrSH;peE)Dqdi*&*=T?HSfEqioLIQ$F4d8KWao%hrDYc~25>jcNh_kZ=qEm19i*G02OS~Y4>)EkY8m17;gjqWR%HH~l%i)yjqs79 zL6vYm{c+lko7+r6L_Uk9%|*<)x^Q+0poFk>n*i#>v{#Sg>$Gesv`fc3WN`U? zl?Z?w;xH-{_vEsP-`9t608`o5J95?-V>ydzNY{pB*)JjwSYzOZANpdZtTJr`8@X^9 z`@;tV5HOLkgI$td^XFYa^|Nc6X?bFWmVg?^M-q&8L-Q)>^11P@<$WAa{ zgL#q429NH9HS=cfue87)8-(W?DQMmDin;AORbuoA_`uprK&@mcDmh+KM5 zEL;m4Rf4dKy2lfk#yZDC_$tgDrUk|BSqr9bPOKBMbD+Y`Q-_hqSljd!v-*cOL@UVo=!5x zCwLoFvc<|*3_}({++=gzzsglK3=;fR;1E62-Vh%ks|A zzxkT2V$>P0f1}QWVP4L@K;h_3WxU&}G&jg|PVxia#rX2NY#=)3rI_FwLX(<3V|Bx$LbVUshU>S!vI+T{H2kc^hf0C0P;3)o}F6E%0z#0*zQA z_PDY*7PxVyUF2IGnALk0z%_whp{!*@rOa*sOiK~GP`8@<4-5Soz9ndlJo4x`#WQ0< zRk#_Lb{Ty%mkXQ~XGLbFqMiRt#Up}A^?uUlHcgMuaIbUifQwy}ZX^dii`Bc2 zeh9KniKPi|Z0T+j50cHun2$qAfZ5>L9o-jrx|-Jj*WwrJJ&W2vbS>>NLE&J|AC(kU z>AOkH2LIq@>bwR={8pG2CstBj>xzq4wkT0$NRhRN8#9ro469VBPXcSd9=*&P^P)9n z5c`cT{Qk5%(I2ZHYg!B6YG4$oadL*#(O%vg9RgKdW?$y6h@7i0$&u{m+Vjt7ONL2U3ta5FB`zL^I%GE zrBHl`H~-i>3^MRl56jTWN`!qv?C#F7&Pq5;W4)T)=MaNHYpusUlF0RMf|k zV2tB`InTy_#a+nC_NA7bPab@-D@9H}`cJf~m`#VhEzfCJRrqEd8fw1F?nv9GNzWOH zuH!bF3G*V<3NwjRni-gO+HIZ@=h0Z&z_AgDYmjB%~RwhdJZ;(cm{_x-$CM6x~d2u4j_| z6`GqEfS8ovlbA1jgX}zb_Q1{*AM_J2r`poOC4?u#=XY%4G2tJ^0oc%E13J4sRxci; zm|FO#d{JJ7fNjN@NNQk-UktFUES#s5EOfZ*-a^*x{Z~fz43ju_n0!j#TjitB)vxRN zY@1qd42;)KCZOyzvdFbb_5$6wQ-b1@jn#+Oc7M;+&(nOVjHp#p`YvXSKZ}O&@IT?m zeG0rlH7Gnc(&)aQH7HMPDVJ+~;P)cH<0u5FM!n!Z0I4;VS1TYQ{U+?$Nd$B}YDyN` zV)OBO3e|9;!Mj_d&v4A&nJ2JoEtmVx*DC+mNzjsbtJdARZeNi;9|o+S{lrb%I+PDz zsFxh;ZSxH;%AjPikppPVS+E)CfGN>~aeMuqo5hww>SxAU8Ee!G^|KphPqp{)+XvjZ z)_1iJzLn5}$e_dc6|yf=%W11EZYFpJU9VcctV6QR?wV5n<~sEpbP9qe=YSUd}#aU{kDYibN$r52qY6! zrv3Z&39?nGJ11I!f-Ej9`Yg6fZa~2#$r=s1*cZX&pf{1k)`3@y&9AjZaSP%gYHH z@jM+8$Irk_h}vX2)kHeqBXzY zHC>N?XRPpzK2lu35^{z+eYpgkIlA9IFzux=0@Q-7KF5}(YMl`DMxd`Vn~%ETPDVvE zir$POL5vBmET`}9*tqh~+OV-rcnS5{L9fEUSAWd#zc=t*O5n=t>wPe9)rA4zg0W3Ky7NfOXiBcNqYODXiGF{B+ioG-mdVuyJI+ zp468E=YL#OED%sRkTcWZM8DswLD!#hYBY+nyE<$sJfYrrTF=c5)C$V#2Ep3t)PGa6 zmH$w)f63BWF`a_GbKY;Q?XMs**Wxt=R&YYwqDpv{>LOlLz_zT;&@1c&n>)6#Q389x z3v7|k`r#f$rk-$^WW85o6q8)|D#?MFD4`r?P({kUW4a`X4%qJ>6z?ACS!q_)t72-X zDLOSUKJg~+$sp()NzM-0Y*FIIlIuD`+1B_(LM;Yu4`0h+*sUD08IIV7o5Iv9kGFnl z19K~id+TwaMt#bK=>rmhn{AFs%smkCtI7ZtN~NfLcc@8n3L&mho%W~qpS}tEzrZ&Z z-*G85MIxDojGIvD(KE%p`H~bbAs<1{N#ZMafpb-mO`=CssQ7)G;-%;qF=Qqd-Y(O! zKd&&a;9TO1wg^gqhi!#zm+~}DeVC0G!7(`o_~-o(=e-G%DcLnxV^(~YOF2_CQb+Cx zC>rp~s8IBUJr7|4-EZ1fx%Tgs;sl z!GtYUkaSwrtTXC@S#6-&Vx=q_a72&M7S%yX63reWu}uLfYrzOL0G==rpP1{X$=|SyCL8pgM#5Q=US1s zX_Fb;%50;g6htkoIk(~A<3o&WL5<-veZk&Z8;r718Xsc_*ILR4)3&Dn*u#WMm; z&i~)(B+waa3d)TV&T(%(4O~E7XwQ$UT<1hMe^a4gHi7WpN8ha0<{~sDfgize?dF>h zt;=uL94`q-3^~XxcJrs~!R3_B2Lht3u-f^2YGX-$;~CaTHJ#C4~*&quhO zLLqXWq|+~>$viuY#bmT6C@O)2NgBW(2`gW_vWi7rdOkbOJV}OHa zQW*`IPym$|IFRGhEy$VVP6i&~`%X}?kt5%WGMr_ahP*64S{iez(p~$^mxnZBs(6Ym zPUe;9t>z?3m>2U50RQGxy+LqatnjLOx1j(9>MnCyU_cSRdoAA&qOo(iX7uat3E{vW zn&(IH3)CYTYONRwJQ_Q@!NuHUZ0@mK;`&Ae#723?&P-@6}fovQ|%_jkH)MjrVudN;oKmPx7NSO`1&MPnn0$3wTbSDYU0 z?p?EGy}a>B=gMXHfTQ>x`exdEaFwWLrbPD#xufO!-iz&=?XxIE>(a**>*e<^`Z_?=|N$WZvcz3oj?qUeOzvp@8jQKXdXWIf5I_a|yyA)qW_#GR^k8tj7aT)RcBAb6! zZ+!OSrmOSSj)eF5AimWP`0G=$Wxwu#v#*1EJL~*8ix*-`7wMK;euFUWWxtWiU8er& zNV_G-R&tu>Xo*n$CUu4_c=@`_SZ&Wz_8r{ z#5z28&p>rC2Bu@AFa9yUXZ$X5gl8@T%<DZqKfer$=6uk2FtlC0jdGM*?qU z#KYiMeR*d7Uq#H3a5&#laJPT^Tc<6`oDpim&C|Ki?gwbm;Q1%5UFPkbPl>itt6}GL z*7lkZL1!3)O8mxPS2F{Zso_F&_9clDg20hS!Q8nfiN@@zb_wB^?|mw{X|B-J*_Fs+ zUa)u7*0L2%2e;i#?`AJGKb;Egs&szb+If!rVZX-u-o4GQen+7+Qb6+*-0q&A;}3nA zRP+3s^wF-Zt6)C=l~aI)6P)y%uF3tff^ICmF#wS;wVxfkH)8!dmNigaqvM!+i}lw{ ziOnNOH__!CLDZEgb1iz4Wh<|%!KGdewcASMW4468{Dg1*0IhjrsxEsBV|^*^eLAr` zAQXUk7n*%Zw{YlJ8t_8^lWAl21pILo8Tw>o{_dD*U}9r@k5r>`SF)r>UdXwD#42Dl z3G1s+-P$&Pmu{?6FKhQ6eMu*`fTbf__l&Eir_0B5znHPyyS(jxq;@!;<@}h1s*ZsS zkyy`A#o2d#6<&J}GfwQ@hIJG!9r_~%aJ3(J8!o1OCKVTyvrS$HgONp zDk0g&BPQ50p76DJ)`1Xy$jRu-ZqM1JwnH3#=!<}&Yn-zb&ini z(Y;^q0`3XD6B!i=+;rH2gdH4pu0%&agd%-*SPW^$@?qzbFF>5#y{bay-Dw#GhZH{Ob@=h{cQd;+=oo- z2ND1Nxb+oev!ge6;T^21HeJH>Wx4G5emckNjJ3x2rlmTe?yb)>Ac*K2`@!Hf?}nq~ zCdpo{1+=RrnHNqAC6iOlLZ3icqr(_wR_h1GBPY}OJ?e{9?i%S_E8i}y_>d^{&o%Qm zR~+T#*=&`0e&gN^bj|T_c*C6!&rpBmPwl^BHk*D<)>@Du%tn9RV<8YW-qJ>0%N8Ld zd${MON7ey6*6%jkPwsck;@~nx?i<djv{=X}?1B zrGmXgstzt>PB;G^yFU{chv3ujUyfsxyNSa+EN9R5_N*SR$~4kZhypSu2jLvK8D48_ov-e;I0Wfq^!DTU{=0CEM!GZ) z?NhCVD2vq(LAVY)3Ja!x>-jq2ao?aGx4TYw%pZC>9?b&IMhl3P(QNE8Ph;SIhGNLC z4A(r$7507o4UJ*RJ0aNiGjFLe&z({oTfSN=rwFLYtw4-Q?0i)d&24b;d@>j_&D zIl{1d9kHlmyo1#S+O!8KceBL@AA}7JEwg%W@9!O?jyCVjor!O$HR#(-G)<>v@|Ozf>BrW2Mpg|MOF28|Rhf0< zI}U!sx!9U~3hcNtfXkc_mX>6@)E6>cDU3arTI|My5`qNqhDRg?Fd}6$v-1(%LTai_ zgLrDn$rv?v;o$c)VwND2c>@2jHp;mQkHxbyT^Mp~lm;a9ZI-!#I~_+%R5A?_=g4^o zOJKb+rDEA>Vp}JKL*B>4$N&}_NY}yrNXH@13F}wLH&0`d0W4k4gB-TX+|5GmS&uxT znN1nt|0K25-rV6Pz6c#iBE(o$PK2<#8L6v7F8OWLn86NrlZHL3&06h1US8({!-9nW zjwTFgt*)Yw4`k2TqG&k=@JG4^RZtiVHASK(RvQMPO&EQwBT0`z^j+XC!5_t&{Y&L9#D}{5>d2H*DZMU zemPY!-~%b*W{*`azFcZuoc1BeM* zL~x$N&xD38LO$C1^tHc5-tIqL^!Ex7%(d!Lot>?}3}vQ94oYU@yxcJ^$Naf_5PnmMFbn*d|!!yMkB<9XTGG`a7>KFIysE|43qqYCEP*U+OpYH)@Zfk&2v+- zu$%vhLkhpAu{iX*R^H0(?(UYqVyrR!ZOzq8r6t9T_BlU!odU3;UBj5s4}f3A8%h}TYnxBDUX_7%Y0eM8dEVAXE7ylb>lBV z&fQ(xqoIRCnE5u3>HT6>6|=TafhOi6_pTxxkH&??`|mo-akTC{ieIn&#+)zzqx)J4&jJxS<*6_od8N-o znGZl%OG&Jf0q$c|tkbG*0NPTFdnhKVYsN*M8B^xbzQ7s1fELXI@RS;E;^GZkYw@oS z_I8x_WLDzX9KL~`w7`m3s&!gpt7tRnP?-~jmG%lMEXj!;qWb&9k0=walZ+`htN5bt zdZFg)gSn|j)?8?~SV7%GaQ)+TaJ`PAB=_wL3&Fn%yJ*BIqeK{Q0~wl~g)~Zx<@SY) zI0rZ7s*4`uu2oPlkQ#_vi}EQiH~5MtPvxcubq3<0#x^RVd7<*`smXdsS(3~4*$Y;y zLA+T^TpQFSb2+zKsdUS(T9eyW+?;ZNvbMy048)$Q$_~%h*@a{LFoRav;po#SsG@v2m-2g}2KwFIvZ@{5eB<8HuG{7y|d)zwvtje~*Ic z)}3>`z`EtF4%GTRoLhz*zqJii=+gvv+A$J$?;q6yJ*(58M=9;yEm+@DnXYzJPfp9iv(7dF)LiIJ zw>ZAgpcVauqB>wwqWjosk2?sDAxO#s&lW*S?Zf@(-@*07_KaIELoRYc0ZW82NHx;E z<#Yer)Yta@st{x?DVV=5LmbhJ1qrC?9qwcJ?FZK%zqq6tk6PK*0Kvq7iKNT;SpL323sY zw}16|_5i3;zw_C`Lg){@d3S zlOXHv1(Fv5fb#$ibU$E@`}3sK4_(0II@s8_1@wRh!rAH-{&DZuSu9Ex|0K~gZ_wqb>_uYay;lL6siF>T0?Gf4UYrfELXX7NWV5ugFu zpTgwWxH7;+qS&m+KOVVy-HCT>Zg+2e1p~Uk@NN?P(SRQy^_(O*pID0iDKN?&+lLSS zH70!=Ak{U70~#iU2D-CGsatdF&qXQVb}ZM2RM@zFq(7(kH65Is)hG$3PReDjW=jUI z9<)~=6@~J$b~SB0ij9{sxo-uR4Mu=}tooxvZ9s7`ztm!}aO4T!yh1b%n{eN+e*d(Q z{IR&Vy>814+8a$Lq*1m2s)wMBIQ^aa$&>TAeB>WR0IK>&Oi&-tMGE|P#SB3`AA(A) z1BZIE74}t&y`q)E?|V9$OchyKufN^oTAnEgzK5-1BBRN_COUBs=!CS$a&Am40OOhD`sC91Eupc#yL2 zEAK6OEhpg7uC^fL<}A6&I?lw2SoN6%^Q+*9XRSg4qdez$e3Jk5TmZrNKB|1s5NI+b zND_iMNpnvz7>gJ%AmP9y_wI&v%HX9c{JB)RJU$M#q_%0T#P zx}CMX#w}p_VbTcRWrj*F01&V8yfmMsl{g+Ug4!aT@+mY`yz8HEAbz0Hv*6{_={B^p zeeFepgeldq)`#KEe=^bE&P8km44F+azWKjg<+sy3jlqIQ*Z~>%PfvhvQX@c&yrj_L z#{V}){!v7hzuO;`{QWIB<2s8y`Ctm<|349N{p}fQ{eJ{w;qU@69d9Jn_UOOd=J$Yr s7r_5EhHJqBgnMI)Bpeg diff --git a/doc/graph/fusion_patterns/sdpa.md b/doc/graph/fusion_patterns/sdpa.md index bead0f80974..0f2e889789e 100644 --- a/doc/graph/fusion_patterns/sdpa.md +++ b/doc/graph/fusion_patterns/sdpa.md @@ -44,22 +44,36 @@ optional. MatMul with a scaling factor. It can be constructed by [Multiply](@ref dev_guide_op_multiply) or [Divide](@ref dev_guide_op_divide) operation in Graph API. The scaling factor is given by users as an input of SDPA. \f$\sqrt{d_k}\f$ in the formula - is not considered as part of the SDPA pattern as it is constant. + is not considered as a part of the SDPA pattern because it is a constant. 3. The Mask node is optional and is used to apply an attention mask to the - output of the previous Scale node. It can be constructed by [Add](@ref dev_guide_op_add) + output of the previous Scale node. There are two types of masks that can + be applied: + + 1. Explicit user-generated mask: You can explicitly create a mask tensor + and pass it to the library for the computation of SDPA. In this case, mask + can be constructed by [Add](@ref dev_guide_op_add) or [Select](@ref dev_guide_op_select) operation in Graph API for different - mask policies (eg. causal mask or padding mask). When Add operation is used - to apply the mask, the input mask is usually an upper triangular matrix with - all the elements above the diagonal filled with `-inf` and zeroes elsewhere. - The `-inf` entries will become zero probability after Softmax is applied in - the next step. Alternately, a Select operation may be used. In this case, the - input is a boolean tensor (for example, with `true` on and below the - diagonal, and `false` above the diagonal). A `false` element in the mask - forces the corresponding element of the scaled output to `-inf`, while a - `true` element leaves it unchanged. + mask policies (for example, causal mask or padding mask). When the + Add operation is used to apply the mask, the input mask is usually an upper + triangular matrix with all the elements above the diagonal filled with + `-inf` and zeroes elsewhere. The `-inf` entries will become zero probability + after Softmax is applied in the next step. + Alternatively, a Select operation may be used. In this case, the + input is a boolean tensor (for example, with the boolean value set to `true` + on and below the diagonal, and `false` above the diagonal). + A `false` element in the mask forces the corresponding element of the scaled + output to `-inf`, while a `true` element leaves it unchanged. ![SDPA-mask-1](images/sdpa-mask-1.png) ![SDPA-mask-2](images/sdpa-mask-2.png) + 2. Implicit library-generated mask: You can use the operations in the library + to generate a mask by constructing a subgraph. Currently, Graph API supports + generating an implicit causal mask (top-left aligned) using operations of + [GenIndex](@ref dev_guide_op_genindex), [GreaterEqual](@ref dev_guide_op_greaterequal) + and [Select](@ref dev_guide_op_select). + + ![SDPA-mask-3](images/sdpa-mask-3.png) + 4. The SoftMax operation takes the masked output and transforms it into probabilities between 0 and 1. See [SoftMax](@ref dev_guide_op_softmax) operation in Graph API. @@ -97,7 +111,8 @@ platforms follow the general description in @ref dev_guide_data_types. softmax primitives. The reference implementation requires memory to store the intermediate results of the dot products between Query and Key which takes \f$O(S^2)\f$ memory. It may lead to out-of-memory error when computing long - sequence length input on platforms with limited memory. + sequence length input on platforms with limited memory. For an implicit + causal mask, the reference implementation is only available on CPU. 2. The SDPA patterns functionally supports all input shapes meeting the shape requirements of each operation in the graph. For example, Add, Multiply, Divide, and Select operations require the input tensors to have the same From 9208bcf782ef56b16046b3401a9ac1441942f2f3 Mon Sep 17 00:00:00 2001 From: "Bao, Yixin" Date: Tue, 7 Jan 2025 22:08:08 -0800 Subject: [PATCH 16/16] benchdnn: graph: skip unimplemented GenIndex op on gpu --- tests/benchdnn/graph/graph.cpp | 20 +++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/tests/benchdnn/graph/graph.cpp b/tests/benchdnn/graph/graph.cpp index 7bb314586df..b3e3fd3b2ce 100644 --- a/tests/benchdnn/graph/graph.cpp +++ b/tests/benchdnn/graph/graph.cpp @@ -404,7 +404,10 @@ void skip_unimplemented_ops(const dnnl::graph::partition &partition, // A list of ops that don't have DNNL backend support so far. static const std::vector unimplemented_ops {"Pow"}; - + // A list of ops that don't have DNNL backend support so far on GPU. + static const std::vector unimplemented_ops_gpu {"GenIndex"}; + const auto &eng = get_graph_engine(); + bool is_gpu = eng.get_kind() == dnnl::engine::kind::gpu; // For an unsupported partition, retrieve all operation IDs, find a // correspondent operation kind in a deserialized_graph and match it against // a list of known unsupported ops. @@ -423,6 +426,21 @@ void skip_unimplemented_ops(const dnnl::graph::partition &partition, res->reason = skip_reason::case_not_supported; return; } + + if (is_gpu) { + const bool has_unimplemented_op_gpu = std::any_of( + unimplemented_ops_gpu.begin(), unimplemented_ops_gpu.end(), + [&dg_op_kind](const std::string &kind) { + return dg_op_kind == kind; + }); + if (has_unimplemented_op_gpu) { + BENCHDNN_PRINT(2, "[INFO]: Unimplemented op on GPU: %s.\n", + dg_op_kind.c_str()); + res->state = SKIPPED; + res->reason = skip_reason::case_not_supported; + return; + } + } } }