From 9b183d909eb3c8dbbb32d6da79a7135bdaa50017 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Tue, 19 Nov 2019 20:18:09 -0600 Subject: [PATCH] code object tracking and v3 code object support Change-Id: I081ada13f6364ea4401a97a485dedfa9bf8a45fe [ROCm/rocprofiler commit: 60043d198b399eda45336814ea3fe1dd6dd22a84] --- projects/rocprofiler/bin/rpl_run.sh | 5 + projects/rocprofiler/inc/rocprofiler.h | 3 +- .../rocprofiler/src/core/intercept_queue.h | 25 ++- projects/rocprofiler/src/core/rocprofiler.cpp | 12 +- .../rocprofiler/src/util/hsa_rsrc_factory.cpp | 54 ++++++ .../rocprofiler/src/util/hsa_rsrc_factory.h | 14 ++ projects/rocprofiler/test/CMakeLists.txt | 3 + .../test/app/intercept_test_stand.cpp | 2 +- .../rocprofiler/test/ocl/SimpleConvolution | Bin 0 -> 132704 bytes .../test/ocl/SimpleConvolution_Kernels.cl | 175 ++++++++++++++++++ projects/rocprofiler/test/run.sh | 8 +- projects/rocprofiler/test/tool/tool.cpp | 4 + 12 files changed, 293 insertions(+), 12 deletions(-) create mode 100755 projects/rocprofiler/test/ocl/SimpleConvolution create mode 100644 projects/rocprofiler/test/ocl/SimpleConvolution_Kernels.cl diff --git a/projects/rocprofiler/bin/rpl_run.sh b/projects/rocprofiler/bin/rpl_run.sh index cf0fae5914..0556265ef7 100755 --- a/projects/rocprofiler/bin/rpl_run.sh +++ b/projects/rocprofiler/bin/rpl_run.sh @@ -179,6 +179,7 @@ usage() { echo "" echo " --trace-period - to enable trace with initial delay, with periodic sample length and rate" echo " Supported time formats: " + echo " --obj-tracking - to turn on/off kernels code objects tracking [off]" echo "" echo "Configuration file:" echo " You can set your parameters defaults preferences in the configuration file 'rpl_rc.xml'. The search path sequence: .:${HOME}:" @@ -394,6 +395,10 @@ while [ 1 ] ; do convert_time_val period_rate errck "Option '$ARG_IN', rate value" export ROCP_CTRL_RATE="$period_delay:$period_len:$period_rate" + elif [ "$1" = "--obj-tracking" ] ; then + if [ "$2" = "on" ] ; then + export ROCP_OBJ_TRACKING=1 + fi elif [ "$1" = "--verbose" ] ; then ARG_VAL=0 export ROCP_VERBOSE_MODE=1 diff --git a/projects/rocprofiler/inc/rocprofiler.h b/projects/rocprofiler/inc/rocprofiler.h index fc723c5d32..a0d186526f 100644 --- a/projects/rocprofiler/inc/rocprofiler.h +++ b/projects/rocprofiler/inc/rocprofiler.h @@ -64,6 +64,7 @@ uint32_t rocprofiler_version_minor(); typedef struct { uint32_t intercept_mode; + uint32_t code_obj_tracking; uint32_t memcopy_tracking; uint32_t trace_size; uint32_t trace_local; @@ -222,7 +223,7 @@ typedef struct { hsa_signal_t completion_signal; // Completion signal const hsa_kernel_dispatch_packet_t* packet; // HSA dispatch packet const char* kernel_name; // Kernel name - uint64_t kernel_object; // Kernel object pointer + uint64_t kernel_object; // Kernel object address const amd_kernel_code_t* kernel_code; // Kernel code pointer int64_t thread_id; // Thread id const rocprofiler_dispatch_record_t* record; // Dispatch record diff --git a/projects/rocprofiler/src/core/intercept_queue.h b/projects/rocprofiler/src/core/intercept_queue.h index 5a6234abcd..f639b3e525 100644 --- a/projects/rocprofiler/src/core/intercept_queue.h +++ b/projects/rocprofiler/src/core/intercept_queue.h @@ -148,9 +148,20 @@ class InterceptQueue { } // Prepareing dispatch callback data - const amd_kernel_code_t* kernel_code = GetKernelCode(dispatch_packet); - const uint64_t kernel_symbol = kernel_code->runtime_loader_kernel_symbol; - const char* kernel_name = GetKernelName(kernel_symbol); + uint64_t kernel_object = dispatch_packet->kernel_object; + const amd_kernel_code_t* kernel_code = GetKernelCode(kernel_object); + + const uint16_t kernel_object_flag = *((uint64_t*)kernel_code + 1); + if (kernel_object_flag == 0) { + if (!util::HsaRsrcFactory::IsExecutableTracking()) { + fprintf(stderr, "Error: V3 code object detected - code objects tracking should be enabled\n"); + abort(); + } + } + const char* kernel_name = (util::HsaRsrcFactory::IsExecutableTracking()) ? + util::HsaRsrcFactory::GetKernelName(kernel_object) : + GetKernelName(kernel_code->runtime_loader_kernel_symbol); + rocprofiler_callback_data_t data = {obj->agent_info_->dev_id, obj->agent_info_->dev_index, obj->queue_, @@ -159,7 +170,7 @@ class InterceptQueue { completion_signal, dispatch_packet, kernel_name, - kernel_symbol, + kernel_object, kernel_code, syscall(__NR_gettid), (tracker_entry) ? tracker_entry->record : NULL}; @@ -243,14 +254,14 @@ class InterceptQueue { return static_cast((*header >> HSA_PACKET_HEADER_TYPE) & header_type_mask); } - static const amd_kernel_code_t* GetKernelCode(const hsa_kernel_dispatch_packet_t* dispatch_packet) { + static const amd_kernel_code_t* GetKernelCode(uint64_t kernel_object) { const amd_kernel_code_t* kernel_code = NULL; hsa_status_t status = util::HsaRsrcFactory::Instance().LoaderApi()->hsa_ven_amd_loader_query_host_address( - reinterpret_cast(dispatch_packet->kernel_object), + reinterpret_cast(kernel_object), reinterpret_cast(&kernel_code)); if (HSA_STATUS_SUCCESS != status) { - kernel_code = reinterpret_cast(dispatch_packet->kernel_object); + kernel_code = reinterpret_cast(kernel_object); } return kernel_code; } diff --git a/projects/rocprofiler/src/core/rocprofiler.cpp b/projects/rocprofiler/src/core/rocprofiler.cpp index cbfbc432de..61fd4619ae 100644 --- a/projects/rocprofiler/src/core/rocprofiler.cpp +++ b/projects/rocprofiler/src/core/rocprofiler.cpp @@ -144,7 +144,8 @@ void * tool_handle = NULL; // Return true if intercepting mode is enabled enum { DISPATCH_INTERCEPT_MODE = 0x1, - MEMCOPY_INTERCEPT_MODE = 0x2 + CODE_OBJ_TRACKING_MODE = 0x2, + MEMCOPY_INTERCEPT_MODE = 0x4, }; uint32_t LoadTool() { uint32_t intercept_mode = 0; @@ -188,6 +189,7 @@ uint32_t LoadTool() { util::HsaRsrcFactory::SetTimeoutNs(settings.timeout); InterceptQueue::TrackerOn(settings.timestamp_on != 0); if (settings.intercept_mode != 0) intercept_mode = DISPATCH_INTERCEPT_MODE; + if (settings.code_obj_tracking) intercept_mode |= CODE_OBJ_TRACKING_MODE; if (settings.memcopy_tracking) intercept_mode |= MEMCOPY_INTERCEPT_MODE; } @@ -432,7 +434,13 @@ PUBLIC_API bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t fa // Loading a tool lib and setting of intercept mode const uint32_t intercept_mode_mask = rocprofiler::LoadTool(); - if (intercept_mode_mask & rocprofiler::DISPATCH_INTERCEPT_MODE) intercept_mode = true; + if (intercept_mode_mask & rocprofiler::DISPATCH_INTERCEPT_MODE) { + intercept_mode = true; + } + if (intercept_mode_mask & rocprofiler::CODE_OBJ_TRACKING_MODE) { + if (intercept_mode == false) EXC_RAISING(HSA_STATUS_ERROR, "code objects tracking without intercept mode enabled"); + rocprofiler::util::HsaRsrcFactory::EnableExecutableTracking(table); + } if (intercept_mode_mask & rocprofiler::MEMCOPY_INTERCEPT_MODE) { hsa_status_t status = hsa_amd_profiling_async_copy_enable(true); if (status != HSA_STATUS_SUCCESS) EXC_ABORT(status, "hsa_amd_profiling_async_copy_enable"); diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp b/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp index a47062ddd4..4c63b8abd7 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp +++ b/projects/rocprofiler/src/util/hsa_rsrc_factory.cpp @@ -193,6 +193,8 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) { hsa_api_.hsa_executable_load_agent_code_object = table->core_->hsa_executable_load_agent_code_object_fn; hsa_api_.hsa_executable_freeze = table->core_->hsa_executable_freeze_fn; hsa_api_.hsa_executable_get_symbol = table->core_->hsa_executable_get_symbol_fn; + hsa_api_.hsa_executable_symbol_get_info = table->core_->hsa_executable_symbol_get_info_fn; + hsa_api_.hsa_executable_iterate_symbols = table->core_->hsa_executable_iterate_symbols_fn; hsa_api_.hsa_system_get_info = table->core_->hsa_system_get_info_fn; hsa_api_.hsa_system_get_major_extension_table = table->core_->hsa_system_get_major_extension_table_fn; @@ -231,6 +233,8 @@ void HsaRsrcFactory::InitHsaApiTable(HsaApiTable* table) { hsa_api_.hsa_executable_load_agent_code_object = hsa_executable_load_agent_code_object; hsa_api_.hsa_executable_freeze = hsa_executable_freeze; hsa_api_.hsa_executable_get_symbol = hsa_executable_get_symbol; + hsa_api_.hsa_executable_symbol_get_info = hsa_executable_symbol_get_info; + hsa_api_.hsa_executable_iterate_symbols = hsa_executable_iterate_symbols; hsa_api_.hsa_system_get_info = hsa_system_get_info; hsa_api_.hsa_system_get_major_extension_table = hsa_system_get_major_extension_table; @@ -681,10 +685,60 @@ uint64_t HsaRsrcFactory::Submit(hsa_queue_t* queue, const void* packet, size_t s return write_idx; } +const char* HsaRsrcFactory::GetKernelName(uint64_t addr) { + std::lock_guard lck(mutex_); + const auto it = symbols_map_->find(addr); + if (it == symbols_map_->end()) { + fprintf(stderr, "HsaRsrcFactory::kernel addr (0x%lx) is not found\n", addr); + abort(); + } + return strdup(it->second); +} + +void HsaRsrcFactory::EnableExecutableTracking(HsaApiTable* table) { + std::lock_guard lck(mutex_); + executable_tracking_on_ = true; + table->core_->hsa_executable_freeze_fn = hsa_executable_freeze_interceptor; +} + +hsa_status_t HsaRsrcFactory::executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data) { + hsa_symbol_kind_t value = (hsa_symbol_kind_t)0; + hsa_status_t status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &value); + CHECK_STATUS("Error in getting symbol info", status); + if (value == HSA_SYMBOL_KIND_KERNEL) { + uint64_t addr = 0; + uint32_t len = 0; + status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &addr); + CHECK_STATUS("Error in getting kernel object", status); + status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &len); + CHECK_STATUS("Error in getting name len", status); + char *name = new char[len + 1]; + status = hsa_api_.hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_NAME, name); + CHECK_STATUS("Error in getting kernel name", status); + name[len] = 0; + auto ret = symbols_map_->insert({addr, name}); + if (ret.second == false) { + delete[] ret.first->second; + ret.first->second = name; + } + } + return HSA_STATUS_SUCCESS; +} + +hsa_status_t HsaRsrcFactory::hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options) { + std::lock_guard lck(mutex_); + if (symbols_map_ == NULL) symbols_map_ = new symbols_map_t; + hsa_status_t status = hsa_api_.hsa_executable_iterate_symbols(executable, executable_symbols_cb, NULL); + CHECK_STATUS("Error in iterating executable symbols", status); + return hsa_api_.hsa_executable_freeze(executable, options);; +} + std::atomic HsaRsrcFactory::instance_{}; HsaRsrcFactory::mutex_t HsaRsrcFactory::mutex_; HsaRsrcFactory::timestamp_t HsaRsrcFactory::timeout_ns_ = HsaTimer::TIMESTAMP_MAX; hsa_pfn_t HsaRsrcFactory::hsa_api_{}; +bool HsaRsrcFactory::executable_tracking_on_ = false; +HsaRsrcFactory::symbols_map_t* HsaRsrcFactory::symbols_map_ = NULL; } // namespace util } // namespace rocprofiler diff --git a/projects/rocprofiler/src/util/hsa_rsrc_factory.h b/projects/rocprofiler/src/util/hsa_rsrc_factory.h index af03189585..06cae59322 100644 --- a/projects/rocprofiler/src/util/hsa_rsrc_factory.h +++ b/projects/rocprofiler/src/util/hsa_rsrc_factory.h @@ -95,6 +95,8 @@ struct hsa_pfn_t { decltype(hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object; decltype(hsa_executable_freeze)* hsa_executable_freeze; decltype(hsa_executable_get_symbol)* hsa_executable_get_symbol; + decltype(hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info; + decltype(hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols; decltype(hsa_system_get_info)* hsa_system_get_info; decltype(hsa_system_get_major_extension_table)* hsa_system_get_major_extension_table; @@ -323,6 +325,11 @@ class HsaRsrcFactory { static uint64_t Submit(hsa_queue_t* queue, const void* packet); static uint64_t Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes); + // Enable executables loading tracking + static bool IsExecutableTracking() { return executable_tracking_on_; } + static void EnableExecutableTracking(HsaApiTable* table); + static const char* GetKernelName(uint64_t addr); + // Initialize HSA API table void static InitHsaApiTable(HsaApiTable* table); static const hsa_pfn_t* HsaApi() { return &hsa_api_; } @@ -387,6 +394,13 @@ class HsaRsrcFactory { // System agents map std::map agent_map_; + // Executables loading tracking + typedef std::map symbols_map_t; + static symbols_map_t* symbols_map_; + static bool executable_tracking_on_; + static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable, const char *options); + static hsa_status_t executable_symbols_cb(hsa_executable_t exec, hsa_executable_symbol_t symbol, void *data); + // HSA runtime API table static hsa_pfn_t hsa_api_; diff --git a/projects/rocprofiler/test/CMakeLists.txt b/projects/rocprofiler/test/CMakeLists.txt index 970973d39f..72708d7d30 100644 --- a/projects/rocprofiler/test/CMakeLists.txt +++ b/projects/rocprofiler/test/CMakeLists.txt @@ -94,3 +94,6 @@ add_custom_target( mbench COMMAND sh -xc "cp -r ${TEST_DIR}/memory_validation ${PROJECT_BINARY_DIR}/test/." COMMAND make -C "${PROJECT_BINARY_DIR}/test/memory_validation" ) + +## Copy OCL test +execute_process ( COMMAND sh -xc "cp -r ${TEST_DIR}/ocl ${PROJECT_BINARY_DIR}/test" ) diff --git a/projects/rocprofiler/test/app/intercept_test_stand.cpp b/projects/rocprofiler/test/app/intercept_test_stand.cpp index de3dbdaf72..4f46f65efc 100644 --- a/projects/rocprofiler/test/app/intercept_test_stand.cpp +++ b/projects/rocprofiler/test/app/intercept_test_stand.cpp @@ -73,7 +73,7 @@ void dump_context_entry(context_entry_t* entry) { const rocprofiler_dispatch_record_t* record = entry->data.record; fflush(stdout); - fprintf(stdout, "kernel symbol(0x%lx) name(\"%s\")", entry->data.kernel_object, kernel_name.c_str()); + fprintf(stdout, "kernel-object(0x%lx) name(\"%s\")", entry->data.kernel_object, kernel_name.c_str()); if (record) fprintf(stdout, ", gpu-id(%u), time(%lu,%lu,%lu,%lu)", HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index, record->dispatch, diff --git a/projects/rocprofiler/test/ocl/SimpleConvolution b/projects/rocprofiler/test/ocl/SimpleConvolution new file mode 100755 index 0000000000000000000000000000000000000000..be4c1332a279c55c80b0461daa2fb94f1f7344e7 GIT binary patch literal 132704 zcmb?^349bq_J0S0Mnoqny77+g=z=0997;Ua2}z)%K_kbSRWS)UNFx`{^=1-!`r`>m>;>2zjTcmKc1CsX}iRlR!e z)vH&ps;hhE=j4s+myn=YpMKgOHA0PJJpv`aHLYl0w2Vg^qaCDqv{dbQ?EsV-fPbr; zM_pOzURBLXlTenQ{`gPC|CoM+n0lB*(&GfAX;x~r&Xny%KOS`{)8rRqNh{T~RFO|C zdEK;1=5Jb6EbvzP0WJu6LC{9CCdA7$hJWR*X;p9c_O4-+k0 zd`DB3Q7G|u;GBFEYh+*GjTBs z|J4d{jY?~gdhuV6|JnF&!2h-Qw?5b5;(Gkg#s3ZXZ^Zus{NIRw>*L48P557k|A0aU zk%rXuZ7OX>dI$dRQm~cWgNu9dzgQveL;6?z-;e(#_j6EKH!{xjk>7&?{}QH_OA;rJu~H> z?-tBY>|Fbgi#}50y@vG$?iFz7scnF=^fH9aleiTuti_JAb(OtEcYXQ8oY1N6t(+GkeU9kAKv; z`>21eI-~BYL(aMP;D2veT;1|j>Es!|yjnWrpsaTu9rx8SkKteYs`He!Pj7zy+T0m0 z&X|AS+_5=tfBErenTK5CD|oeT?!&J?Qf(i|AC=7Z|4m`I)Bdg(I>2^yDEL)lcR&b+qrM+H3i2V z@!9B250tFh0JX*Z1P{@)!(-$;76-p`K5Pg)|Mr2g%U8wG1IKABJ_pBXcUv4fug1Y|jw9D6 zaqySH3CH56J5Ils$HD(2j$9iD#im~nhyS66#xDP09Qt?1p|duQyjgJ0vE&^Tr~LRh z{oWI&{K0YfIW>+RIPb*b=chRM6XW0)#u+a=Gg*VZ`YcgE@W9dY*{JArAi0ICS>L;pfCS{rg!QdwvUwiPgV(aqxG>;peD0 zdiyp``LS{Mc_)rPxjat!Yvbsxe;hiO#<9cW;_!cV9Q@cg{azAhynGf%pFhOuugl}~ z?~FL|J{m_qH^kvH4g93y-}$MH)9$%({O$X3_`e{I-bR9-RA-}r6VWfR@Mpy7uao2G z`Jp&?R~&v;q3u}qe@q-bB*gKPcORx{CuzyrlD}jM=~BeEru{>-d#ILny+>e1pCU8X z!=Im|9jmpiFBXVYiPr8{cH4NKM`2`gCGafYxv_WoJHY(O+M!y__p&_aKYm_Q?RH-7 z5pXlcAwTts&l<(2RsN4IO*;(yx6P3JLyh9|B=Dr)a*V{^qf7iTpmUgZjMk|5cdPRA zp&#<8ZI*mCD*WjvPdXhHlFmeh9}asK_^Ty;xWZqJei!m?kaP+ZKKDrW*P&YKryha1 zSP4)HI~3&=KMktwOCO@p}r4M@-F$oBQS3(d=>OZ`>fd{%Qq@~0D7i9w5ond zQFsk>l4H<62K;l}6(~DLcZkm)pa-7_u6IuCMQUIt{O+FPkIERLeY0Cxtw+~0{US;wyFL?b@6G&c%i&$ zO8*M29f1NYzr<1gZ@`n!v`Q)0{feI+^ds?3Km3l>UrRiK^wVnGnMz(w`G>Pq!}lxw zIQ#ue;Kew7PSW{SmA@B?qWvc;`?vhLhJD0KYRYD1<#F#YT`Q>?TVGjKnqONzy|#EJ z!0~0?tm;|bvIcMNtnzBDFn98V!qT$Zvgws|-m==s6SAtRXO&GZo>or1I8e{9r%gbt`#dgh+EU8Cy-luoQ zR?aG}y;k;VoF1K2=9Tp`Yo~*v66&(rN^hBx3ITwCa^h3Iv;B{?~xORB2t z%0!LHe?tCfS*6Yk)0`Qvl2%way`iD7rmVKEdRB2&rT5yx*=fJM%~91DK{@#sk)M%; zxnQ0{YG!q5nV?ZUvZmgfQ-W*3(Qgn!6z!#~5l~!RWlVitG|4-D3A9U!2Cy{mk`^M|)DyNGMApz*;zot-E52LQlnK=`dQPNOc zSW{nHRyez|)>~g(1)8N*`83JNliA_QCgUrURz1nf9`sh!R?jJ{Dw{RkTTxh6TU%W# z6@-ytWAeZ2BvPkc{Ra)B)GRo3DNb=!RdosUTUgdmQdZ-wtez!UOh2ozkfwo=Vh!I| zLu_JHYE?RSa^d8(!kkHG73NF1Pnn z)U{wN#BnW-fmAu#TUJ#yz0bJ(jl=5f+A=uR*=0Et3%!8;k|s5+xU{gOxXvpa<^Q0q zvrDRH))d#4MF;F}v1_Y)COJHI)_>)(BUJ3uM*c;!qZ7yfGK6LSqdz%@X#%6`%F1R; z*|m`Q+3D{Gl+*Vu3G)sb`&a%!@(J~PVr`R~WIz4mW3k%%_`Yo~{~c|u`r zSxr@O2@Gr|RVRnSfAB1$#7vV9y+&0RnQPMvE&ude3`(~5-+F6C8`1fOOwXF5SjTk2 zKtja~nA4nmo^Y1jU z$Q(VUv|3KlDkQLXa#Vx=$)!-Os2S1fenKKkEZY?G+uO-bqlpWJmDiT_o-yp4|3>p< zUzNZad&_brijj$Vq6A*p(ewWWe^UMbBP)nL87WM>XwoP~s0dw0mvdzzhqc`;{6_gk z*V^Zu>Z-b&N-nXcmy{IN$rV=Nlv#5yH6eVjtDjlM%n}tj7uFys^OhHuRLs!oFeR00 z<+M1hoO1^H03lJy4CWS=7gttkwPmEEVV=ZfslgW^uCV5Jtk){NS~*>wHnVJI$;=v& z)Lg5Ti}+ZZscO}g6wfLLsp3+cBblL<&tYV*p{I*WwdrMEZ{^Ig>hjWJED{RqZ0v#N z%<9=?z>-lrnq6I0 z&(P)Ux-xHlO;(=lvls@AkAqO25`+N4Xy=HwNQ z%gvjdGqG@F!^pxKSk4HnoGyXvrIoV^>+8x&dzTvBFq#a(b9c08xjI5o>Jpby-eje!d9*V>FST##y5hYpv_lVlDT77{X$1nH1GARFgHc zl-%T#ms;T@RS%qv%yICNk;PgsQFUIx zET$*tMuZyKOUrhN($8{M9aT|QRTF<4jH<1lb#e8qNoD9M?$~7I$+f+*uo%o+MjDq{ z6&+&%V-<`l7u81g${5kBoXjd5 zF?>|yDm{8RV)*C?8ckwvxgSoECgN`Z(gETs5$ARLvS9DwpxNw&gLxoiehcc9j#;A=$r1GL==_71%hcUsYRM}ZhVlyiaNm+;`qNdFrj>{*`DNY5XTQ#} z;XBoRG2(e{*!o!aJ&C`~LL%){_d|*2d1LEi-KQg-XZNg+b$^riJ1rzq>zn}b)N|yc z?oSc_bOZ@qsQWU+x7zU5{Yv5&+wj)?ed6cX@F{IlKe(+A-A6yEHat2s@)=^o%aI_n z(rkF^ECO@J*zj!I`gm-3b?ZuG={7uhvOf7Xyt*|nvI=Z?bbRDfWW(dALF7|m!_zBR zpBfup-SQJz4K_RsG4g4&;c<&D@>yuZ`jj;+sP8y>gxBA+ETJnth~pQScD zZrw#bEjGM$50aU!Hav!02 z{yZCATN}0iKicp~Hv9!Pe6kJivEfr}_)Hr<)rKEy!w<3HvuyY@8$R2HA7jJk*zg`3 zew+=j+wkLU_!&lhw`8Ir|4PRixUv0w|+3+)L z_zD}o#fH~?Unc3n0eAaP(sX~DcYyKc0h*@!)+epA6~blD=o?X_=aeRVyH51rim>3+ z?CinklNMPfdUs%9$Mk{*Ew1II~T_I|(y%Gg}0_gD^uj zbBTc8C(KaIY!>hq!VJ;OMghMaBC}b*6A7P6xKY3t5@rZw)(Chk;UR>J1bhzR zp@j1VJc2Mo9n&M=(+D%fG1CNmGGT@`W~zXXA3`0l!a}A&J>6;4Oq1ikOW8 zeuXeY5VJ4gr5hm{Yjf zCg7cfIdz*Y0^UKGQ?|K8!0!|0RBbj3cne`p#b%>`Um?sX*sKxoTEd)q%_0G>Aj~P( z%op&pggMok9sxf^m{Y8oCg4X1b80nH1$-Z2PN`tymlNg`YBmdaBH?nvjRL-qFsDqj zM!;hUa{@Ap1bhzRO2YX99zpnO!X5!XJMSc|ycw|gOqmq;Sr7c6FZgmR<|sYzMX4T4 zJcaH=4>Y+q0hiupyn9jtvkX_Arup0IFVcfAyW2$Zoq><^;DJ5Kz>d~~O%m&CZ6oda z$*gnweXJ8$SAPtl3AYjQjw^Sa*u*m0O3sve=w&w#AREKC0tntJ+AY;b4ibz6w$`7c zFIZ3hu9&)xRHnoI2G(_L;`NLh1y3*WdbMA*^*2J_KCEe7j~PF~YOMzJz&hiVcPZO; z<79aA@@CHas~6zVw9yaVZN)-_2VsJ%d2pHYzqu2=UOwF7H~p*Jz#CWTfgQT<)9(Dq zx$SLy0zz+J!v_-d_I92(H}-(k_Ye`Ix^)>g)TZcXEJdRG_8zJSKAU}PD zjpygybP^w})F5TW}n(CP-L)5?Wc2hRfE0_!5N5n|(b$x|;Y}l1jGBxkK6~#ze zN8w9WwEEZA^O~#aNfe!5=IeRI)pVt(TP_;@S8L5I)E&FxW+bgw;!AOle)jcjcQsug z3Lh;Bf85&2hK_7~9m(87eLXu`pTw2RkMW>U_tO4-J#V<0-bD@9ZJDCpo2{=wE#41T ze2JuWWmH90`CnX3H;eL@it_tfx5Dzg|7h(-GWQ8zPfzQ5T-Bc?`*%Ul?=Wk*no3c4 z{)h!VA7M1Kc0l{Cra9=?1wGq|Xc35VAX?voF4sl-dqGbM>bEWw1#(cp*R!JaNtCI7 zn+<5KX~?ma_w~Hf+6<;$O(~-AyRA0>s$aTG*`u%L!`38`zZc`r*Yi>9570{e!f5`c z)_0&;SJPIJzq$2gLjF=vLmBk%n$~O3JbG2+uWh}OkUu+;|9a~PQGc<>?`TaG_1%&D zU9HJNm5k~`pe;AldqVXd1Yig=z(w2x9EY-jl1^3Z3o)*Kn`YvL^sl8 zJ@9@Wax&1U9z4!X@?B53PHe=0(lT8yCb}0M7{T9W^xG}`yT50W_hdcrvvKDx68uOH ztX-|O!{2`gGOfJ2f_99x-@}*~^F4t)9`CEj>*;WXySvsH51kNfZzyx=8|pw!#;2D(6f`IU5A2K60*&$k`f|4^93!H&#D= z0_n^d7{x{a(O?=WRBCGD(kd%|Y+c*-HDe|uneFeo0?^7ziUiEKqs&-K@ z`}-uB`&Nx*OBw52<-WwXDK#K2_286F;~>=S&BGpaM6PdNvddo&g@Nu9q}!_(T}@SxOG=nz zmvBO;J~=P&l`--gA>jm=OG+Oh??D67FQjkF^j*+1cc?YrNk-na#5Jl=GBbD}1xse`${!`H(iJ^i7BTP#U4wyq8O19+&XaBZt$PNCf)3~vZ zmTG{Rknw$e5G36(#M%kIrC875Dk7HA`jr%A3A+H}DKGF1eQ4&?%&RgB&HKnB<6FV2 zb5tR+LU3x0EOWdBj2~#P!M=r`$H<8keo_L)rZ|N^uv30rQT}V3!s)a;<;YCV&c+)k z+03z+7n+`o!RY%kB`@%Y7*mj|=^p3WiVc6sa}ghs2~qwH?c`5m>JW`pao23)cRH9@v=|_}W-a8^UCi z7yNzK6Vkph?@ZAaPD<{&N5XPmJPHMk1?Zk;%%>}O<6ig{4Q7v(mIrWNsMMWildZ`; z2H5mAPH;)0U-u$v+XHjnn4Or06;if4H81dq9>{hl2anKw>k>mZ_UzlYZ_k#KI!bkG^t-;jqb{RUt6R5w)M{&`U$3FsOU;}lI!M=s6WimuBJS+ z0H>{o=DJf|epBf6N_PrpXX29e_O_PS?0{FPDxBE)CFw#f zqAl}RF;}3zsLJ{IAgeYbk!+NUh9b{I>)xQ5W1%S5i`nipq0NP8!dxb(VSOOulbmGp zcodRSVP0jgO6i^_C79o>=zC*4L|vT=z1(_kQYsNOQv=`Vfyod%kfbli6CQZrvJTqS0gc1K>$1|+IVQ)$o zAENrda`^|)e1qp2FBlYprG!$`!p#kt+KLzPCD}iPv5#PNyI~@t6BP9jna+6Edk9(^ zWTb;87#IY%aS{zehm{>Cb|75-!`Y5+eF35tInGyb2R(dpWFC7Al~+^|mx@MZ+|vwP zcsPoP8Aame61Nq>g9N7vR~{osbTvH&V)XvrDtdnQYuguK923@36Co-4)0_fhip|g9 z82o54`1umOAATGBh}a4HA+RHZ77=4IXrSp@E_((5<3W@(js{tqKG6eE=>TH~xFod| zVvNwRRNtTn>p$-LRPt@bU$D~!UmB3fW}i)NYzM(7`avhpo$P9=AAlDVj)!}G2W}rx zE}dwN9;$K|!8A5vU}AK34HoMWc*rFETspg^Oh^YC$)R)octt5lhWWM~;|^gnUiT0h zPJa3p)W$NS%AHTsAHq8b#-W1u$tVJ==d;aK#cU*>=3XSsT|Nyxt6Z9Ubt<1UC@bCQ zNf|Dz9riSY28|lo?gAQgHf-6Q?jng6 z+hMj{Xn>yiZ4}K#G_41FSz~NN%ncksQy631z%faK6u|-by_J014H7fmmC!7Qs4&MF ztbq|WL=Rh*#=&YcOgTFh=_{gkS(D7c8^})*VsP4Vj8bF0Q?4VFTzo;BG(i7+gt)oB z4M}+!RrE6-KySI&lj!4UhJ5OffkQJQg>W>FqkE+k4o4G?IvC64SWZR{)IbJ3SvZ@% zA_S}4Zg#{(n544C?Job7w9-Xy0YhcFU z+{nqV>(#E!MknYRSKzC;?|y_>|5Vr}Lrfp(ZNc$Hq5NNcJAQTf&!r>W+dnt351a5D zxb5Q$I&t>`X+@h%YnC9*432YW2CquR@Je(oTX%jB9oKCSqDJUE+%7@I62rI&OkmjM z2HJB2&-&4GTy5oz(_KyFiO2wQN(7njYFZEE>gz!&C;h9O^iPewPz>xb z(cd<&BrkwC8apojZS%5oOEzF{BsVnd=BT{JjO`3zzG?f;bsGNvJfJ`J{3EIE8rC*8 zko|&&MR-DX#;*AtnZET2nF$|epIhoqtbZyG&Lt@)IIDqDhQS8Jy(%^Hjju3^Lg1vV zQ1%%brq`Unthov830nhm-NP~i^M<*W?ad5~8$zAvB|XO3onm%@>|?gmQk-6TZz$_()Iq5%-63u>G>e=$DDT zm2n}r=ymRaI_DUeriUcFlAU3CFCH5jG*Bo4oIUi1+>*_C zC7WgY(mJ*L>11FO=xIFf=&j!qBKCxf$K?->&3L*In!NjPE?I zyEbMfeed#fXFq+jiZ6qs76wM$qWji>>-pcBv&BU0(N>&{w!`@WYkvI{W{7Q=2P#g& zjMaOquNxg8C)9tUPIMb4{R58icH`)e;AGfIWFPCRk>{yJ!?QqQ(NJNVOlh#mu4Pxc zhY4%dOMW(1e8^d3p02_CJg%mlu$}a_uIDn_$EA8e6seo(Y7O73xRrpZFOis){1S&6 zIQIEsJC1$6XbWkLj@0boy-&EIql9#Qw|7QD8^{;C$>ym9FwLih8=w|es}_4^+l>Pm zZ9up*MAgZ1nooq#PG`M@xB@U0UJ}vkt-v|gKn=w4gom6M{*2K_Jst*q%4N-S`z%hh zTO=o+iFML;86WiKeUcw(3qsd2! za|nMs$vYN+*Ath)p4%6ZvPd=Aqc#EV712PiCKo|DPASu3GEizu+8czv_39{K-aeLikxV9{ki^?%PJ;(UH$oP<^1aCe@ zUv51Mh={$@GBe~$T6A_kc+VdfRQoA4z6TpZ2}4I_71`Ph)Scl z8RBK`nlJhruf8QHjpBy<#@xV0;~3ST2PL8nz6BR)HOL&}2-P3~*Ntt)9xCS?>a->CyZ{m;GXfWCnf~lZh20_w;=(`y-}{4 zUSUH97JGW&y$K=r#^}ycZty&}F;aBG%kCl+TQLKA3Qxorm6aEGT40T4OnjCMEM{E7 zs2Y7i01<4S*ax}~gv5SPC8j9n6mSM4b1q?yCJq$ri?0E|h5rJioN_-$*P7p>Nb~Au z%ss|C;6!Ytu=N(@z=Z!4`um--k!F>%gsI5+^=-*NbzBxke8xz){? zBeIbn*y{4{!@X$iKX8h%g-o&cH%Eabr{Jl?IRqaBT=ew~63lgYgIkh-GceFg|A7h5jaxi925OI=GY_ z>&7t50ek+R**IvmAq{V(1b^m81Ev!g<2WpE^Ma3w*$y|#jA2`lxt!eNGhaXbR=0s- zdd4Q#%{P*DRS%~yc&F>ZH^{>Cf|S*WaqZm@v(4q_`K=Yl<11W+W~Jet54QB+B1u$L zlsUg=4tn-gkc^meC)V!L1E9B?nAgwyVp z5(mxM2izjy!pnh_+|{GEO&^+ww2if+$Jo~q>H9ovA;8&s!X$+GgY#X!aR#M=&Xxch z?nF;wiEeC#FAp+0gtKLa@w@GkFagz*Ex42tUydGOPc}Ua!cK9%LZ(xk_lQGvyEq$w z3%?4a6z2)BW!?tEZFUm8KL){tF$j7|@b(D7DI{0|#nT$3n|ax-A=~J}W`gc}*)1YO zmp>0~n_v08ow=V3$rO{i86<+*iu=oRoNT6h|UIC8{%MR|n5 zp2QH;z6JTj#h*SWz;utou(lS)uBGrxCQNr45 zHg>!YwSs#s{0>UdRBMFwXS?;hkb46%z}*+d&qQ)7W}~a=T+B>jf?0#@>~M+X-#AF| zj;Fk}aE-6zx5rk{QK?>t~jHSE~*nm?d2j>be z`Mg+x5Wa{~!boNQH^!^Y(pz~zd&RMo<`zpD<93KiZu;Qp@$N zPtMEu0(Ejc`wec4;-`+08mq^dlHaM4)m20v$cDM(jlI z2$w0t>0l8!6A$I>ungyxkD^A;9Ah(ej@Ah%!^y{`G@Ll| z8@u7m0WLfcG^H{M&^P8!#CC<-*{3pXfMW_JJ{8(!DT%vi0wf>qf8?9&;;Z%hNdw6~+j2 zVuos&%~8R#V^Xir|EsTU=R*poI?p4{q0S+|g$IL_ROdnHX0wjI6wP*`Y}kXoHXivP ziZMnjMwm@Tx|xaz^7I&tv=isJN%b6Y4n`gVF1!~=$w(vAV(#P?6nX*2E1+B*+l#5# zzZ3ld^9+|O+Z+WM$nO#=t^$e^Zu)DiyMxwDhaBS!u^5uYteGyN!L*B&khMP3H z=4wjhH6-jnG10Zy7ABlbJQ^8dihG}-TiqQ3^L}_c$1f$r;fUimNwzpXW%P3eSi-@* z&xCQZO85tqCna%bGxfsYpibVdPBH!oc7!M2;A)cR6ehC_H;@Bs#58Z501eUT-+6^f zl5RjhcdmN~@Ba!MMmdcxFQe0S(?Jj+FNh@?8HKLc`ttYCx?FmY<($CB;Q9}n$>VvS zi@D>y6&j|I@ac%YjE|H!ZJ-ohY-xER2E3(Z<{0e=n#m3UvXhh<*Z_`_!BGxxDW##Q zyx=hRMZUQh94>#D4oPV+v$YL$5i*DO+VuGLMwr5I_*-W{S4An6ELckUHUoNINcP<` z!ix03CVfwb9?I`AJSZzI>+^P-zO2r~?Bd_{6WW>oWAywBX{S=!B^(0}ifsakG13`Z zK}u%RBGQNmy=j5rzx&%Z5G8tVi*bhv?aa*4VJ-GK#t(-Ji3 z)LRgQoqC%?oI`IVz=fwqit$VVPMQPJ|H9$o&E z;G#@^3Gq;sJrDV1*Si#L%l;Xc`M-wqGx>#LT_3YvDw|ZBUb1zU|0ojX7sp`{S9;j` z!Liww^Jo8mAq3N%IwG+PB$vH-tGL5mhh zwhG+G0{jw92rFCUZegl)b{7lqONgQ}U~e6kG74CLUsNg|gl0Mo>FO0YQ)6scBQZR4 zKoX6rjb>XLCCD@R#R48xAm3JCEDP|91yWUkG+Tj_Sb$$FutXR>#9*q#;1kv2v1-AC z%77TnyySeg9bud)z8+=K2l@3EMX*SOorm=N3++zuUc;T*euXa1CXN`i1ccVs#NgOT}7h#`!#P zdeaSM3H=7a`U>9Q2@gY5AcLBhgeLE1({VROMag)(q9H%)%FR(c=Cjc#14T*uV04@w ze1WRT3;Bip@a$q>d%FO@h-O^IGNF(#pImUoF2k|GHnF~qbY_B(Fm7j#g^qIk=*;DZ zSQeCtqtY~1!I2*(wIiSjb`cn^~`j9uVD7RB8a|0?KIK7%$JL`Ba|K$G)i=aLcG z+y4d+;tWgp3b#v;cYeZDXeI4e##l!jY2(Vq?t+_CIkC& zWROKa{2|C0UqP*8{_zz^mF4b*gD{#fSQNrYVGtU}W(Tpu9Q%b!iE|u%SOlE&=tCoM zH-JAWaIxa$Xz(J=edxidn{3BXbMev`Jyf>!3hU7(oW1X}-#m!2q-=gUynw!|u^CvfVxWCo(H|RiHqD0U zSk@V-Vy!RQhHsp<-_{>Pb+-O5&S*j5-EspTcRUJV*<25wLUHwnlcQ(@G%74$!3+>(l^VKX|81*`kEd6jDu}@ z>j;UKMNl}YPs?|#-=6~Cd`Y(aM~Iyh%o)aKY;%JzkQUBT<_tmelxRY50v=YlTDw|4 zG;_5!p`b7Lvg`x)%Zjd8`7x@5)ILf;SK;*zN(l#` zhNb=?w+SUILB4TupO$Z3w?8Gk_M%k62HCP3>hIe!-F~FyTT#Pmxj?qOlr6u7(bHSz z!}r_rWYzK=zq#eZZ7ny-mRs2J<$YRSjAJPAN6z~%NSUXi<;>ueqRe1UX++x-3Y6J} zVHQ!T??D|)>K55{E8G4E#@$=$txWhd>K3s$?*0+0<;#y?rKJm zt{-tSP=-e|AcI^iG(KDU@{*x3=@WJ zt^YzV5T*0r(#evng#99kkS*^;R^wRqjsw}fl57ph-r0*RQ6QUYy!ua}5()c75+UmX zS#XuSTj`?R|JL_%#xt@7wN+`zGLonV->SvX38ti_uxw1_&48$O7C4YYEBFGh6**&3 zgP19r4nq$*_sDfN8@ziAi_2OvfJ!4CLCn^c$PHIPpUPs{2;Y~fZWitNlY8Zzi#*H- zO`VwL!Avvs5PqEIQj3wL5XJ$FDS{#s#91av{3j^P842A3qRvR@M&j5Lw$1Wcz=aP4 zQVMhlhXVvU5CX}FP)>PjP0)cIZMnoS-vA5p+zv*z{W;9z*h&>~(dRIeVzm4{xDJnq zwEVVGWoQ}pPr3y0CJ^bno&7SUEg=2q8Rw;>wt#wm*K zhyMV5CGshV-(@gHo{VRsdmzR*1`lyoH0Pkpom{;`oTHc50_W`IXNX(OULFd~S&Zli zPmGba=Gz$y$7y~gG~;akGB)opiY($B%083Lr$ZmoD3W0#RQ9H)q3qAGwBWk&NZRTX z)*{^a1=DPUbi1y6Ci>h(#pX@qVTS}aRwMD>%$3*O)UbS=v zSqpyG*IrYpv!3YE#kh-cXB?f~j9z!@?DsJkbH!lnXVeXkjxhECR7e#p`a2l=?prd} z2#r~@Xt;)aO}Fv&^Gosh`UcWD`3jO(hbhe=4-SV>Lfl``viVWB7=*F)e_jl7Lt>B{ zOx!ahcRR>Anc7W89lf}NI7ctO23(lO3#EZAgFdOO=}$pf{$H$}ivoZv_=eB$VrWYi zU4)wQ>Kx?e1|Q(Q_=*%3>W&y=H40e=D2=r!N8$sK%9@&}DV9r_0kL>8Uxn#7g&9i5 z9U4B8j5{=(KwK+TGWVOlR6cL+5A*wgew-;}Ld=DLijrC+O621*R z&wDDjy<|zs)5skRs<0YD3EuX_E|$FQn_`K_ooWUWQ^)~XI^Xa=>N(hU4S!2=Bv+Gei7cwek9385j8M3j9CMMJ<|3tc-6?F4Y@ zG_--lISg$#XoOu6Lwgt6Hhts(%L_=)Kn5}zv`5D}seTuts-#w%3m*;yzA?lpRe0)X%OqP%_0(^s!W)QQilLw#; z?k7a7letC;j6~kGBM60wB%-iC!**z$O}zc<)YzBEbk2lZiF0UdJ#gWfk(uyuFlRPm zGUAn2F2Y6VN_8Kt;d!E8AUYdka5-W(sa|5B zUXeuD{Ru>bdcTB8JJlP4EKc=aOPoW!Wx$0OMbtYE*2Gv~-@|HUE20glfS63wQ<9h? z)9*zzYx2v<+qbCWK86TVehHJ+T`XnYNR#dh?1dECBt zw?tjmgg|Fr;M-_N%zi|wIOOlP`U2wftpAs=?YzJ~_Cq>D)!-_Yplb^3jjq92U-SW; zI2~IKcsAEUXsI*SiB`a_I(5P?Iwa8|Jyau|twb1%AJG^y>4BI;dOQiOIhFYsagKeLC!)-$=y&t*2%mkMbNTsJe(`Gl)$=%CpN=SQbu-yL30IW$XrwfjHy*Nb z4pcc8pfNM)Cm{)aO7t)f-l_v=#$ubemb zE#dh1z*SqJ%T@?uKz2ww26|5J-be%wu_V6u%DR*C2i!$gh6< z5(l34lI>mm`kG%Rzdps6Y_(aMk4UUmN3vD%^)d^;#II-h^$&h6<(K17j}RbVk`jqU z%J>n5`6XLRlSM>g71_Wd{E|hIWf75BMIK=heo+ngsdRRvN{y@{601rztMH3e&P3{n zY;;C2Dfc5EwGuN1fm>dUVculJm4Nej#POM`DPd|V<;-)$u z()^{XubG_lAaDh2B)S84Dsc`oIUKm~5kSgevWvF&5a>z$$hQ3=v4p)5rfc%c{yvt; zQ}od7i_sijr1!IUkSKjO?p+n>!PT9>I^WP@%zyxT2nS~6`|3R5#ZbkHm+{3j;%UZT zIkknn$p1#gSJATaXw%=IPM)$}0~6PSE7^)33brDV&oNO>BhhLm+y>Ht4aPcpK}GjK z5w|(%UaIJt2!x|_mqzFUtr<^?zC_iqA?ePr(Orf(k>?jJ*(AR|{#Ek(Pw*SeN{x`l zE88T)!1FkNaxlw{WN7EJpnQD zx5CbGe{d6WSYV30U25#0Mc}4rxEZ{OlU?X2<7qMCU~W89$MuPjTV!XuJ+cGGfJbrG z>vpM6a$@{`v8BKa?5pDL{WOab!qF!zLhx94bXC#m{)?(>pH)y&*B^4JPhN(mQv*^u9!2I_2F8oKxQQ#9c^wUBB*+zx!j* zTNs0$m$*XG`zz=fQvk}5&}&3Jcn>XZyD=NR*R{RtLp50Nm@sECyh`lDuFkIZt_@x5 zy4G~9=DATG^HmcTU!^QOBHsb~vzWCSFl%w=q1UWcfhUGuU_+R-TXH*a68?s6PF~O1Ca(#BA1B~9YdD1^#hSu0tvT!j z3;%XifOdDTdB@)D&d4*!Qp0fDma}-dD^#<8^~0qw`s} z*s5LMV0#rujB}_@Q!K}Px)?a82u;NCB`^PBKD#moJv|1!QN#sFZ~ybzpcwS_!W^9Z zeM;Poq__Y1>}8a5%DXfMy+y>`NqYOA&#Ge3D~Lfao49|F-u~ya<6_WDib3yt^m+I` z(%b)h_Aiuk%G=7ij+<+b5!XU`oX_BqxD14w7XAo7h%p)Ya;(r>V}dz?W0;N(1L`s^ z0s2fZ58o~)f`g<76$C@UO`vO=V17X%Nq4cNyOIN0(d}Wt)`xDcqPs_W2t~IgLYEVa z@mPfJCP}x%M)%2>bich#^81;cZm(&@xHv*pt#uBuQN1)K)xRpLZ%3)}HP&KL!Y|_| z6p$L=90JWQ$HLdrz^{mi`)XjkqS+p$$ya6ACJ*DW2vyD)pju+1`eaP1--e|8FN3OB z2a1J@TIvL!6MB|2g>62FxunTY-<&yxXF!cZV7igXfZ+$->sGo~)y2ofWTjlOGDiy7 zK|M_jvKH!x)q>`Z2$}81+X%3vp2UJi`Ro%(2B&3bgA5)2GT1B~|LZ?lq=e(w1tD_w zkqYDNj{p5zF@hY9e>Mm^N7XdooTKVu;`VS<<)BH%cTP&fVo*wnK`D_qzAeT4`NzKe zd`=!6UVIC1PJY6~ts|8eVo+HUgUYQjs5B7A*Jp{HC8tD}$DlGU29*)Sy+JBNWAHLK z1{Ey^l`mlx;jN_d$zH_a{|4w}XftsR<-EXZ4j1?^aob61F-SS7_+n70jzQ%L;=U!7 zi7}|0AA`!!7*viVj;}j02Z4&PY_1rh;Yju^Kx1Wb9PENCi>_^5?{&Q^3yPpht}ov0 zdY>zcu2o$tyOwvo)b&Ewb6x-JTGsUsYkeCW@5YF@?PH8EEVCWgJisSUiB`6y6$3Gn7Toc_)W<`j$z4Ldg{@DpBN z8_XG!9r!5IwQOy6LTBi_{zEs)^J#%Kd}(lIa8MvL_yD}#He>HD-F4cUU74;wCpPbL z6>O^iFw?gdFPz<$lks-_a$sKB|lK=A1k1`6>jDzjIK|u$0~)y+|bDjKFP-sGOsFj8;GUYF%C+4Nrt6 z;r2v|-kz0=`r5{<6il*Nc%ens7~s+*ZVaBq$jaw+=AL(vJp|bW8?y?Cu8~C=G{`s* zc{=hM^mgA8;_=z&Tgn8OHm3OG`hLZWQL-xZkjUET;|$CKolNoB=;I;)eE1eK35f#8 z7sqZzKDWGBBaJC919+LDD9V}}eLN?|{Ct(a(N`ca0pC_xj08|=uY_0}&=4jeQ529a zp$3^(A)&=g+I=-h_}S=VHZjc#gNh4T223N_=wk)ux2b$cD==UfDT*0PAbTXA2x2sJ zIzOPeRTkyXW^MHGoB63Kf1{7DJrvT&Vg!+&AIUEe`QTb%Krw}7ASx2bpz`@RH9z25 zV4wq}D8>Sy*b>Qa75U&=VL-8hWkRAR`cLFT6%xtfD6&*S%%o5f60b5sPe^D8Lo>*Z zL~_DxA+5rYca{nHbR>{n<$F~AMu8`No-!lvOdxwCp9o@_6$V8sSO&Z!f$S>3Mdfew zwF(TRL5gAoQNJUSzg6UeYlcrG4pC)M=Ag%jM~zgK53U6!;M*jNu>kTtk$hd`xBH3| z1{C+oqO7^mS0VDD3Zzs86G$u}W&$dZ&|(QKlu#=Z$c}^`NXsi=S`>!7vrNdh6$xZl z`5h{Mqrj8?CRuDRlXhQsB%cW88%SMEjRZozA@E8lm5PKPXn6Uc^&mM!KX@gCZ4~t$09H zTYf2MU~%$0UKjcd2(F|3R5gPl@__kslS@Sb)r2%-GnhyDLqW7=k3iQg;nm zN&$)uLV2mUG%k`P^JN8HRS+7xhAf>f^J`>&h05p5v0Nuh2h03snZFQ;jis-BlBF^V zElW^65tgcjmX-pgSZcwg(TZBAK%WUI=^a@?u+%`77R!8kP-;Qtiw>I4{07x(If&3r zqL7Xx%2&H@3}Dry2bV^UBuN*F(iKS9B)e{Th$K&u`E;SkU#Rj=VtymFFhJ(hg(80` zlFU{-nV~Je6SZM6pWP@ehFwn4WI6gsJ6D?kD|6|xiTYzyeuOXM62%n_4eij0q8#0&ovQ-CimM`A8b?d2bg(E* zugO-a&E=1fs~;B#uIXTrPglwOKQh0AT)ijr>0pshAIbbESId#l(M9v7^k^I9=pb1R zy{#>OP?)j-lwzzKmqw!`oGL4%Ac+v}xdsgv$$XE@AEWXokYMs6;OZPbZ9gdfsUN7`GrZ;?g+&hqU{W&>PCUlaacb4#9U0ULl$^0l!j_ zx>`7+hv)H>Ndx>_0e=3e2*F}SJ`SMpg_ugsX|vTc_6Q(7Nt9@#Cm9~LC>8J!38w+B z@Br2&oDaC62ylgjYc#`{j}8s~X-IBpUMgvmv*rJyMqiUTI&wVB;cwe4-@=^dk)wxl zQgPeA4R)W|ep4gq;gc8Gog27mA+C&v&f?><^NMmadhm*mhi_mr^E|m3zXG~TLaDhK zzX0+J=#P5F&j1<#aGIz=Dphe0F|kyk6fs-$1NJ`k!2h2rK7`JxhVYjZ@&a@nMGx;H zl+Iob<+%4N^3$X3B0n?NF7gryDUt6KkP_JsfFj?3REqp4F?~f2ANqgN=j}Ok6_{_? z!IU={WQ|KC$5=(aI&{3AfO%6ad0OuR-t-`<1n0q+Ql6*V$!VFCiuCivm)z{tW2 z1Ha-oi`L=p*I?{+;IMlSUPaT-_>zj@t7!Pkm;rwym~u6}hTgz0DfzeIt^CR4liMsi z-$K@H%K`Ci@weTN3%vyS+|x*jgD$YDwI8nFm7leV3fTci3~G-v4j+ zUDl9^q_Kg-CvgNbAt&RD`3o}ni%tpry`K!Pd$7wtz6X0|a1MCQH9AP1*}9|w+c8t% zKwoG?hHKdrOdmJs_+2FNW`Z~B#cr%Wh&*{m;8Iwuta&7A;^%x`(@w$=?fr8;w?k(6 zb3O&QivFC>G%u>)ZrFvh!3g8waQ@vExs7vTFp_UEq6>s(ErkQK%O;maee+=9Bc zn{KD0u=&ZamOZ7nnVtJ+t#jR};S@Az{B)B7jb;(QAcOmMY4Wc7t>n(Jn>++9I5#K< z69@V1Pupt1g?GUMZR~tNMM3ICklxG#TkI(`D1OdW{gP)2+T>ngH_94+Cj&_&9M$%d zz)8JNAejLqgYrvQFUHTy-AaK2>Q{&612En}Njbj6Sp^I%J=XcUF*vEvRg22dRo} z?lF)|YfLiz6q)NVF@j7tRAO8lA!A4~XV}PG6O+tJXpG!sNp2Xr8@~f-A!9h@9>LA$ zT$CwpPK!y#tH?Ypxmipyo;r(+N0NESM&=U{gGTizb#sOyGeB~)gk*k<2vs4;6xhhz zACt`IXqiG?CCMx$nP($p7E3aR*~mlId9f{kB$qU4ZD(l z>uvNc-C(HNdh-BYR6Qp0>m)A!Sel@3y^bG%j``k!$3V~7ZZq3(+$5Sy+$a$1x+x}| zD`U{nW6&8z-1(qmp4+p(^Js%&(Ax_$aq{&kap#a;7aVeY{l1KHPC1vxptp!PZd90m z-7meW81xEa(90%n9O+%OUwX&IpqCVb-uLL}a2DzP%(-~DbwA5#Pya$Wr@XDK>*&|V zh|49tzk{BW%1tq-)W)E46>)i_66*m8PYf!=GC9~ztO1LTUTj5V z1Xoi!uGrx59dO2Qd_i!S%|i?FNH!;|Y;i!f=8VthT`_MB|(jv-A$S=`u= zyJJXGxk-tKiLraF|`-()TeCN7q?RTz%HQ%`o^P~muqUEQ5<=-;@1oIU%7WH@kkF_^} zkD}PxhkKGVge}Qc)Ch>91`R4xQ6VC%EGZiVN;NMgSoalnytFIJaYs@nCM5Q03P&mQU;SenXt;;VyN8B5SaZ`r%(BJ4 z88CZ*xsov65X+j%XXT>ed!xq0tT4Y7!F%Vv`sY+v!)Rz`dB(jnI^tx%v0>;l((&$J zY2hVJpzCf;kAo8H{(u9u@5X_3?bMD(U(ekHtp;f0Lk&%hO0Nae3eo$;Mk`rlemlM|BWNEJU@G9kgttdQ$AUFd3uBuZ(Io@3VN_O_<6MUUy`N zn~b`#I$ZPCzvpm`y2EX}?gZ~6kP5-NiB7Iq{Ou-63x9LJ<@dK~^tVH}Ph|zQap)K5 zZ>_M6y>lH3?d|?{3}AMDa}c%*RqOSjTQuJD0zQX8lr2ktjjR71t}pz7jFx@!{jHDV zVrCm3*xUhpyqr0q8~!+;072pp7x>~f4lOteo+S-BS+wA@w}`s1G-A3GF@iSYOCXdU zjL@p#c4*d#-V6PD*p1SkhO#w61u(l&zMBcy8#Kxs^v^TfO!u+1r$wl&5bfz@3^!O^ zpXr^q)9;v{l!j|hkAh})@t155h+vbUT;r%XivRLprU9`*g_-Q(`{%) zY)_wRA7OkP;A8#7UB?gzuHO_j)uAHo=6x3VcJtl~nBBasglRPIaiD_z(}oT6;ppun zaF8*n8NDZMxDvLs#iTnSK1`$jz7bSz2NlDmo2hrlzbQn+pSmy)$CTB&f2mKt}_X{ zlU95Sg4n4H4MOFZAXFTLJt5jZsEEE-^a_-|S??20=&r>8Xc4MaoEo}SJdN&*lxZJp zjm9=r9>PF)>*bgzD?ti4LJBQH7vWA@0raxB138ok?6r-mv2}DKbP-%Oim^?XU_M(* zp$Pd%nf~wM!`l16@KiEH>*#{Xj1jjsQWyf=zh!o?{$06g(21hPEM@)1&MjIKD0Dv5 z`9SA=ovS+U?z}6z_<(!U*OA`Se+7K0an*{7uPL@th@TAUIjdMv$>b&}10yG0xqW`*<{TIHFbP5g0 zNrQ3np+GIR`4_mGoHpQmKi*%hSEC$*?wetZL8B7>_7F6Jr2p?j&|P$04nfZhIr_vS zW9aQBpeALQkz>!RUX(+Ds04lJnHbn>MF$K$X`p#4bpsvs5|}F9_Yew<@j#XwuDCH! zb_a~#fXF5q=O4Z4Ue16Sr%ghC%H~ypfeOtY=O2@n7ZMa|5XaSm(47v)V$t5Sg@6sf z#KJypkWnOU;KrDq5REb<1Tk znEt#D7M#qoV*be0@;M{*G5^scbr*&Nk|qbdr1_{$bG%8@95-G7n!)4Bg0i47B~JSY zV+c|`fWeg2?)V-BjHZ2mDPl@AQ;mYQu7Vj@UPnG%;z4-63@q8j-?=w^73sYZf}?-* zudlgU{%H9HUnxsXtmJ+aN8LmyI zo)E7MKo{2I@>XxhEvdm&q|*|nRc|8@Ur1+JCd!ONwUL-76HN&XOyEZ%i)9k?)i4-8 zsKlWVClscU$h5eZixglj2QW@A!Fvz>_&xGfJaf-CRAoV2%&XQOVI5#_lgBuY18mju zgDWym^Nf#g-1}_fj)<1=@!A$qf6EdYP%Nvp7Wc%d(MV#;F?WodOq|ifti#G;oM+rmYeFCiOxXxwzd|B*sJr@&+LU^(0ZNw2==RpWp zYjvL1F%2BA07uN_xIAOsW4bUV5cKFQ+Ulvy*{AKjgy)XVrP^)eg}iBK=-SXlXj8Pm zv^(l&BULq;1nvL6pzx`7x1g}4^Fw+B3?ZIA<)mNTV(mUy!h0lNRNjicMQ{L#J&_vnm;E<{#$>ov~yhXBydmsiQ{Pw?tYc{QW`d0q`>3cQ zjg=ywK@It+XZ=fAt?AWF?+0A_9F>kG^A(>t(TX`;w~CW%Vz$nWc|Bvc?j3CGle2Yl z+U5-OMVz)d8dd74g?dM`HbFMN*QV8NF0&b*7 zSkj{{>10dVrKh#yM03->$cN8mk^u?wtal01nAG`{>E&3}Gl<$MWrZQCnI(FKTWDgm z1uEfcU8bLC_I8`{PFxrLDO$g>O;e6=Yp&r zQl5#D66%tIwo%JMjRgf1>zZWxQl8euE4UtnL*FYqjR|0aG2BE%SbTgAQ+gQlWhF#w z5x`|*0t5y?hmIN_dC+tcV$zpr=iS7?59pr>bilZalRxGgyX%=gWPyu4oWP+Ik$e`M zXdvAk>~`Xs_=FoyBzWphv>b55iB=)aAMsaP(rYZ~wU+coOL~W%*4`Fv3c>=wZ?2L7 z3G%Ev(QYPnKAiXFk506gsHl6&3LB6Dt}(R8VFkj8m_a9!XK72=oIRA?kM`((BvBV& zEI~hdRuIB>4s}jsaU|>b{b-hO0>4eElJKJ+CRzMQa@2)|;${s~ld~@rel&xARINS7 zEBFzFljAe#O1^x6TYg|4@C>@5b*I7(?Hi4VuC)p8Q{5jq^@zGwMpS%G>SC-Pq%>T| z^@??>2&i-Q1&ckoq>1$puoUrB4lT7W4p5+H2*Fhl(1kpLb^MJYpS2Ehc`U7~BpSsB z@Fah5)oBYMfz^O(QODHe#2k`faK$H@A(eJF0q`U#A_4r--tX z5M}SmmQfZESO@Nl+P6kozzvt$O&t7S%+CZm#L;vZ71(b{AF`wsjNAAVQ5NLsP|xbh z@5D9BlC9xTly0XHWdS!FY8BG_5r4HMy~dIjPqFAY8!hl1dRnUxEegT{K(j6983cJY zqAZg-4Gn4jM3f~e{D88;2Bd&%RD(tbVU!gJh_d1=Z5blU)(icOA}=QF7qsA0S$Ztk zD`@C)9>G9J%@EP_DJh@M99Y8O>`JN3s}qG!&7g58l{O(u`jiwA z*chBjpK2pF8E(_X2xw&k)oAQ=;wHt_nyRmcl@o1DXZ_di*5C*(AY>82>-*Wb|KYmM ze!zpS>m2n9Cq49a9TgY;`$2q*QClm@_3XABbOuJj$d}7OgK?o%0#ozifR`W!$Sz-k zh}%VUk=klu0d|%U@*}$z0jSp)sMO~p_8Svnd2JzBbgt;Uqx1I8TRU&*yvgNRz&lFr z`!cp%q_SkoZf^~4p}ezkAkGNlBukw7k2*hP3XPN_ELN{E($bEdAPVbRc+X*R;E8c4 z?f_4lZ{7jQ;acYl%oYPTWS6-EO8?6HTkbld1bLKThLj+-8Dk^7S(VhKAnc+kh|yJR zt%0N-ClzoKV;OL|819wzIg+t|)vk5xSDq}LIFhjnP6 z%6b@AaSeUPwdo5r%h3sUjBVnZUd`4i-S!*l;oDhl9(Z`?`*@~)i}!o1J$f}9MCCD< zz!aa&v6zA}#CyK(XCj~guuXhAB{ z7O(2?AzryByK}8SzF|h7MX;Qw;cW#9W(4q7hsHM?S}j6=7kg-Y+VLP~uJ7aB5AK#f z3tS$ow{-N}aR>CIVhgF5C-0D}^-OHu|0>GmzzmL7SB|Qy1`R8}tI-vtVp`&38T40B z3D4EZ0OmIcE$?Y`E-id+A zSfCO%!rizvnJ-Q5$2lKw0>GSE;)SKg+br|b+$zS`%3lK)@?=K4TOenZ(~U{%jt=PL zYn|qBt%U^meU3*l_PEAuiCy?78jJQ#Dk}&>BbR)eL-QV2Ye84tW$-85{Wx7%kHTvI zA(q?@yLBYTGwEunMhve`-k$6vY_!NuP9_|2yLtv-(Hajgs!xdY#n;@=9klQeH_`+9|k7kuMW5c(^^{OQ!~#l*Y_xhu9pp zTnLa3?;Erc0SeyN=^FEK9lnDuTm1;g#4i_&W7KiI0Ag06pPx}SZfMop> z!q1LF+vNukwnlfGKXkQ@+k*#tqFjv~(c*Q|BR6)wzwUSl*j|W-U}amM3i4QL!uK)j zbx1lO>jn^OKCsxh&jLg=?9SUateuylXh)&rUrWF1TP==GDXql4>1Per| zd6t<`tVeKKRK5Fkr-&||ws>x=>u89>36WQ6_&m4q6=#b8i+4|HPg7%$T7eSncNmg$ z;xM3WoV_fGX>(=S5iI*|$g(Hf%d(%*#-gmW*D9l=b{b$Q>`oABmfnTt$myODBfrPn zXx!~I34YhPTCh~t27h-)UV?E@wyewKo!S5n*=SZqU;m;LteNXKEqlM`99U6 zghFcTt4`PlZ)uZW!%df}=mhfk+-#hsdI9iEN^of_#?ulA4MH%{qNtCnMRP||_r`f; zu=~fTOPA-EBtgDT$u) zWiEV7B+H@I9feR=-%lFT14X#+28%IRk4T7*UBvxpxWt0YH%Mi3K_)dWws`{VPYl<1 zF2t~|U9nBs0P59TBMsGwv4YUlZx-!gDkyYr+=82sPcgy$CD`&cLayt#F(jss!Sipi zH(tvT5nO<8T8z?H8?FKigbLex7A^8#WXj7&*Gzfu!-YlOfckw7d9r?Q)9d$aA-4L` z%VF!&($tOmd{sC?Kz&B*Ag62WxCDHN21;oMF2Hc03@Kw@okzV8yu>F%P2UL!4Wf1} zVFl}1g&^!bu`6IafcZBrj*WcV)xufY1)}=Jl9y;q@H&9$?^uof*46TcHUN^a(IE#7 z&N~$3Q2(*dIk4=7yNlE~@uGi!`vQFu|5Q7v8);HyHAjQAwiy<*(7j&LC6nmb=MdqJ0ueWUJoQm^?bqp67|D+Rd>nv)M*XQE8^9PPM!Ux5M+6;LRi%H0MdDjfBc5^GmK6Ss*$I~~8F*{mhl1S^gDD6#Z=NUEvj4467 z#-&w^nHEH0FRywM^ribaCs!h} zSTn*Wb&F^hlN<@G!?Gb-sD=V94*& zRmb8w^2Nl@64LXiw{E3zk9m$K-B7UgJpxKI_m7g*lf{}r?Rie9WnRl5)T-ZOwEvQu zXPKB-K**FYr%h|o>Rni$rC)X*Us;SU?o^_*mm!jOUl+pwo@?gaOO*C3Vf-kNwg=wp zZrR$|k1h0Cq`V7ZJ5$9$-Bo;q)$39Jw}mS9dxs}eMcswydY*o3nAKvc*so1o({U{- zzB*b~#4R)&Dr6w88I2lwO}Eloz%Q$1F9=>BSYAllmq=Ru@st5C(OXq`BeK%D<3K606MgQ<|W}6HIhxu6Uuk9bOGr?_-_C7pH_}MkE z)yRXwc|nBZ(yN&T`F}oFoIzDfJWs0VWT4=)?Az*Gu9gq7T0S?vourjA3#DF1g>&SK z4p82nDi0OiT7}og%o>3m z4lv+oSIcl$N2Z9gmf(-kGHw{q<-~2zxE8UY-`);LAGkRIb1hlML#%A51M;W_fM^=c=m29_WVZD99OWtLLO!ApIG zDN7r#QI=9Of|q)-d#OJKdzix<(bdPo_M!ck+8t|+0a?dapd4JvHjeDe_0eYuxauJtw_i}=b~TPRSwuNeG+9?4uW_yHGf>yoI8Yrdh& zIFy8$`)jpq{xH%&4k&Pb2S$^D%X62gLpbCh{TAb1-0q>nWW8JbQpCG;!$?2gr2i3Z zPWngd^w+Zk1pQire&^k+FIHkN4r%seai@j%F_0lV+jBt{+An$?j7dknj1UOzF|6Yq zwVrS`;kGdgFRx0}meQrM1A7uV+&dB%D3;uQ4ytZlHV$yVzFBy~RkBusjtahDegYuD zxGiD3kIC%UyV@17u$S3lUp2P;J+!`oBu;8=3LaWTY6G<>$t6cPjY)M z(5{A_x-&HmM|EwW?d5o(2^C_u*W+L+w#s3yLHKdQOD$v(+svn1BojXa*8KpJLH+$3 zM!*MOI05+Zu%92QkA{TYpzOe+4=1hKb6^Dai<3vN7Cvpy$Mt$h<8K^?`?3XYd9Iwar(y zKR~yDqa@2cdwDWm|9~oN{=V*F`y(0)p|kzdZZ(9#M%Fi!Tn1SCW&{=4Rxf}vQj%X$ z5#m{9?w3wW7(~Bj>5ba5*YgcOp zVK%iE0oFdpRI3a=L$#b#%LNIv_ko~Yh(qHaj=Bl*Guc<5swEp2GFof(H>|8cAr}53 z3BH*C9PN0FkGKWpAKw82XIHeTD1cEwr&E4ecZx}=3s#aeA3*WW4Wb6w6`=mu86~mG z;ch)3VdqN}CRWZm_5ulSCH!vPu{6O?rU@Q5DdvNs7rXJhk1yE!81Ho{Ynj{Eg@qJ! zIekTR;%)u3kuNuPInd7g#x`9I+ch7&_H=j4dnqjk+EN<-bhtgUF=Z5i$%%$*lwmTJ zHl~EwLx9~*D~KsFp&C1tIIX{~$)DfC^bcU~B}ZZRT!3mX=-d4EwHbIJQu@g6JNu>J zfke#wN2iQ=@0yXg6^WnVIosLruVM9Hs+t68}wQ-RmPJY?35#}Z+&egaLLkT{$Qw>lPUROT|S#;@SqxDsC7{Bms5TH*3&2f})$!7Be{7%W>iSbu$f$6|&McDwHk zX*WJf9^P)zCYOtJyB+JFkN6w*8=-Uj(%@*Z9|+qFIDVL6K=XNs;Q*yK;X&FRpwP~D zKhgop+4Tso$rDC6*m<0P2rcWE{_}naJKTA?w8PY%mRc1B4FBOfBk4bi$q4>)5FLu} zpAPuXcF^Bt`p<|E{xd@7^t-{yEV=;OxC6W=9j+$5$5eh4KtlPkRGx3-r1BNG`IO2} z)RjL<#!FEz;q<<5sMLEf+O4R6)<%Hde2FLZjy8Fq{)OJw`Ztw0vi>^9eFjIf=z?t^ zY%{@lc#F*h!vRWfIvCr$5V2akN>0>2;_LrQ{U=I0{D8I?+zxz2ov444(f_gjbxvc;A%tM)##qC;(!}9BIC5|eddro`kBkVG6${b1u$ZjeL`&WLxS|y zGRtsP#xcp`H#$+wV!O9El zvAyL$5b|;j3AzY<-i}?$&whm0=ac|_u&n6R5~WV>dQnfH+5iU7BoeR2K&4Rdx+`3s zM=3O)dUtiCqrv0Goui+6Om+XN%>n_}Iuz#bI@cMKu@h-$>>}}YsFuxmVJn-(=R;7h z6yBl+KQ1flN!X{Euis|xy^cfN_7~9J_4_^_a`BuA?Lj_n0h23~ms1kP(v%_d9TjNI715X-%2H+}a|T-F(OX#EV6-)EjZ^CSH>{6mBDQG{kFvhmx|IDg^4x}P5f$9`f6 zdYeAlTKbtm{|uBg15Ao|?h}KQ*ydZ&K*<*`fCoi~_ZD26Q8-1;2lAn(YSAEw$J;4~ zw-VRu#!!en1QqE;BOrF5L9b!$dQiIGo@UIzbv7MYz#GC@QKt8748R`s0ht*W6;=Ngx1dYRni^d%ix+$b&7sZ|OoInsmWiw?EU{rCi-waUgDrOEgMu)R4#NBZ ztlNGmn0H<;`dwJ5uW6+S-AYfM7-%JpR=OBB*TT5&kXAw!d< zMtkNqaC$sH;kxR5NU7UEd|g$4YLWdzJj!@^F#0=7xw0_jQdo`$y_{Y8k5DUi>30BT zm;M>TDk=THK+JzXXqA8Od!gkYjiB!3A&C5|!<4(9^7CMpo9Lu#aR}EQrMuRzg958} z23_ku#9rZA^TYY=neYsO#R*;bZKCP7J#^xC+K4YF@pOYYUJen$Z)2MnQngo4kPMn8 z;L2|O%fQ`k{Rx2Ct$!|Ix6%3=_3?zAQvV>7lpvJ8gE;N2q%;SVu>NN8Bj>xJ>rZq^ zJcRJ0WuR-&>-fNWm_|8|GUPlyWP6nRrnTiz6B=QYh;9CPco10! z29dQVVgIG9H;FqdG0*3Zv!mY%T@Syq9(aIaqfNzn>xwOn4^(V96&u5e1Xcj=e?vLg z^>nZeunQHNYgcRvG`3g5c);wHkVM#XRBSBZ-&-s7j@zRIwhMrovrKi z^|66E#a|1ZZo;khU(@NhP&)k=8FroS1kA3}Lc+FFr#PEVFTN34r;n%;5723bsZ(EF zr=x5-O`}dPVansL>6D6U4q8tiBEzoJcEId9{g1FWsna;zi&^VQd;PEJRATDXPuJ;R z$B>d`%uk)XxCIkBop7Y~R0Wd+)u~4ioqhtcU8he7+eMwq4g~7d|1auPMxDf9gqwqT ztMTvtfjVuZPRAHJT@^A&>%ngYOhTQy1;0^{)m~47gXq+ou)Wl2+0TJGCBKI25aUbJ z_q27QP9|2Hm36GH*iA|-4Kh9Vus&q6R(3I&cf}S zFhrk_HC&gcy#yPwb~Effz$945Zh&OK>;^cKupeoF=YAC8%`?S|)y4DvBT&2<6mJ=3 zxP*9v!_<@Lqz33BcI%{`LWW($m4MkrY$1$mGTx(wi2jau!^@B?gadv?lk?ze^)$dk z%y#t;U9F>RYAvT)FX2{KsC7@`5&GU#R2nVcO|1`+VOMKAV0N|sM_449x_8`uO6gZi z+YwqVej7x7Af(l^*s9i1x>{r7NX+j{JE<1u(}Y^?FlxoZ(Xy$fRV4)@9LXpoEFM=FtO1G zghf^CI~7>19>Tjp#OuUzt$@M@jCs3eFNo>zaI4O9jr@sZ#epA81#D+tJT9BdNZUB z(COc1qn}6mDF*$TaP(h=>cT$PNctPC_8F+tKi5XTjr2b|K`%LeXc+&z!k}L)>F==8 zAEncOzfYijc98z{2K~#z(eGo>|6bDHX|>NFo&FpfedPwwKgpoKDjfaCpt`V6FJiTN zxHH+b{|P$%V{PNDC+hU4*yyhzeTPB+(s1-W2L10Oea%YWsnb8$D^Pxo^zX%8u#oJYaP-eG z=tmQ))w5rxFXF>VI{n*h^b;FFKgFP59ghC1P+drIwWNQ*O8;b?{<${#Gf4k4ISmlo zoP0(Y`(I(uFOu{RS?zO*PXGPrK>I8w{p$_-mxQC=$DrRS=_}}GP5s<39-#p;pk5^=wBu2$64h+ zRj2=4&p`X+k$#au|AI5a*uN8^3HfJ8`u(l+Pt)m7vC(fJeTPARdN}$XgZ{UYe!P`_ zqE7!{k3jj?kp8{$E-5jH(NAm!{m*a{D@uCc3-(cux%PF1LBCMaA7Q0GQm6mE8fc#xq<_6Z|J`0; z==U+`p9?9t8#P|+ay~;+=0@jucq+axVq?)p7K05IPlZ!qHMAFsHpNg7%W`Hr{l>$G zf*y!u3VZm`_PD}h@W)ryp|JK#RG@-_`zqWC3%bvTQ}A?y`)PQv#IN9)y6INfD6gen z83yH5J;Ru8Crl#vUk1w1D;bi?W!kS<-$Jj~A_Dc&+2TW-#MVv-#(f8Ff(6}c!dW%n z;C_2=t6sx&y|&PxmLai5y($gL#o_ci1fc}~#h?tmT&7+zrd|gefqLm|p%<~W^kCfg z>oWw+<>BHlU#m*Y!HuM%hKZ+HprE%G?l6FPB02U!>(( z&vhy3hnLa4gkEV%V14Oqp_f?2x|M2MtSY!C8{8L%)9V!oE@b=?t=r$eM(BEtwNai& zy}p(=>#@zRh109rp!}YU8)c?m3+Y}$uL@rmQvSH1vxQ#7)?T987QF=bGK2dYeZ$n( zPY_(_HB_|jr-fBp@f3)bi1IF`UH}UPkA(2?Iqe468lNb4_;Umtj}Gx^>k{O;e|?va z8#sx9cYySbV#Xd}A3h~=bPcGCgFf8&_cPb5TPLQOLy>b0|EIzrFr$4nEQw`g?GAm4 zp!wgl53DkSJK!wkc^nnMCli$Izz0v3gQ%!^FBV($@VQ*0Ap?XykA z?nCW*&lCdtr@p>_mb#e}E0%em*yf^NQGf@#0A)c0m_%l_wc=#LK1DS2zVZ>+v~2M{ zBD~(uzA_B^VZVbt57w!=e}*ZLbqED`Ag84aJeE1Ws4xY`Qc)hP1>SiE*C8zPZ;Hl) zwZMI03bar(9*pSSn#D)X@K7c8=N-qm@W7v4*q?{3w|gNtKC;~QXCU>oFx2A@f;taY zagGjCAn!L6;K5qp^QS3Za960Jg?O+Qcq&W*CoRN-wLn9d0_RW%9;^j&!xSi?4m?;3 z3=UJEkvj0e0u5+!Y}?=e57&Bl(YNd{!iW~motkpHX6wW}vxSe*TX^d4GHh)GdV87P z!hbMs#&0o3tR#jOz7uDog>OfT(CYB$q-f!M#>L*idcqm(E&OK`u($9}0c)=}TlmJq zUD$C5quc2{&$1O1)H%=C;&?yH+8XSQge8KW_xKN_T{a@7Sc5{-0WY5XdLFuazeUQ< zC?^P`^Mf!tm9Qu>TC*FBFmF83PAW18seRyUmw69ieMzcRCxy@2NG}c$d@Kd;`VE4Q zL@FpsJW=)#B_@b&HDF{n;$=Yunn2jG6zJVuA#^(}2&tokkn%y6_Q538qLT`)TcHNQ zPp06LkqWBYIHLFt0c8|Yc4Oa$vUc5WAak38R1)?RBl@BDC8Gx6;CNRm&*JS#cY=vN zQsC=o1Z<%Wa~f}biEVo98wgZMf$l}hE|3~TpdW#37wBWcey2dQC=dewRCuhN+VjA& zt9O48YBv(5k=pU3hB&bbYB5f1`>#}&*Ls{-gQSQPISlSRkiEtZ7H^h-3lf(uXd;?R zWO$%8&Vh$NM11sALt=&(@KxFi;!p@PQeXZe!^8FgrT%`9Jm;;c0N8R$ARUj3-8arNCR{W2w!}KdVJ_0 zP(DBk6EO3HEc)irJYAh*bahy~J+1{$N^j9N;bs$0=r@b+c$!TXu@X3#B;YXeKj12c zk?CN-VJ3I}*8^x*?LFk%Rog_EO|>TgYk$#HZ7C|;8&B0(*T7|6fAF|emDfgHBPr_o z9*+8WFnS#Wjm5FS{)&Td7=*W>zxVRsr(oGV z`Y>~OFxXPgk-D6;sV*l?C*&ObGYarvvJDjWNW2iXAWyLFM{4q|t^^t{F&? zNN^R$6Hg`HXyTcU<~MKCiA4GQXP_K|l-;-HqpV%UQo!sgUP##aRB_K+QlM^~T2c^F z1A>t1LD+bbYST#t*KOwK5PT>FzXYkEx}_4O2T{f#W!G&bGVHpw0A|;%p0E)V$hTAI z=KWkmw~2gMK<*PiUvKKC|0pBZUIvxb`oQSD2>oVz5j_-jz+6VI{Wu!~BYY`L(Wd{T z4~!nim0jE)kZBk9L&8AS8o{;$)_x7dl@aV=nA1D*M+nY40-&<_{IpuP7_N=xgQRFa zmtr95&BR78e^<(j1%@tV<-@?UDJ)R19xH^`!(O78-uwt zMjXzDF`k7iS`L_qQ6g{h?Z)^FnRa8mP8g_Ijqw6t?aih!?xZgy?3c#a3sf58tB0ko zyw;7uni0k*XUyY47=utgzhUI|W$(v>#SWL4b{MDI;TXKniZ}UA$kv{ycV`=M=Pm5e zO#R^$%%vSh;2dCw445&b9X>#&-45Fb0~M zbvv+9gdN1|wt4XTFWx_I*ddCy7d+4oUkNLDqEOT{#W}hudgDDQG(|Ep+skxQJd0ap zVTz;(X$t1j6nTP#oaqrRF&35#X^J!a&7piiZGezs)qoB6x`RA3s2Lv=1k6 zX^OWWl$!EdHwEiOnBoe$1P?TYLMU9qG}LeCP`y0R6rUj1KWq3;-3aRU(g-t<+rCOS z!WvAP2qSO}kd&EABUIxo{Q(;m&WHU3*4fPUtQSb)X{tPE#V#aXJGYZ5%o<{zkol+t zE;<<87kwwWHz3R8KFZ*pE9hd%iDz(c69i-xaO;q^7Q%}vNWzu9RZfI>_9{4!Fi^2p z!6|^XuQsb73eM?W`kkx-ip(m=eLz~6*LoGu6ru_c(tKVXSZB$4fx&tKUh4(&i~^ev zqW}+fiSK|qc8Qw_LrJT|R|3|4w<&QxtnK}%Q%dM07x=&kgWE40fbakx=scG)HXw^i zU3I@K!fRbdT39G`FO}kfr9_uEgIOSK>U|dp^FS0amP*!9Oikw#g$Mc&*^q$AV>)wr zAOW%MDR|%7dk?DmbvTk&L5id{>ZEq^%qKU$`4A|N66Is0jzDi6k*^{08X^aycRNuo zCQ9p(>HV9?!-za16unHM^d`zbk4$d>kw5+x$Z?_Q{fb!uZ5>g*=lno;`@KTsRw8c< zHJR`*QT|Pol}Dy`6_Ni*Lr?{p%6?*+0m6uq8Ad4(v4cSBze9#Hjk z;%8U+t5A1h!z-FE#J#$gha(v15zy2sNc|pC59SaHiQ*>80-FG5hY{dm3c!QCr9TNu z_LjZ^F#Bk45n)I3KKLIUQjNJ*<=5Wh7uT-*63FT0!Or`vAiPfu!n+S)1Ihc=?S3|F z!G?jGs-c>87VfmYGo_WiKeNNIRL`+tB%DAAd9W7v>TW+TyGrv}fCsxuGa!Q9oLPX` z%{iK|Q>fDBZQZF12tuVt5Gp@?11e{b$^uZKZzTrxnS4xLXvYBA>*6E^@Pfy;(~0-} z*T8cUFPQuFB+4s9Is6__ECzmdrN4;XUXRfIc(B`R1cb8NYY?@v%`!z3b~ZJ=d8^bQ z!2QOA;oVB!JlJ_Z11;>lR{>_{y@arTlJ~LX9n2~}z6+~7^A)V}4N@?0H}|{!E`L$G zbtY0m9<1h^8K%HhEWm?ZrOzOOU8UCnv#az1VHs2@bxU_DD}qp26og7GVL7C7c(c$t zg!{eo4y<$Am$1&uNQHL4$B0)+ynBfk%+QO8GLb0P62;~J>Cw8T0q(c%&d|J&EU> z1_hxK9fZn(&q3u%Qn?mXLPx(zm@mFTI?l(3J?;FXR2RrZQpoN|H zeSq0{FD2{-@{T9(U{?A0O<3i*y|9WGsbJCXj@$i3?bgYrggjW?uPjV~1uVdWU8OG| zf?cII0kf;Lo-ppf^QO0Vr?N5#m6jk>>Iu7rRD3T9twTh=_uhbYmVXNCyoyxl==TKi zs)=_$@q!upKSaqUN&`_W?l<~2e-XQZA4iwu<-u;Rr$NbXuX_Nq+si}P3Tk?cu4#b# zy?ASA-fiT~gPr$zL3p1Mgm)}q50LjAFZ$U8M8CS%VU_bffmQyER514&5JtirO2~uN zoZe-AUUqX{!vZ|mRVsl9c5~(dW;f?J!k(Z?+t+ueGB^m8-a)7w{1{Z8C6xwH2_5~W zyawwW4`go^QlZ^1iFiML1iS>|1#`bXM0tZK4x(7x?}c0ZMeO!^0v)xN2fMw_hER5U zIjNN``t>F3C2D%xI;lZ``~5Qv@0-Y*2RrYz(8A990l@6M+X&<8ymtV32eZmAufQs6 z_rNNhNCk7h6*v2f+O3mI33;%Z^WR|#aNHvXKX#SAf(Uk%-U7_7(o2NBNtH5R=x&&M zf>7}Up)!}SU8LdwmC)|D>t$Hy&JSUo*O3bCeoqpwhIkJWFPNd15G9u=Hz4JIUy0>= zX5xMC_^i*OINX%;afgmY@ksI4v1pK)a_T7{Ci;^j%>2=2YMhx$HdAS4iu+zA?JP5u zXQn2ZscB~F5;HZ!OqG}^?q8JjD$G>1nX1!MkeUms60yNdHJPc!e#AC^a=Aab%AZ{A zPpAQ9cJn^fA&r!u~ihW;9ImC$qSI|ShUx~{L+u9`IGzo$wMZpf@yC4 zIu=DCMKEX@A51}wn%R6O1AoZo*E;#@SmZQQ31*64*WeF`BaqZdj`l;7{YjTUndeVV z^CxHclV$#7wLdw}pKS9dm-~~e{K?h+!3p`aRl6_^&G%6!C#S5%&N$4IKe$?) zp4cHy&uUbtC4EFz>r2rn;c88v=W2avIiVftqj6HUrza~N=|YQ+^hubURbJpbw>#3O z(E=Um!Uspb-H|>+SfnMr3@;Y#NUxS<=aG(+Y5k>&9qHIIE2xX70y@&;uEyo|^b-7+ z;(sO!D0}bV3qT)~ax||327lX|DLgWsSc{XZqj@7Q@!Q@^b&&DwZk&*@mzVf$Z>ECC zcgGN`ViCGGC=ZEs$Jq|(vc#uI)UoH%rd*I)B;dvl9CEtaSM$dgB&;-%V; zZG`2u;qP- z#MO~Hf`~Y~^mDR)&ePA+OBM~=V-VRCdF&pN1`X=Dx8F*`T_R#%Vj^FC=m4ygbizBNA-*XhA)mOTP zc1~=~>Io{l4cK^>lU+xo_!xQX)uX3- zN1=viwuQv6B|Z@KGFl{Eq~+7;-iwi8>+j|Fy|P+|%*l5ulU=RDuEn{)-Et5wUC$c& z6~CL~+H^Ql`{oVM!2A52Blj5ZXTuhJy#F-~6nySrX;6eE<>tirVPDelzt!kpYmg&Db%bJA=u_^dlCWqejT;i>;Nfra9VnX zJI75fR*Evy3Q5V&EOh6jrBAa1yVVgf?JDcckQJcI)ZF}ug=y~W^qc~l03s(dKX-Co zH(DagJqZT4bIHug9iIX>&d4Z852%OqDd{74T7(%Lhf?#kjMC9k5W zZe~@@Z0F3X`pOa(5bTnisYR8QRdvpyk`iZCwIHS>IkQV=SJljS7FARj9E8BVm0@LP zI+bC=%1SG$m9+BO>WZTI&bqSlT4#ABsS9a(E0-75mKPV+*430(&N~16qU!3>$`ZvX zekiQeTyDdtDFvb8QUymLq`SmfTUJ$HQR2M3)Hxyq*jZPFh+bPXtF*UYJEgcHzqF#X zsJ1k_} z%v6lLB72^0+3fP#S{T3?kbkmrG74AB#DBdqM=3~~=$unpQ>#}Z46f*x7OhR0m61s! zlw4RdHOo1qBxy*Al9HV^EF~{*80ZX}GcRd=(p5>dut8>Nol)WLv|7In0y0ozW=@9i zD5_jhRaD}xDL*5=Wr$c2yL}3rv9Q@V@fcj0LNIW!0J>gDs!o4N7?o)0*dXkPG;7{Y6 zrf=v(4Cy2p)J`x;`K7g`b#?TZSv6Jl)y~@Tt4f6v7FVRA@z=qnXU{IGEV;10w7%5T zyr6t`bp^eBPE|!c+i&4SQF*mV#TE97vs9~L7Tu7d76bHh_U_fIm*~};S-F`sy}+>> z{&H5obV)tSj896f5G|*wvY@oOsHW)h05PR0jM}oWR%ZE~*H5cN1{-P5T0b1f-C)HP z85Q-lWqza*(dL!uhye=_0~Vl}U20xks$8jDSzlgS$Cg(dHzH zF{|tAoCujU<@1!X((+klb&5-zf?~{8W*607L7xaDL#RQe9AXt0RTL132NIfFRddBK z;ZRCmQEe@{jo!Ljuujt}ucSBSqi-2sKXWG9tkI&<=am-2pWuy#*Pt;e&hom_8WH7d zr5L)%LO7>jEZuKX`Rvlf+S1~o3cFSRaH}vSD=RlublBzy5@TilfQOM$hC9oh!?8<7 zZue0zhOQWBX6B@)VW2v`Q18+G0|<_&3JWIj8%%}i`T4o|9P<{6;Y0yQT80#opE2eN z7$2vl7o_I9^Csoy3$X%+1%(B!l>Br#(#cQ9cp^OqgL{qxh^G$#@=`E5<(NdM2w^#l zwbJn;q~exflFSe6H4C!~i>oT(KlR0R&dU0V3Mad1(Hlvzl;LGnvrC7U&Y4wOIe$2u z55C84;H;z(!)uEe6l#Z8p|2>e7>>V@N%Mwx3l+nYk^=t9fm~lwIvivs*OgZcPtD>f zsjRx1EY0AXK8HP_-YLv2hfB+$ODcvl=VQkt2Q)zj7H2whCTC?yzbCyRwdXqv3#+SY zor!bH%8SbwXNO!>Kh*gzF{BvwFFHFKabT`a|9oetUd}d-3W0A4;r0STex1IIE&g63 z$&L&?r1IyV@9*eEK+uN@)QUAX6A_~hFBkZ$_$5{K2$(GJ*XS2j!m@Rxvr20OlfPUJ zA)=xRt-PeDuBdK)H8=_%Lw6%$b4}Ikl*$rm2BVYLLqu;UnBb?^w=o=&!B!tk)+RZX zs>+J_&XW;)Pj(`NoQwc06NthmJ8SE!tE+0zarX|2XUR7B3>(H+>_kWC95zghR^+%< z4nAhr*VTjYd}rxA_;PLe9AQGMoWq96@yp*PU>!!leG>ico%grKsGRKt9lY*1fq1yB zNkcqxr}>c~S!n$1lxbp0VDhB&!pXSnEBG7b{I^}^4O?LGBk;A|(OEc$mN!pC#C!z2 zDY&zSrc>Q%lU#+V2$}ycM)%u+|P!_<7tSp+Q@nxBwwR9YeR z6RvOJ7LK2_oGaa(>9Tms5zC$Rf8?i^`@ix_OV69+`iuI}9{Mbrr2<0Rk@y%QVFqlH zK4D{J6s8<}YJ5sTdLbPW-i48(rC^wH7#gOfOiC$aAeZI3d1bhID8Ps{3-haah}Jn7 zZfmdz8{cKa{`=+h*(ov8o095E7jm1LhSt+w4kK+jFJ+cxNBP&w2O8-wmE&Yu0Jm`b z=m)SU2J{#sTc+;<354bwP~O)1!|~-Eio9tF?Ad9Zzg8}YnqllJZ&gxr^YhVd^0q;r zh5xI1`fr;iVRANKtP*5RCM7F{(00PG#GR6Zxz=Z)8Z*6r6`ynUe%#P><@nE!DcmWj-SXX7yum0c zpA+r20-@L4`z9NgzmNe??13|oJ}%G(Bmf4jVGF!;C2 z+XjPwyS!~M__xd327`aQylpTTw!Bz5!;MFV8$){ZHdtY__+sdK$C^wI?v7G}J8Q3jCn+8`gd>Nh!FZ zyt-N}fYoxs8!JO%D$H2*;e->HH|Ex0PWdFz&CHuDXUe!#l!W=k3M_Zda$a6kd_?ZgLqXaH3rgClWZ46(Z;R$}64b^!!wsNMab_ z$ATY%IVMGsq&csT~x|m#I87f1U zTs8}~X2u1ak*1w(^lFVoHc3q23X0+Bjrtq(MhSYw6$Pbrk~x;ejc^{|$9h}|j@R8E zg5z2Jj0D5-O~ZOqC702y@jN(P^0%UgqFaqU5v3(!%}&=xUoFc?%P*>&RT^M_RNd5~ z^16(w8nHTpIdgp5)h0PAqrJdO6QsnOgF65%Eb2b*8lAHxt zMyM&R#Toz>%EZ18Nu4cfidfclBD=n#4j8kF%CWA7Yb80 zTxQU>D3s5}J`=H`Lk2Rq*Xk9S#T^p?cwAb*mL060LIa~OPzpq6F<)$);36`%`haD5 zB|AClT3j`|8vcPLSh2qXU7iR)y7G!5HxW3E5g4`{5T(>vk44?-BY2TBY@~Cp=xK9A zUn_TE7zLpE95`SRmX$80k~%%9%wMLa>u#e(Q}{Q1_!o(V$@JQj7C{r$?I!Dm+9MPW zGvl4pvX{ZBw=r3F&S6d#J!0Gz>-hRmG}k|X_4E7dyvWa7_Y7?C!2&qZL*fSw(0>nq zl#GLVXwf@3BLXK@%S|z|Ra@l1Zm!(R|BdZ2CZ@Q9F}hh(zwUTJLrWNZYm+v_@Y|z1 z{eZ3@4Ben1F${jtP-Nl<+eed#)@N?6G2&wItzvx3Ct_{M6(i5ZNI!LU2?j58$aWEO>(kT3}uy$x2B&}zG?&MUlMsg zACXd&Bc#(F?(&_D7tp7z?(#L^T!-`HfY&_&I35(4 z^aOV5GJ)_{O7xTLXNsie7-+%PAJ496)N;B_4yvaxuhI$ocsS9a6D_4c7@Nk0Oy3+KHqCN zKTzxQCG-UU2EcKC0Oz}KPQz1zU*fzK&q{wuz3_1QILNmZuPdy8JO^+tfql=$L;ZK* z`~c2h;>^c?_;Xy9uQ=w#IR^Ia8QtdS8Atda{BL}y%Xc~v^j~m%@IyZF+wSi2ZR9g@ zF>x6&@e^YE%#ChPE*LQO^s^F9CJB~L!~e8tKA(83-eKWA?Y?oPKj|wr1p-9M~Re}XqTi% ze10-$^;Dm)C(ai90^r{W{0xDg?#E99Ui9N#e+ggc4}1;yM^hgCN73t#BwTbVI?qav4w6Baau z{-8Ib*yme=Ym0xuZn}SEq0wFo|4I$;ue6v%xCd-09gc;iKP&;gji7f5Nt#kI(K#UGRwTKfMmXK-VkX^auLXdGM(qb&w+b>j3EG&Gh+dY5y4k z^q^BWJGo=pA|ek3CzxTiKl`zH6+YjzzMy-309}9E4)9w?WKxfyf|`EIxc4IZm4R%j zcm03$IOgit!*yfi(x_H-Nsoz*JtLp$(HnhnYM)5=(MfR^%h9y7Zw1RA#yE-g_+pK) zhu-(O7)Ntr8XS?C5iyG$5veh4jo)y3!)zBKj^sqdtaU_8h}j78J3xM?Bb(u6ad7ue>k~OM3I)=DKVHWt8H{_t ze-g$kWwAcrk4P)oW?b~g8o#ecyt)Okc{_Tr*U*K%j+pQU6nQfyTJ&-Ydn{2~qn1YA z7$M@#wQLMOm)O7wrG^vM~p`v$9O#tZ?O|?N3_E-jPT@7w0~SAkMa3N zAr3^Xl=Q>&$yqUNQIYk*`?-lm{ANAzYaEGqRyhX)+^D@rq?;MDGY-98kZuOi`%KWg z1oWcs?DBOX9d-B_tDj`W>{laJnSPSeXAp26fG~^UC+8w9<9xNLcYlDkB7ZIDoyvTS zW-mo<8j>FY-iC6eg!}!=-ws`j0WOXwvtyELf=MjK@drfsd@>5jH=@+9qLjT+cn1KG2A*^V8iqq%s48DOhA7I-5$ZQl$~O`BfFQC* zU6rUPiz3vm{glHIV=&B%Qh)8Iyc2aVvN!Zpf9t0#jXqCNZttaj(NFoJ*GkkbO3j*f z6Pisu3@QJ#)e7xz)#r4Wr#V;=3J?2l4k>!a*cv3S0-hk9opG<9`FUu99WiWkpp zj8-4)t32KdTp#Vtg17b|%7cAWlzg@iQMUC_LH3V6PQhxziw^adp33JA6|WMzCqn(X zr}9+9Sw2;{A@bk9tI9o5>MyGDc$A87%(tqL;cZm~_-j=K%Ap|T7bkzlnjWiQgJUS2dDAl%@X zi#xp~4)wm{luZtG`*BLALw)!-<*^8LO}z3_q`Koc<+n)nmE)BAqtunijZ#+}r)*Z$ z$KsXid#De^E06V1?~PYFd#HzxrHS_gzo)t(Ub!Cz0z6uMHeT7>OMM6Uz125>-&=hd z@ZRdyc;%fw9qgthouaHr`2GY%xnumet#0Mh z@ef6IUZUKbq5gb{vOQzD0~bGJsvo+P`(04udzX43Pr2Kz?wqdda;q<2rYxSI?#x$S zn4qq`RJnbk`rr&@$3%6*h01Re)vqs69?eo;nyl=}QkNAfOS07!1LHNKR$rW= ztj-bi-_3mjJ<@Gc)U8vLMN>N?!19KR)ekOK{&%tZ$;HZ})7AH;D^FhX9lT)E3{k>c zr2g-6<&De5#ld2AU5WC3iMZ%0RX5L6KAR~nTFTVd%apInZX?MhSEz4QC@)>1E}E_U zbj3J8KC4h4ovl1nqrO$6JY1{ZQm1@W3mxyTQ~y`5?5;y{TfMq^u5z$m-8olzY!22K z9-61#JYU%}4|(^^SD(5{*)<=@tyifJFHnBDN_}&I^61q-d+-|drfZcCuR-3u*Q#r- zQ{KN8$t~BZ4_)t24qT_cfzu-m4q!fTgF|g=bSNL(-~h&|Mu+-jlS6s0kyo3W9O{D$ z9m>x*zmD_jW|p{rkwd+4u|wIth&cBwcBubFRld8JSDRWK>H{8!@)ORl;r#G_SmHi3 znWe1`W!Dno+}-L>pIGWp-f894#-$GR{x*m5BhIhl{Lqao0WUx9RaL#AxAJ{Yby08S z#ppDo5A>D|b=0ptc8EV#?N*eV`lt>4l(l`3Z0f6i-dB07 zZ;>GQYXn;B&1Wf_Bh=Qjlsgc*&QkVAst1yk|3<0bB`G(n>RU<57pl4@NqM6OA>Z^+ zzZs=;^kjN(Pj&k!r7e0O($6E1jZ*eUpNaJ9UTV`QWo_?7q`&K}-hP&Hcb_#;xcD$u z{bHnYa~uob6Q}MOseBwa?d6fmjYo|Gr0b{wNI!72>Kmbaa`ZT)pX^U~!!at-YmXVT zX{2)a7aJnR&;e9MbDf9`VO;<@HlY z=fElId&8CO2`u|*!dSwGF#X{W^_k(y;!_VhaIqv&Jvdz1nD{?l{5DiQn4~;$`ZT0l z&rtt7Ls@?YrFeT7>gxTG)CQG#e^lxf=*!?-8OpB?^uAvcU4CRZyZvlIKjQVJX z(iw9<(ktTBmQ3YP97^8t549;%S$q_dj~u18XDEk{nu!UXqt(|ll#TrnX};~RzMY}$ zI0n7UFUP36GL&t{0=n;5^}bAH!*NJ@k5eDcR9=Wj@{4$Nd8YF0@koApyt*<|d1L^R z?+j3H$y8PjM6zR`dRvC_+d!!MpF!%aE@kB)^`}gwQ>1S=L47k**?a=wADo~*o~hh1 znCUfx)dyY5HzNJ>V0EQSS$5*29WLbsCv&$s)z31NUz}-3?>p&cxccG*^}!6KBjI^y z+I*_IDnohpR2LllH1(ql<(bnC1H5*qx+O!=h8CggJzd?Ap){P)g{Zb_)Cr*S$|&`| z@k-NK=VIh+u3Soy3%ltiu7~mR4mO<9y?cEI9_QvkN8iYr@lH?`TjiOzc~6% z2avuPqwYOlx##?GNN*ntZSNX42&FcUQ(yb1(lu^-1TJ2=Ky64?9!pk{em+^njK??0 z_c7y^)c>!&tB;PWxbCwnk7R;n%LdCrHfGbLK1mKNyDMoWQ%asCud%>58%sDYZS%5T zt>gtstFyZjHvJL=oG;=~<2Wt}i7~{4khD(Y7%(4J0vOXY1oGh$LJ|?E@oB;#gkVT| zfd1~xoz;79-ve7s|L8fMbENy;Z|CdXxpU{vym@!z_LSJwb#v&I<>IR;%H!FURM4lp z<<^VD6Wvr0n>96su!6~upKXF;x&`2NSU%Szo(SXeH(`0CNqoHqm;YWPC!54)XW(-040#v>zC^Ci zN%^BDv5h45u3GtYlX#ZOy0%XKyh(hm?rM~^w_g6BNxWE(GOnLZ!~7$&arxYAd0&&b zc@8f3&yi0yiI?bdOM`q05gI5$gWTIB9-G??$($?qHHm+k2Xx;&d3Td|dLGJn9kSgd zzI@te&=z=}dA|<-0b&4~;~6P#l#k#V=*L14aCDWttwD?-`He888pQqdyibMX?gnu! zG4vMjwn2QCE{}%gt_HD{c>6Ckaz}&s30)qmk=q-@PGa-@GvrMT;(5BfR?06lh<_$F zLsx&XLA*kjTWjUc26104E`L%hpPeHP%zO~@Bh+PA=z^T$^G6qm$3nQgH!NRl6#pH5 z0sR^_(H9not#w~Uz4y-~3+o4`fTu5>B5#8Ubt-cE>8bK)qqt_)b%_1RIr2!O*hRhl zK!e=5Kpbnh6xPIC`Naib&U?{>r^)9U#eJuJ2H8D!y8LCM*fAe`+BskDULYQ&%O~c` zhZ@D_8*zDSqujed{H$@Fs%GL}Lw`V*gFXGdI`KN$mDknE7wW{mTIdTr(y!{om&l~N zy-vPZCyuan`Tjj2`CHhkVY$Cfe7r`&$LkMkE^*_lve?SuOrj)PNiLcfE zHLmZZ)EMGa!8^hwxjQ1R4#_ttcV~YybR;BxAHwCyaFPQjIdGB#CpmDE11C9fk^}!= zIFRnLR8`ZRn#TBEf%pe|RRu5Ygh@)l;C=P(`>Xk$UeQokJ6lsM8LF`EJWElICiY&j zOx@#^oeGEU&nSxNlL}@}38o(^JjuUf;-SJ!rmwv5W`Qnuge@Y%m}cXR`W_P?72YHD zJzgcKZ~-iO3Ro&v0hYW9^O+g~V_$~}T|Z`nn8_K@Iy?nTF;u|pOoepJQqw6Yx_{to{X4_|=j15!-&@djZr9yR_c1-l z^f1#SOph_G{WD$vMy5-cCYi2ax|V4_(=n!-nQmvgo9RBL2bmsbdW7jQrnTpD`AnBE zO)_1CQ>^a#^qOl$v~%V)ZTX_Dy*rfZq@GaX~Pndx?> zyP57|dXVX1rbn0_V_JIwm(O$w(EpXnIW%}lp5-OY3#(}PS8Gd;re7}Hvd%V)ZTX_Dy*rfZq@GaX~P zndx?>yP57|dXVX1rbn0_V_Lh6%V)ZTX_Bcp;V}Mw>T|KjfzCWid7PL_PHZmxNsg+!_}|Bq(O7G=?Y$b0Bx3PoEE!wsD#3eGgVC-dYTr###wK>@p@tT$ZL(5Ez-XDoXWWG&Zj|4usz_$=ZaZ38S$I?rAEvXJ3sG- zpC)KKqBqpwZMLy5oWIqG`NG&0ZoE-U`iyrqVuA1*2Wv!=cM~@zR^#AIRPciR${JCXm|MVo#@qtEoWt7U93?w5C8@(OwI} zz4jsPF9?Z+V*F8EPpSun9{h!T_<0p@%E#E96qmwg{Dpkl<;%eRp6g!Vl+SpVg)6+; zsZm5c4?_4yF0>Tzn1;n#RLl=Q$arj}j=zfIpBm=;8NZS7HH_czI~`x`4aRT(#(1RL zjklm)RBr4djh}|N6t*#*W?b#{1^yc2lb2W&O?+Sdhth??NBQhoZ7IYY)bSz4V;|Hw zy@;gnamF9-vD7{7m8I|<#wUMase9TZL*ZN)%9MY))1q6kMEw`y1Hh?$ulBSrwX-pB{kIUSGd_tN{Ek6%M)j8%q^Rdk!hkHwAwA z&>l-a{A--gkf*Y=dYPw|Uwz0emj-m5h_VqcCf(e|@?U->;ru2*B?Pz~^G3=9dp%Ac+X-yy>_(;j0AH>>zl_QM+* z7P|uS`L6(c4kq@*hjkt~Spr-=SG_Km_GVM~5b$|;-W&E>>R#F2ii+d;Z@FOob4Z04`Ln!)vund0XXfU z^^5 z^gQW@@lK1PiTA1h*lrMz|Ly?%o&fxJ0eA{bqjE=fTB<}EzbR}~xW0D&-x7e|37noc zoptNmgTU!|C;53JoLM6bSBk&$6-yyMX1pCZl2nB?0eCI|e+dmY538@<`R3ID@qZqG z&%pw0q^dSTybHKrJ(Ivmt{!K7YviXVAbt)w@gXup^X*+BUD!6pN4Q@Y`8+>B4(|xa z|N9D`EhZnaRL1Xv{XyY30r6*Hf!go6;=t#k9mOLVmA--NxiTRBUje6fZ#bmm8$JAw z3dch*f7(=TnC}GSbB5{{9LYh>3Bdm$0RL(L{zw4+0&ubqy>{0hfcweqJut zAbuQzO6@hm?KQ4)7b36v5#Kn-mH_;Dh0hiiueanl&8q?N6Bq~l_;5`C{sZ8|&+%iH z%2|CQiK|+?$nnc38i21U>^Oxd~ASh<`0`zjpjKa2g+XvOSD$rNWb(4}IL~ z4A{8jThE2)TBL_3f2!ltw=5LSV?1)ZrS9q5D+-q@Twgo?e=`98Q2_qy0K5hUj2}N2 z2jC+C_%`4XX9f9sPeA-b0r(dIM;%wlgZ}0Z#+f~4+y<&OYU(~lyDEn@e zPj!VRIsb;oE%hUP>qO!8fcT3R>vo^;w0jG1YRAZPmP$`{359jQBkIciZ$QPz_hOdi z`e8G0dM@K{0b;MhtpWMm&-u7A!OhPD@N>`dZ{P6%d|Lp1R{&lY(fUNNJ{;$Y^#G@K z@y-*=j0?`6e37;h2VKds|1Fe=Q|B2E6C>w~f$AQ$6H;tW?ftOR4m9Sk{xY^Nxq; zh)+*KyAsr$)?D1vVrx-v?QpWaoet;9lsBZtoM<9(rLXV1R64yBdc{^Q@SF+a9qQb! z{tP^_sbgkvc3|rAbUNNf(T&40JKwjaCvFcpQg)OJ==AV6MDjHqHuc)+73eS>`9jG? zpUL$swNnnZCEJC~Uh?hI=z6=`?)g9`PUFO>RCE+U%f=Q|rl-8B!?yDUQex0+E9jsC z8$Q&bQnh|H*B7q7)K0DB2rHMbA-!2D*xJdHu$TV@9Hw@60qwn{tCE3BHTL002T&@b zTU2VO&%LtM38GZAlB0e*yWtAUx^Pu@XP2FbwnSS+e>pceI0Yr5B;fEs$I!ZUHpL?S zS!jW?zf{)!L&a~F@@!~X?A@8pB1D>adzI>LhhwYJk*-w`l49=4{BWPU>`H#xpW1O` zm`JD`s|s>SVLYka<*`$=w?&mxz<$hq|L`9|c8qq;2OKih%Tgv$)zZm5tY?4hJp$S)wIrz0|&zEfSgy<99p}zSU3#c$UNN}~a ztE{TJM-xyPIr7RYyfSZ(!0~8;S{t=HQJ-AxWINzoNYX1VW~l;+)s@Z&<>SZC6D@Hm z#&mt?D3$3Jb1JkVfbcX(z9L-`$cG2IFeMO`S}w1urQ5@+J8Aa^a1d%=e%K~AjVRjb zCA8#<>f(CaM&~G0E!)JZfbJJ>Q4SC{tjkq!7;-elr`(F;ItP8bijRHMYxl-r z;UsOBKtTm&H)d?(wcW-MU^#UzsbzJoTZi)_?0B@rU%1t1S$R^Izqfc2(o5&bD*u41 z^251QKAnas?fX#i7^&+#4voDEy<0t5B4*p=4RB#-gINRR4YsnpyuAUU;W{?VB^pWq zr_~K`1nt0Z3gdaJn^(FYR2~25DYdtII%6!_p%gS0*_V+zHumaUDOY#XI=Qpl-ES)0 zx#C~a^Ayu~27m3S>)xhyvsSP{`rT7|C6iU1qNhNr3f^dQ#d7yh3YL7K*V9Z(NzQr- zZT)njX|;CjUAsCC@q&}Y%=-LTJWet)2xV>4QGUo7ecC1b)YIZiEAm(vz#!uo4YV0X+|G-&0X1K28|jN6fnNc9&cw0O&ca%d2Zu(O~VE zslcISu2l`C=zql8N}qR4V@Sno#yc7)eeqae6ZQ5^N2nd;R~s)>1yZIXo2D8euFNza zT^eu2DWILo%Xv4sD%RDBbFNs4e&(=zPR!DBTbHDG`<3zMQ+!bL%%}4AKt?(iXudo$IuU)IOQC|Yruo9sTtam?WMO(A{c^cf; zP|&@xxA%|P25;5ws{6a$qPwcuU&`eib1c;@VCCVU3YvV`y!<}$;hy0@9ez$ zQn{-0z#PrOm$kP~^<}&Yh0Gb(2u#LaDgSnZFvdZLIyr9V#?$N7tyq`sOZm)$h61dT zx0@}ul8Kv2L#TaBp}6hVcS?)p+o;D3qy~{Dp~iw9-1p^+ZVu~FUVAaG8j^gZLmwC9 zmJXX{u+$l8p=4@&uNpnn1O^ct?J)H;m_G*?c074D_CYzRPF&Ab*^7>{O^1kgkW6Zc zro)iP7NH5MBWm0|h_xhKXE9@u4ChMd=^CHKjIM;fnDMNzVEF(hpy2_QvRjSC2cfnL z{WgZdft=kxR7UYvx!NF_pj920js6y6yela^hhs9|+6^i39hn>~KTgqyjeCjB>I9<` zYV9*Z*@OA@Svy-ARhm!N=6DPh-?lsPI<}KT_u^PCr-ls&Locr($aG-2H>2xlz>XZPzcM ztgwEmMk&9cCJ~3x<7@G#4)8Whjac~0Ws0eaB7|X#m4Cx7k4+q0ab=&YQOxzoRD;4` zDV4A8nfW$a1|IN;6UoIiEJZ!0P(Q4ecB80yMrO`^b2GlX_Yy2KnfH%a|E9N(y>vE- zurD2^MPnE#rO_c`Ojj2csMTV#Lxzr1bI|-P3k_M$rPlb=)t*us^!5_ma9BN6eJiDv zN^%TmX%vdxX=5LCVB&yd5OT#f^`8zSfT@;BDPv4e@98aSYBgrw1H+>AW%5jg8-mxxox2;H#0rvWSwwETXuK4is=fFC9cwDV%7rpkBE|a~pVr#D+d(V(yT8 z7AtgWAsgEdNIj~|^xg@bhq=Byi-agRGnCIFU7<_`0--lLQ5b{*?>UF)m@D`0kZIpx zo0JIQ&+I$u2wS+rc@qp^+mb6#TxWO}Zty8eC4Fz$!X11UbNYz7oBC(I!!Sr37tZwM zu%&(~b9%EcTY>oJO}`x2F?`!^>c1mw0U^HziijJSVjH6hyiJtvgjkI$+7e~bo9{~B z!wJ@5e=!BxmPpTT(qr473TyG7wzQe_=KIre?5Ct){3Yl+ork$5`eFQ$uL6_ae5d-r z**aE{8E)!t;u(4cqLJ?db8o(9jhPLZoSy1WyfgLJxeIX>IWlzK>1EG+_KaQ2lAImpACJb@v%aF!&%k#t&fuwz_$hvzjX&yx6wLRChwj%gz5Mo; z|2)#*Q;kY`^Ih|koZcJXd;c=hQc0$K^ZoPWai)I_X-Py)dh?z1QB!|zSR=-!eDhp8 zJms73sb4#wpWnnsG!@Lfp&>+b*1z)I^@L78Vk#Dp{vDq5W?%XEgF5}YHR04kL@HPM zg7(`}T}(0Nd&%1#(g_cF1m2r|KB7^~n9i>9-=uG?Ao@b*s?!ic>fMyBW9a|qzpi60 yHVODg&#fxk^gU<#C%OJ3*q2TL+eroKNW*udKdFEFmA~?7UDEF&h~Nm`^8W|42Fy(W literal 0 HcmV?d00001 diff --git a/projects/rocprofiler/test/ocl/SimpleConvolution_Kernels.cl b/projects/rocprofiler/test/ocl/SimpleConvolution_Kernels.cl new file mode 100644 index 0000000000..08dcde881b --- /dev/null +++ b/projects/rocprofiler/test/ocl/SimpleConvolution_Kernels.cl @@ -0,0 +1,175 @@ +/********************************************************************** +Copyright ©2015 Advanced Micro Devices, Inc. All rights reserved. + +Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: + +• Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. +• Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or + other materials provided with the distribution. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED + WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY + DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS + OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING + NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +********************************************************************/ + +/** + * The kernel has two implementation of convolution. + * 1. Non-Separable Convolution + * 2. Separable Convolution +*/ + + +/** + * NonSeparableConvolution + * is where each pixel of the output image + * is the weighted sum of the neighbourhood pixels of the input image + * The neighbourhood is defined by the dimensions of the mask and + * weight of each neighbour is defined by the mask itself. + * @param input Padded Input matrix on which convolution is to be performed + * @param mask mask matrix using which convolution was to be performed + * @param output Output matrix after performing convolution + * @param inputDimensions dimensions of the input matrix + * @param maskDimensions dimensions of the mask matrix + * @param nExWidth Size of padded input width + */ + +__kernel void simpleNonSeparableConvolution(__global uint * input, + __global float * mask, + __global int * output, + const uint2 inputDimensions, + const uint2 maskDimensions, + const uint nExWidth) +{ + uint tid = get_global_id(0); + + uint width = inputDimensions.x; + uint height = inputDimensions.y; + + uint x = tid%width; + uint y = tid/width; + + uint maskWidth = maskDimensions.x; + uint maskHeight = maskDimensions.y; + + if(x >= width || y >= height) + return; + + /* + * initializing weighted sum value + */ + float sumFX = 0.0f; + int m = 0, n = 0; + + //performing weighted sum within the mask boundaries + for(uint j = y ; j < (y + maskHeight); ++j, m++) + { + n = 0; + for(uint i = x; i < (x + maskWidth); ++i, n++) + { + uint maskIndex = m * maskWidth + n; + uint index = j * nExWidth + i; + + sumFX += ((float)input[index] * mask[maskIndex]); + } + } + + sumFX += 0.5f; + output[tid] = (int)sumFX; +} + + + + +/** + * SeparableConvolution + * is product of 2 one-dimensional convolution. + * A 2-dimensional convolution operation is separated into 2 one one-dimensional convolution. + * SeparableConvolution is implemented in two passes. + * The first pass is called Row-wise convolution. + * And second pass is called Column-wise convolution. + */ + + /** + * First Pass - Row-wise convolution + * @param input Input matrix on which convolution is to be performed + * @param rowFilter rowFilter vector using which row-wise convolution was to be performed + * @param tmpOutput Output matrix after performing first pass convolution + * @param inputDimensions dimensions of the input matrix + * @param filterSize length of row filter vector + * @param exInputDimensions dimensions of padded input + */ + __kernel void simpleSeparableConvolutionPass1(__global uint * input, + __global float * rowFilter, + __global float * tmpOutput, + const uint2 inputDimensions, + const uint filterSize, + const uint2 exInputDimensions) +{ + int i = 0, cnt = 0; + + uint width = inputDimensions.x; + uint height = inputDimensions.y; + + uint tid = get_global_id(0); + uint x = tid%width; + uint y = tid/width; + + if(x >= width || y >= (height+filterSize-1)) + return; + + /* + * initializing weighted sum value + */ + float sum = 0.0f; + + for(uint i = x; i < (x + filterSize); ++i) { + sum = mad((float)input[y * exInputDimensions.x + i], rowFilter[cnt++], sum); + } + + /* Transposed save */ + tmpOutput[x * exInputDimensions.y + y] = sum; +} + +/** + * Second Pass - Column-wise convolution + * @param input Input matrix on which convolution is to be performed + * @param colFilter colFilter vector using which column-wise convolution was to be performed + * @param Output Output matrix after performing second pass convolution + * @param inputDimensions dimensions of the input matrix + * @param filterSize length of col filter vector + * @param exInputDimensions dimensions of padded input + */ + __kernel void simpleSeparableConvolutionPass2(__global float * input, + __global float * colFilter, + __global int * output, + const uint2 inputDimensions, + const uint filterSize, + const uint2 exInputDimensions) +{ + int i = 0, cnt = 0; + + uint width = inputDimensions.x; + uint height = inputDimensions.y; + + uint tid = get_global_id(0); + uint x = tid%height; + uint y = tid/height; + + if(y >= width || x >= height) + return; + + /* + * initializing wighted sum value + */ + float sum = 0.0f; + + for(uint i = x; i < (x + filterSize); ++i) { + sum = mad(input[y * exInputDimensions.y + i], colFilter[cnt++], sum); + } + + /* Tranposed save */ + sum += 0.5f; + output[x * width + y] = (int)sum; +} diff --git a/projects/rocprofiler/test/run.sh b/projects/rocprofiler/test/run.sh index ede510bc28..437db01ed0 100755 --- a/projects/rocprofiler/test/run.sh +++ b/projects/rocprofiler/test/run.sh @@ -74,7 +74,7 @@ export ROCP_TOOL_LIB=./test/libintercept_test.so export ROCP_KITER=50 export ROCP_DITER=50 export ROCP_AGENTS=1 -export ROCP_THRS=1 +export ROCP_THRS=3 eval_test "Intercepting usage model test" "../bin/run_tool.sh ./test/ctrl" ## Standalone sampling usage model test @@ -129,6 +129,12 @@ export ROCP_DITER=4 export ROCP_INPUT=input2.xml eval_test "libtool test, counter sets" ./test/ctrl +## OpenCL test + +export ROCP_OBJ_TRACKING=1 +export ROCP_INPUT=input1.xml +eval_test "libtool test, OpenCL sample" ./test/ocl/SimpleConvolution + #valgrind --leak-check=full $tbin #valgrind --tool=massif $tbin #ms_print massif.out. diff --git a/projects/rocprofiler/test/tool/tool.cpp b/projects/rocprofiler/test/tool/tool.cpp index 1b0ad02abc..c2fc493151 100644 --- a/projects/rocprofiler/test/tool/tool.cpp +++ b/projects/rocprofiler/test/tool/tool.cpp @@ -881,6 +881,8 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) } it = opts.find("trace-local"); if (it != opts.end()) { settings->trace_local = (it->second == "on"); } + it = opts.find("obj-tracking"); + if (it != opts.end()) { settings->code_obj_tracking = (it->second == "on"); } it = opts.find("memcopies"); if (it != opts.end()) { settings->memcopy_tracking = (it->second == "on"); } } @@ -901,6 +903,8 @@ extern "C" PUBLIC_API void OnLoadToolProp(rocprofiler_settings_t* settings) check_env_var("ROCP_TRACE_SIZE", settings->trace_size); // Set trace local buffer check_env_var("ROCP_TRACE_LOCAL", settings->trace_local); + // Set code objects tracking + check_env_var("ROCP_OBJ_TRACKING", settings->code_obj_tracking); // Set memcopies tracking check_env_var("ROCP_MCOPY_TRACKING", settings->memcopy_tracking);