From 9be4de53268d49665bc0d700f12f1192207fff79 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Thu, 9 Nov 2023 16:00:43 -0800 Subject: [PATCH 01/44] Upgrade to nvCOMP 3.0.4 (#13815) Update the nvCOMP version used for cuIO compression/decompression to 3.0.4. Authors: - Vukasin Milovanovic (https://github.com/vuule) - Bradley Dice (https://github.com/bdice) Approvers: - Bradley Dice (https://github.com/bdice) - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cudf/pull/13815 --- conda/environments/all_cuda-118_arch-x86_64.yaml | 2 +- conda/environments/all_cuda-120_arch-x86_64.yaml | 2 +- conda/recipes/libcudf/conda_build_config.yaml | 2 +- dependencies.yaml | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 4d5c56e4a7d..a479d517c24 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -60,7 +60,7 @@ dependencies: - numpy>=1.21,<1.25 - numpydoc - nvcc_linux-64=11.8 -- nvcomp==2.6.1 +- nvcomp==3.0.4 - nvtx>=0.2.1 - packaging - pandas>=1.3,<1.6.0dev0 diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index 8606932dae4..d1779aaeeac 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -59,7 +59,7 @@ dependencies: - numba>=0.57,<0.58 - numpy>=1.21,<1.25 - numpydoc -- nvcomp==2.6.1 +- nvcomp==3.0.4 - nvtx>=0.2.1 - packaging - pandas>=1.3,<1.6.0dev0 diff --git a/conda/recipes/libcudf/conda_build_config.yaml b/conda/recipes/libcudf/conda_build_config.yaml index 05b2135184b..fa06ed048b7 100644 --- a/conda/recipes/libcudf/conda_build_config.yaml +++ b/conda/recipes/libcudf/conda_build_config.yaml @@ -38,7 +38,7 @@ spdlog_version: - ">=1.11.0,<1.12" nvcomp_version: - - "=2.6.1" + - "=3.0.4" zlib_version: - ">=1.2.13" diff --git a/dependencies.yaml b/dependencies.yaml index 35d08239a4c..3850347aa63 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -244,7 +244,7 @@ dependencies: - libarrow-all==14.0.1.* - librdkafka>=1.9.0,<1.10.0a0 # Align nvcomp version with rapids-cmake - - nvcomp==2.6.1 + - nvcomp==3.0.4 - spdlog>=1.11.0,<1.12 build_wheels: common: From 87d2a36f04f431a8c5236d2aee723ec79b9dc5f9 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 9 Nov 2023 16:54:25 -0800 Subject: [PATCH 02/44] Remove Cython libcpp wrappers (#14382) All of these wrappers have now been upstreamed into Cython as of Cython 3.0.3. Contributes to #14023 Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - GALI PREM SAGAR (https://github.com/galipremsagar) - Bradley Dice (https://github.com/bdice) - Jake Awe (https://github.com/AyodeAwe) URL: https://github.com/rapidsai/cudf/pull/14382 --- .../all_cuda-118_arch-x86_64.yaml | 2 +- .../all_cuda-120_arch-x86_64.yaml | 2 +- conda/recipes/cudf/meta.yaml | 2 +- conda/recipes/cudf_kafka/meta.yaml | 2 +- dependencies.yaml | 2 +- python/cudf/cudf/_lib/column.pyx | 3 +- python/cudf/cudf/_lib/concat.pyx | 3 +- python/cudf/cudf/_lib/copying.pyx | 5 +- python/cudf/cudf/_lib/cpp/copying.pxd | 2 +- python/cudf/cudf/_lib/cpp/groupby.pxd | 4 +- python/cudf/cudf/_lib/cpp/io/orc.pxd | 2 +- python/cudf/cudf/_lib/cpp/io/parquet.pxd | 4 +- python/cudf/cudf/_lib/cpp/io/timezone.pxd | 2 +- python/cudf/cudf/_lib/cpp/libcpp/__init__.pxd | 0 python/cudf/cudf/_lib/cpp/libcpp/__init__.py | 0 .../cudf/cudf/_lib/cpp/libcpp/functional.pxd | 7 --- python/cudf/cudf/_lib/cpp/libcpp/memory.pxd | 12 ----- python/cudf/cudf/_lib/cpp/libcpp/optional.pxd | 50 ------------------- python/cudf/cudf/_lib/expressions.pyx | 3 +- python/cudf/cudf/_lib/groupby.pyx | 3 +- python/cudf/cudf/_lib/join.pyx | 3 +- python/cudf/cudf/_lib/null_mask.pyx | 3 +- python/cudf/cudf/_lib/parquet.pyx | 3 +- python/cudf/cudf/_lib/timezone.pyx | 2 +- python/cudf/pyproject.toml | 2 +- python/cudf_kafka/cudf_kafka/_lib/kafka.pyx | 3 +- python/cudf_kafka/pyproject.toml | 2 +- 27 files changed, 27 insertions(+), 101 deletions(-) delete mode 100644 python/cudf/cudf/_lib/cpp/libcpp/__init__.pxd delete mode 100644 python/cudf/cudf/_lib/cpp/libcpp/__init__.py delete mode 100644 python/cudf/cudf/_lib/cpp/libcpp/functional.pxd delete mode 100644 python/cudf/cudf/_lib/cpp/libcpp/memory.pxd delete mode 100644 python/cudf/cudf/_lib/cpp/libcpp/optional.pxd diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index a479d517c24..adf4fcad32d 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -24,7 +24,7 @@ dependencies: - cudatoolkit - cupy>=12.0.0 - cxx-compiler -- cython>=3.0.0 +- cython>=3.0.3 - dask-core>=2023.9.2 - dask-cuda==23.12.* - dask>=2023.9.2 diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index d1779aaeeac..a69ef587570 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -26,7 +26,7 @@ dependencies: - cuda-version=12.0 - cupy>=12.0.0 - cxx-compiler -- cython>=3.0.0 +- cython>=3.0.3 - dask-core>=2023.9.2 - dask-cuda==23.12.* - dask>=2023.9.2 diff --git a/conda/recipes/cudf/meta.yaml b/conda/recipes/cudf/meta.yaml index 9b5c5f3d14b..27edde1c98a 100644 --- a/conda/recipes/cudf/meta.yaml +++ b/conda/recipes/cudf/meta.yaml @@ -57,7 +57,7 @@ requirements: host: - protobuf ==4.24.* - python - - cython >=3.0.0 + - cython >=3.0.3 - scikit-build >=0.13.1 - setuptools - dlpack >=0.5,<0.6.0a0 diff --git a/conda/recipes/cudf_kafka/meta.yaml b/conda/recipes/cudf_kafka/meta.yaml index cdc547b4d68..9440f8bf124 100644 --- a/conda/recipes/cudf_kafka/meta.yaml +++ b/conda/recipes/cudf_kafka/meta.yaml @@ -51,7 +51,7 @@ requirements: {% endif %} host: - python - - cython >=3.0.0 + - cython >=3.0.3 - cuda-version ={{ cuda_version }} - cudf ={{ version }} - libcudf_kafka ={{ version }} diff --git a/dependencies.yaml b/dependencies.yaml index 3850347aa63..a16b51f4483 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -256,7 +256,7 @@ dependencies: common: - output_types: [conda, requirements, pyproject] packages: - - cython>=3.0.0 + - cython>=3.0.3 # TODO: Pin to numpy<1.25 until cudf requires pandas 2 - &numpy numpy>=1.21,<1.25 - output_types: [conda, requirements, pyproject] diff --git a/python/cudf/cudf/_lib/column.pyx b/python/cudf/cudf/_lib/column.pyx index f751d73b142..0edf9f8aa95 100644 --- a/python/cudf/cudf/_lib/column.pyx +++ b/python/cudf/cudf/_lib/column.pyx @@ -24,7 +24,7 @@ from cudf.utils.dtypes import _get_base_dtype from cpython.buffer cimport PyObject_CheckBuffer from libc.stdint cimport uintptr_t -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.utility cimport move from libcpp.vector cimport vector @@ -47,7 +47,6 @@ from cudf._lib.cpp.column.column_factories cimport ( make_numeric_column, ) from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.null_mask cimport null_count as cpp_null_count from cudf._lib.cpp.scalar.scalar cimport scalar from cudf._lib.scalar cimport DeviceScalar diff --git a/python/cudf/cudf/_lib/concat.pyx b/python/cudf/cudf/_lib/concat.pyx index feaf75ef237..1ec4719631e 100644 --- a/python/cudf/cudf/_lib/concat.pyx +++ b/python/cudf/cudf/_lib/concat.pyx @@ -1,7 +1,7 @@ # Copyright (c) 2020-2023, NVIDIA CORPORATION. from libcpp cimport bool -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.utility cimport move from libcpp.vector cimport vector @@ -12,7 +12,6 @@ from cudf._lib.cpp.concatenate cimport ( concatenate_masks as libcudf_concatenate_masks, concatenate_tables as libcudf_concatenate_tables, ) -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.table.table cimport table, table_view from cudf._lib.utils cimport ( data_from_unique_ptr, diff --git a/python/cudf/cudf/_lib/copying.pyx b/python/cudf/cudf/_lib/copying.pyx index f57bc15ed57..ea6ee76c14a 100644 --- a/python/cudf/cudf/_lib/copying.pyx +++ b/python/cudf/cudf/_lib/copying.pyx @@ -24,12 +24,13 @@ from cudf._lib.utils cimport table_view_from_columns, table_view_from_table from cudf._lib.reduce import minmax from cudf.core.abc import Serializable +from libcpp.functional cimport reference_wrapper +from libcpp.memory cimport make_unique + cimport cudf._lib.cpp.contiguous_split as cpp_contiguous_split cimport cudf._lib.cpp.copying as cpp_copying from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport column_view, mutable_column_view -from cudf._lib.cpp.libcpp.functional cimport reference_wrapper -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.lists.gather cimport ( segmented_gather as cpp_segmented_gather, ) diff --git a/python/cudf/cudf/_lib/cpp/copying.pxd b/python/cudf/cudf/_lib/cpp/copying.pxd index 20725c252fc..5637b55ac1c 100644 --- a/python/cudf/cudf/_lib/cpp/copying.pxd +++ b/python/cudf/cudf/_lib/cpp/copying.pxd @@ -2,6 +2,7 @@ from libc.stdint cimport int32_t, int64_t, uint8_t from libcpp cimport bool +from libcpp.functional cimport reference_wrapper from libcpp.memory cimport unique_ptr from libcpp.vector cimport vector @@ -9,7 +10,6 @@ from rmm._lib.device_buffer cimport device_buffer from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport column_view, mutable_column_view -from cudf._lib.cpp.libcpp.functional cimport reference_wrapper from cudf._lib.cpp.scalar.scalar cimport scalar from cudf._lib.cpp.table.table cimport table from cudf._lib.cpp.table.table_view cimport table_view diff --git a/python/cudf/cudf/_lib/cpp/groupby.pxd b/python/cudf/cudf/_lib/cpp/groupby.pxd index 2ecdf76842f..0266404fc50 100644 --- a/python/cudf/cudf/_lib/cpp/groupby.pxd +++ b/python/cudf/cudf/_lib/cpp/groupby.pxd @@ -1,6 +1,7 @@ -# Copyright (c) 2020-2021, NVIDIA CORPORATION. +# Copyright (c) 2020-2023, NVIDIA CORPORATION. from libcpp cimport bool +from libcpp.functional cimport reference_wrapper from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair from libcpp.vector cimport vector @@ -11,7 +12,6 @@ from cudf._lib.cpp.aggregation cimport ( ) from cudf._lib.cpp.column.column cimport column from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.libcpp.functional cimport reference_wrapper from cudf._lib.cpp.replace cimport replace_policy from cudf._lib.cpp.scalar.scalar cimport scalar from cudf._lib.cpp.table.table cimport table diff --git a/python/cudf/cudf/_lib/cpp/io/orc.pxd b/python/cudf/cudf/_lib/cpp/io/orc.pxd index dd6f919a74d..d5ac8574fe4 100644 --- a/python/cudf/cudf/_lib/cpp/io/orc.pxd +++ b/python/cudf/cudf/_lib/cpp/io/orc.pxd @@ -4,12 +4,12 @@ from libc.stdint cimport uint8_t from libcpp cimport bool from libcpp.map cimport map from libcpp.memory cimport shared_ptr, unique_ptr +from libcpp.optional cimport optional from libcpp.string cimport string from libcpp.vector cimport vector cimport cudf._lib.cpp.io.types as cudf_io_types cimport cudf._lib.cpp.table.table_view as cudf_table_view -from cudf._lib.cpp.libcpp.optional cimport optional from cudf._lib.cpp.types cimport data_type, size_type diff --git a/python/cudf/cudf/_lib/cpp/io/parquet.pxd b/python/cudf/cudf/_lib/cpp/io/parquet.pxd index a6a7ba034aa..cdd1bde0274 100644 --- a/python/cudf/cudf/_lib/cpp/io/parquet.pxd +++ b/python/cudf/cudf/_lib/cpp/io/parquet.pxd @@ -2,16 +2,16 @@ from libc.stdint cimport uint8_t from libcpp cimport bool +from libcpp.functional cimport reference_wrapper from libcpp.map cimport map from libcpp.memory cimport shared_ptr, unique_ptr +from libcpp.optional cimport optional from libcpp.string cimport string from libcpp.vector cimport vector cimport cudf._lib.cpp.io.types as cudf_io_types cimport cudf._lib.cpp.table.table_view as cudf_table_view from cudf._lib.cpp.expressions cimport expression -from cudf._lib.cpp.libcpp.functional cimport reference_wrapper -from cudf._lib.cpp.libcpp.optional cimport optional from cudf._lib.cpp.types cimport data_type, size_type diff --git a/python/cudf/cudf/_lib/cpp/io/timezone.pxd b/python/cudf/cudf/_lib/cpp/io/timezone.pxd index ba481d9a1d3..927c2118473 100644 --- a/python/cudf/cudf/_lib/cpp/io/timezone.pxd +++ b/python/cudf/cudf/_lib/cpp/io/timezone.pxd @@ -2,9 +2,9 @@ from libcpp cimport bool from libcpp.memory cimport unique_ptr +from libcpp.optional cimport optional from libcpp.string cimport string -from cudf._lib.cpp.libcpp.optional cimport optional from cudf._lib.cpp.table.table cimport table diff --git a/python/cudf/cudf/_lib/cpp/libcpp/__init__.pxd b/python/cudf/cudf/_lib/cpp/libcpp/__init__.pxd deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/python/cudf/cudf/_lib/cpp/libcpp/__init__.py b/python/cudf/cudf/_lib/cpp/libcpp/__init__.py deleted file mode 100644 index e69de29bb2d..00000000000 diff --git a/python/cudf/cudf/_lib/cpp/libcpp/functional.pxd b/python/cudf/cudf/_lib/cpp/libcpp/functional.pxd deleted file mode 100644 index f3e2d6d0878..00000000000 --- a/python/cudf/cudf/_lib/cpp/libcpp/functional.pxd +++ /dev/null @@ -1,7 +0,0 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. - - -cdef extern from "" namespace "std" nogil: - cdef cppclass reference_wrapper[T]: - reference_wrapper() - reference_wrapper(T) diff --git a/python/cudf/cudf/_lib/cpp/libcpp/memory.pxd b/python/cudf/cudf/_lib/cpp/libcpp/memory.pxd deleted file mode 100644 index 2178f1a940c..00000000000 --- a/python/cudf/cudf/_lib/cpp/libcpp/memory.pxd +++ /dev/null @@ -1,12 +0,0 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. - -from libcpp.memory cimport unique_ptr - - -cdef extern from "" namespace "std" nogil: - # The Cython standard header does not have except +, so C++ - # exceptions from make_unique are not caught and translated to - # Python ones. This is not perfectly ergonomic, we always have to - # wrap make_unique in move, but at least we can catch exceptions. - # See https://github.com/cython/cython/issues/5560 - unique_ptr[T] make_unique[T](...) except + diff --git a/python/cudf/cudf/_lib/cpp/libcpp/optional.pxd b/python/cudf/cudf/_lib/cpp/libcpp/optional.pxd deleted file mode 100644 index a78c18f3f7a..00000000000 --- a/python/cudf/cudf/_lib/cpp/libcpp/optional.pxd +++ /dev/null @@ -1,50 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2023, NVIDIA CORPORATION & -# AFFILIATES. All rights reserved. SPDX-License-Identifier: -# Apache-2.0 -# -# 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. - -from libcpp cimport bool - - -cdef extern from "" namespace "std" nogil: - cdef cppclass nullopt_t: - nullopt_t() - - cdef nullopt_t nullopt - - cdef cppclass optional[T]: - ctypedef T value_type - optional() - optional(nullopt_t) - optional(optional&) except + - optional(T&) except + - bool has_value() - T& value() - T& value_or[U](U& default_value) - void swap(optional&) - void reset() - T& emplace(...) - T& operator*() - optional& operator=(optional&) - optional& operator=[U](U&) - bool operator bool() - bool operator!() - bool operator==[U](optional&, U&) - bool operator!=[U](optional&, U&) - bool operator<[U](optional&, U&) - bool operator>[U](optional&, U&) - bool operator<=[U](optional&, U&) - bool operator>=[U](optional&, U&) - - optional[T] make_optional[T](...) except + diff --git a/python/cudf/cudf/_lib/expressions.pyx b/python/cudf/cudf/_lib/expressions.pyx index 8d7545ffe15..01a080f635f 100644 --- a/python/cudf/cudf/_lib/expressions.pyx +++ b/python/cudf/cudf/_lib/expressions.pyx @@ -4,12 +4,11 @@ from enum import Enum from cython.operator cimport dereference from libc.stdint cimport int64_t -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.string cimport string from libcpp.utility cimport move from cudf._lib.cpp cimport expressions as libcudf_exp -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.types cimport size_type # Necessary for proper casting, see below. diff --git a/python/cudf/cudf/_lib/groupby.pyx b/python/cudf/cudf/_lib/groupby.pyx index a26d820de6f..b3778e45cde 100644 --- a/python/cudf/cudf/_lib/groupby.pyx +++ b/python/cudf/cudf/_lib/groupby.pyx @@ -24,6 +24,8 @@ from cudf._lib.utils cimport columns_from_unique_ptr, table_view_from_columns from cudf._lib.scalar import as_device_scalar +from libcpp.functional cimport reference_wrapper + cimport cudf._lib.cpp.groupby as libcudf_groupby cimport cudf._lib.cpp.types as libcudf_types from cudf._lib.aggregation cimport ( @@ -33,7 +35,6 @@ from cudf._lib.aggregation cimport ( make_groupby_scan_aggregation, ) from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.libcpp.functional cimport reference_wrapper from cudf._lib.cpp.replace cimport replace_policy from cudf._lib.cpp.scalar.scalar cimport scalar from cudf._lib.cpp.table.table cimport table, table_view diff --git a/python/cudf/cudf/_lib/join.pyx b/python/cudf/cudf/_lib/join.pyx index 416680aae24..378be978cc0 100644 --- a/python/cudf/cudf/_lib/join.pyx +++ b/python/cudf/cudf/_lib/join.pyx @@ -2,7 +2,7 @@ from cudf.core.buffer import acquire_spill_lock -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.pair cimport pair from libcpp.utility cimport move @@ -11,7 +11,6 @@ from rmm._lib.device_buffer cimport device_buffer cimport cudf._lib.cpp.join as cpp_join from cudf._lib.column cimport Column from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.table.table_view cimport table_view from cudf._lib.cpp.types cimport data_type, size_type, type_id from cudf._lib.utils cimport table_view_from_columns diff --git a/python/cudf/cudf/_lib/null_mask.pyx b/python/cudf/cudf/_lib/null_mask.pyx index 5b4538629f6..1f98140d9e4 100644 --- a/python/cudf/cudf/_lib/null_mask.pyx +++ b/python/cudf/cudf/_lib/null_mask.pyx @@ -6,13 +6,12 @@ from rmm._lib.device_buffer cimport DeviceBuffer, device_buffer from cudf.core.buffer import acquire_spill_lock, as_buffer -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.pair cimport pair from libcpp.utility cimport move from cudf._lib.column cimport Column from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.null_mask cimport ( bitmask_allocation_size_bytes as cpp_bitmask_allocation_size_bytes, bitmask_and as cpp_bitmask_and, diff --git a/python/cudf/cudf/_lib/parquet.pyx b/python/cudf/cudf/_lib/parquet.pyx index d8d363686cc..4acb1ce10b1 100644 --- a/python/cudf/cudf/_lib/parquet.pyx +++ b/python/cudf/cudf/_lib/parquet.pyx @@ -32,7 +32,7 @@ from cudf._lib.utils import _index_level_name, generate_pandas_metadata from libc.stdint cimport uint8_t from libcpp cimport bool from libcpp.map cimport map -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.string cimport string from libcpp.unordered_map cimport unordered_map from libcpp.utility cimport move @@ -52,7 +52,6 @@ from cudf._lib.cpp.io.parquet cimport ( write_parquet as parquet_writer, ) from cudf._lib.cpp.io.types cimport column_in_metadata, table_input_metadata -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf._lib.cpp.table.table_view cimport table_view from cudf._lib.cpp.types cimport data_type, size_type from cudf._lib.io.datasource cimport NativeFileDatasource diff --git a/python/cudf/cudf/_lib/timezone.pyx b/python/cudf/cudf/_lib/timezone.pyx index 4d76cbfcdb5..808d1321b0b 100644 --- a/python/cudf/cudf/_lib/timezone.pyx +++ b/python/cudf/cudf/_lib/timezone.pyx @@ -1,13 +1,13 @@ # Copyright (c) 2023, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr +from libcpp.optional cimport make_optional from libcpp.string cimport string from libcpp.utility cimport move from cudf._lib.cpp.io.timezone cimport ( make_timezone_transition_table as cpp_make_timezone_transition_table, ) -from cudf._lib.cpp.libcpp.optional cimport make_optional from cudf._lib.cpp.table.table cimport table from cudf._lib.utils cimport columns_from_unique_ptr diff --git a/python/cudf/pyproject.toml b/python/cudf/pyproject.toml index 1c687269e55..b38970271d7 100644 --- a/python/cudf/pyproject.toml +++ b/python/cudf/pyproject.toml @@ -4,7 +4,7 @@ build-backend = "setuptools.build_meta" requires = [ "cmake>=3.26.4", - "cython>=3.0.0", + "cython>=3.0.3", "ninja", "numpy>=1.21,<1.25", "protoc-wheel", diff --git a/python/cudf_kafka/cudf_kafka/_lib/kafka.pyx b/python/cudf_kafka/cudf_kafka/_lib/kafka.pyx index 4d732478723..2fbaacff7c6 100644 --- a/python/cudf_kafka/cudf_kafka/_lib/kafka.pyx +++ b/python/cudf_kafka/cudf_kafka/_lib/kafka.pyx @@ -3,12 +3,11 @@ from libc.stdint cimport int32_t, int64_t from libcpp cimport bool, nullptr from libcpp.map cimport map -from libcpp.memory cimport unique_ptr +from libcpp.memory cimport make_unique, unique_ptr from libcpp.string cimport string from libcpp.utility cimport move from cudf._lib.cpp.io.datasource cimport datasource -from cudf._lib.cpp.libcpp.memory cimport make_unique from cudf_kafka._lib.kafka cimport kafka_consumer diff --git a/python/cudf_kafka/pyproject.toml b/python/cudf_kafka/pyproject.toml index f5cbd480e9c..4829f06ab09 100644 --- a/python/cudf_kafka/pyproject.toml +++ b/python/cudf_kafka/pyproject.toml @@ -3,7 +3,7 @@ [build-system] requires = [ - "cython>=3.0.0", + "cython>=3.0.3", "numpy>=1.21,<1.25", "pyarrow==14.0.1.*", "setuptools", From 04d13d81b0bb4c2b3db2bfc9d9e28432e0a73c44 Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Mon, 13 Nov 2023 09:05:19 -0500 Subject: [PATCH 03/44] Normalizing offsets iterator (#14234) Creates a normalizing offsets iterator that returns an int64 value given either a int32 or int64 column data. Depends on #14206 Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Divye Gala (https://github.com/divyegala) - Yunsong Wang (https://github.com/PointKernel) URL: https://github.com/rapidsai/cudf/pull/14234 --- .../cudf/column/column_device_view.cuh | 8 +- cpp/include/cudf/detail/indexalator.cuh | 151 ++++++++++++++-- .../cudf/detail/normalizing_iterator.cuh | 160 +---------------- cpp/include/cudf/detail/offsets_iterator.cuh | 165 ++++++++++++++++++ .../cudf/detail/offsets_iterator_factory.cuh | 47 +++++ cpp/tests/CMakeLists.txt | 1 + cpp/tests/iterator/indexalator_test.cu | 37 ---- cpp/tests/iterator/offsetalator_test.cu | 140 +++++++++++++++ 8 files changed, 502 insertions(+), 207 deletions(-) create mode 100644 cpp/include/cudf/detail/offsets_iterator.cuh create mode 100644 cpp/include/cudf/detail/offsets_iterator_factory.cuh create mode 100644 cpp/tests/iterator/offsetalator_test.cu diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index 35851a99822..b1ff0bbaea7 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -16,6 +16,7 @@ #pragma once #include +#include #include #include #include @@ -442,10 +443,11 @@ class alignas(16) column_device_view : public detail::column_device_view_base { __device__ T element(size_type element_index) const noexcept { size_type index = element_index + offset(); // account for this view's _offset - auto const* d_offsets = d_children[strings_column_view::offsets_column_index].data(); char const* d_strings = d_children[strings_column_view::chars_column_index].data(); - size_type offset = d_offsets[index]; - return string_view{d_strings + offset, d_offsets[index + 1] - offset}; + auto const offsets = d_children[strings_column_view::offsets_column_index]; + auto const itr = cudf::detail::input_offsetalator(offsets.head(), offsets.type()); + auto const offset = itr[index]; + return string_view{d_strings + offset, static_cast(itr[index + 1] - offset)}; } private: diff --git a/cpp/include/cudf/detail/indexalator.cuh b/cpp/include/cudf/detail/indexalator.cuh index 6532dae3695..4d261c54b29 100644 --- a/cpp/include/cudf/detail/indexalator.cuh +++ b/cpp/include/cudf/detail/indexalator.cuh @@ -56,10 +56,69 @@ namespace detail { * auto result = thrust::find(thrust::device, begin, end, size_type{12} ); * @endcode */ -using input_indexalator = input_normalator; +struct input_indexalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = cudf::size_type const; // this keeps STL and thrust happy + + input_indexalator() = default; + input_indexalator(input_indexalator const&) = default; + input_indexalator(input_indexalator&&) = default; + input_indexalator& operator=(input_indexalator const&) = default; + input_indexalator& operator=(input_indexalator&&) = default; + + /** + * @brief Indirection operator returns the value at the current iterator position + */ + __device__ inline cudf::size_type operator*() const { return operator[](0); } + + /** + * @brief Dispatch functor for resolving a Integer value from any integer type + */ + struct normalize_type { + template ())> + __device__ cudf::size_type operator()(void const* tp) + { + return static_cast(*static_cast(tp)); + } + template ())> + __device__ cudf::size_type operator()(void const*) + { + CUDF_UNREACHABLE("only integral types are supported"); + } + }; + + /** + * @brief Array subscript operator returns a value at the input + * `idx` position as a `Integer` value. + */ + __device__ inline cudf::size_type operator[](size_type idx) const + { + void const* tp = p_ + (idx * this->width_); + return type_dispatcher(this->dtype_, normalize_type{}, tp); + } + + /** + * @brief Create an input index normalizing iterator + * + * Use the indexalator_factory to create an iterator instance. + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + * @param offset Applied to the data pointer per size of the type + */ + CUDF_HOST_DEVICE input_indexalator(void const* data, data_type dtype, cudf::size_type offset = 0) + : base_normalator(dtype), p_{static_cast(data)} + { + p_ += offset * this->width_; + } + + protected: + char const* p_; /// pointer to the integer data in device memory +}; /** - * @brief The index normalizing output iterator. + * @brief The index normalizing output iterator * * This is an iterator that can be used for index types (integers) without * requiring a type-specific instance. It can be used for any iterator @@ -82,7 +141,75 @@ using input_indexalator = input_normalator; * thrust::less()); * @endcode */ -using output_indexalator = output_normalator; +struct output_indexalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = output_indexalator const&; // required for output iterators + + output_indexalator() = default; + output_indexalator(output_indexalator const&) = default; + output_indexalator(output_indexalator&&) = default; + output_indexalator& operator=(output_indexalator const&) = default; + output_indexalator& operator=(output_indexalator&&) = default; + + /** + * @brief Indirection operator returns this iterator instance in order + * to capture the `operator=(Integer)` calls. + */ + __device__ inline reference operator*() const { return *this; } + + /** + * @brief Array subscript operator returns an iterator instance at the specified `idx` position. + * + * This allows capturing the subsequent `operator=(Integer)` call in this class. + */ + __device__ inline output_indexalator const operator[](size_type idx) const + { + output_indexalator tmp{*this}; + tmp.p_ += (idx * this->width_); + return tmp; + } + + /** + * @brief Dispatch functor for setting the index value from a size_type value. + */ + struct normalize_type { + template ())> + __device__ void operator()(void* tp, cudf::size_type const value) + { + (*static_cast(tp)) = static_cast(value); + } + template ())> + __device__ void operator()(void*, cudf::size_type const) + { + CUDF_UNREACHABLE("only index types are supported"); + } + }; + + /** + * @brief Assign an Integer value to the current iterator position + */ + __device__ inline reference operator=(cudf::size_type const value) const + { + void* tp = p_; + type_dispatcher(this->dtype_, normalize_type{}, tp, value); + return *this; + } + + /** + * @brief Create an output normalizing iterator + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE output_indexalator(void* data, data_type dtype) + : base_normalator(dtype), p_{static_cast(data)} + { + } + + protected: + char* p_; /// pointer to the integer data in device memory +}; /** * @brief Use this class to create an indexalator instance. @@ -92,14 +219,12 @@ struct indexalator_factory { * @brief A type_dispatcher functor to create an input iterator from an indices column. */ struct input_indexalator_fn { - template ()>* = nullptr> + template ())> input_indexalator operator()(column_view const& indices) { return input_indexalator(indices.data(), indices.type()); } - template ()>* = nullptr> + template ())> input_indexalator operator()(Args&&... args) { CUDF_FAIL("indices must be an index type"); @@ -110,16 +235,14 @@ struct indexalator_factory { * @brief Use this class to create an indexalator to a scalar index. */ struct input_indexalator_scalar_fn { - template ()>* = nullptr> + template ())> input_indexalator operator()(scalar const& index) { // note: using static_cast const&>(index) creates a copy auto const scalar_impl = static_cast const*>(&index); return input_indexalator(scalar_impl->data(), index.type()); } - template ()>* = nullptr> + template ())> input_indexalator operator()(Args&&... args) { CUDF_FAIL("scalar must be an index type"); @@ -130,14 +253,12 @@ struct indexalator_factory { * @brief A type_dispatcher functor to create an output iterator from an indices column. */ struct output_indexalator_fn { - template ()>* = nullptr> + template ())> output_indexalator operator()(mutable_column_view const& indices) { return output_indexalator(indices.data(), indices.type()); } - template ()>* = nullptr> + template ())> output_indexalator operator()(Args&&... args) { CUDF_FAIL("indices must be an index type"); diff --git a/cpp/include/cudf/detail/normalizing_iterator.cuh b/cpp/include/cudf/detail/normalizing_iterator.cuh index 35a695d47df..8f90afc3e57 100644 --- a/cpp/include/cudf/detail/normalizing_iterator.cuh +++ b/cpp/include/cudf/detail/normalizing_iterator.cuh @@ -33,7 +33,7 @@ namespace detail { * @tparam Integer The type the iterator normalizes to */ template -struct base_normalator { +struct alignas(16) base_normalator { static_assert(cudf::is_index_type()); using difference_type = std::ptrdiff_t; using value_type = Integer; @@ -204,7 +204,7 @@ struct base_normalator { private: struct integer_sizeof_fn { - template ()>* = nullptr> + template ())> CUDF_HOST_DEVICE constexpr std::size_t operator()() const { #ifndef __CUDA_ARCH__ @@ -213,7 +213,7 @@ struct base_normalator { CUDF_UNREACHABLE("only integral types are supported"); #endif } - template ()>* = nullptr> + template ())> CUDF_HOST_DEVICE constexpr std::size_t operator()() const noexcept { return sizeof(T); @@ -229,160 +229,16 @@ struct base_normalator { width_ = static_cast(type_dispatcher(dtype, integer_sizeof_fn{})); } - int32_t width_; /// integer type width = 1,2,4, or 8 - data_type dtype_; /// for type-dispatcher calls -}; - -/** - * @brief The integer normalizing input iterator - * - * This is an iterator that can be used for index types (integers) without - * requiring a type-specific instance. It can be used for any iterator - * interface for reading an array of integer values of type - * int8, int16, int32, int64, uint8, uint16, uint32, or uint64. - * Reading specific elements always return a type of `Integer` - * - * @tparam Integer Type returned by all read functions - */ -template -struct input_normalator : base_normalator, Integer> { - friend struct base_normalator, Integer>; // for CRTP - - using reference = Integer const; // this keeps STL and thrust happy - - input_normalator() = default; - input_normalator(input_normalator const&) = default; - input_normalator(input_normalator&&) = default; - input_normalator& operator=(input_normalator const&) = default; - input_normalator& operator=(input_normalator&&) = default; - - /** - * @brief Indirection operator returns the value at the current iterator position - */ - __device__ inline Integer operator*() const { return operator[](0); } - - /** - * @brief Dispatch functor for resolving a Integer value from any integer type - */ - struct normalize_type { - template ()>* = nullptr> - __device__ Integer operator()(void const* tp) - { - return static_cast(*static_cast(tp)); - } - template ()>* = nullptr> - __device__ Integer operator()(void const*) - { - CUDF_UNREACHABLE("only integral types are supported"); - } - }; - /** - * @brief Array subscript operator returns a value at the input - * `idx` position as a `Integer` value. - */ - __device__ inline Integer operator[](size_type idx) const - { - void const* tp = p_ + (idx * this->width_); - return type_dispatcher(this->dtype_, normalize_type{}, tp); - } - - /** - * @brief Create an input index normalizing iterator. - * - * Use the indexalator_factory to create an iterator instance. - * - * @param data Pointer to an integer array in device memory. - * @param data_type Type of data in data - */ - CUDF_HOST_DEVICE input_normalator(void const* data, data_type dtype, cudf::size_type offset = 0) - : base_normalator, Integer>(dtype), p_{static_cast(data)} - { - p_ += offset * this->width_; - } - - char const* p_; /// pointer to the integer data in device memory -}; - -/** - * @brief The integer normalizing output iterator - * - * This is an iterator that can be used for index types (integers) without - * requiring a type-specific instance. It can be used for any iterator - * interface for writing an array of integer values of type - * int8, int16, int32, int64, uint8, uint16, uint32, or uint64. - * Setting specific elements always accept the `Integer` type values. - * - * @tparam Integer The type used for all write functions - */ -template -struct output_normalator : base_normalator, Integer> { - friend struct base_normalator, Integer>; // for CRTP - - using reference = output_normalator const&; // required for output iterators - - output_normalator() = default; - output_normalator(output_normalator const&) = default; - output_normalator(output_normalator&&) = default; - output_normalator& operator=(output_normalator const&) = default; - output_normalator& operator=(output_normalator&&) = default; - - /** - * @brief Indirection operator returns this iterator instance in order - * to capture the `operator=(Integer)` calls. - */ - __device__ inline output_normalator const& operator*() const { return *this; } - - /** - * @brief Array subscript operator returns an iterator instance at the specified `idx` position. - * - * This allows capturing the subsequent `operator=(Integer)` call in this class. - */ - __device__ inline output_normalator const operator[](size_type idx) const - { - output_normalator tmp{*this}; - tmp.p_ += (idx * this->width_); - return tmp; - } - - /** - * @brief Dispatch functor for setting the index value from a size_type value. - */ - struct normalize_type { - template ()>* = nullptr> - __device__ void operator()(void* tp, Integer const value) - { - (*static_cast(tp)) = static_cast(value); - } - template ()>* = nullptr> - __device__ void operator()(void*, Integer const) - { - CUDF_UNREACHABLE("only index types are supported"); - } - }; - - /** - * @brief Assign an Integer value to the current iterator position - */ - __device__ inline output_normalator const& operator=(Integer const value) const - { - void* tp = p_; - type_dispatcher(this->dtype_, normalize_type{}, tp, value); - return *this; - } - - /** - * @brief Create an output normalizing iterator - * - * @param data Pointer to an integer array in device memory. - * @param data_type Type of data in data + * @brief Constructor assigns width and type member variables for base class. */ - CUDF_HOST_DEVICE output_normalator(void* data, data_type dtype) - : base_normalator, Integer>(dtype), p_{static_cast(data)} + explicit CUDF_HOST_DEVICE base_normalator(data_type dtype, int32_t width) + : width_(width), dtype_(dtype) { } - char* p_; /// pointer to the integer data in device memory + int32_t width_; /// integer type width = 1,2,4, or 8 + data_type dtype_; /// for type-dispatcher calls }; } // namespace detail diff --git a/cpp/include/cudf/detail/offsets_iterator.cuh b/cpp/include/cudf/detail/offsets_iterator.cuh new file mode 100644 index 00000000000..3eb77b32353 --- /dev/null +++ b/cpp/include/cudf/detail/offsets_iterator.cuh @@ -0,0 +1,165 @@ +/* + * Copyright (c) 2023, NVIDIA 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. + */ + +#pragma once + +#include +#include + +namespace cudf { +namespace detail { + +/** + * @brief The offsets normalizing input iterator + * + * This is an iterator that can be used for offsets where the underlying + * type may be int32_t or int64_t. + * + * Use the offsetalator_factory to create an appropriate input iterator + * from an offsets column_view. + */ +struct input_offsetalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = int64_t const; // this keeps STL and thrust happy + + input_offsetalator() = default; + input_offsetalator(input_offsetalator const&) = default; + input_offsetalator(input_offsetalator&&) = default; + input_offsetalator& operator=(input_offsetalator const&) = default; + input_offsetalator& operator=(input_offsetalator&&) = default; + + /** + * @brief Indirection operator returns the value at the current iterator position + */ + __device__ inline int64_t operator*() const { return operator[](0); } + + /** + * @brief Array subscript operator returns a value at the input + * `idx` position as a int64_t value. + */ + __device__ inline int64_t operator[](size_type idx) const + { + void const* tp = p_ + (idx * this->width_); + return this->width_ == sizeof(int32_t) ? static_cast(*static_cast(tp)) + : *static_cast(tp); + } + + /** + * @brief Create an input index normalizing iterator. + * + * Use the indexalator_factory to create an iterator instance. + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE input_offsetalator(void const* data, data_type dtype) + : base_normalator( + dtype, dtype.id() == type_id::INT32 ? sizeof(int32_t) : sizeof(int64_t)), + p_{static_cast(data)} + { +#ifndef __CUDA_ARCH__ + CUDF_EXPECTS(dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64, + "Unexpected offsets type"); +#else + cudf_assert((dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64) && + "Unexpected offsets type"); +#endif + } + + protected: + char const* p_; /// pointer to the integer data in device memory +}; + +/** + * @brief The offsets normalizing output iterator + * + * This is an iterator that can be used for storing offsets values + * where the underlying type may be either int32_t or int64_t. + * + * Use the offsetalator_factory to create an appropriate output iterator + * from a mutable_column_view. + * + */ +struct output_offsetalator : base_normalator { + friend struct base_normalator; // for CRTP + + using reference = output_offsetalator const&; // required for output iterators + + output_offsetalator() = default; + output_offsetalator(output_offsetalator const&) = default; + output_offsetalator(output_offsetalator&&) = default; + output_offsetalator& operator=(output_offsetalator const&) = default; + output_offsetalator& operator=(output_offsetalator&&) = default; + + /** + * @brief Indirection operator returns this iterator instance in order + * to capture the `operator=(int64)` calls. + */ + __device__ inline output_offsetalator const& operator*() const { return *this; } + + /** + * @brief Array subscript operator returns an iterator instance at the specified `idx` position. + * + * This allows capturing the subsequent `operator=(int64)` call in this class. + */ + __device__ inline output_offsetalator const operator[](size_type idx) const + { + output_offsetalator tmp{*this}; + tmp.p_ += (idx * this->width_); + return tmp; + } + + /** + * @brief Assign an offset value to the current iterator position + */ + __device__ inline output_offsetalator const& operator=(int64_t const value) const + { + void* tp = p_; + if (this->width_ == sizeof(int32_t)) { + (*static_cast(tp)) = static_cast(value); + } else { + (*static_cast(tp)) = value; + } + return *this; + } + + /** + * @brief Create an output offsets iterator + * + * @param data Pointer to an integer array in device memory. + * @param dtype Type of data in data + */ + CUDF_HOST_DEVICE output_offsetalator(void* data, data_type dtype) + : base_normalator( + dtype, dtype.id() == type_id::INT32 ? sizeof(int32_t) : sizeof(int64_t)), + p_{static_cast(data)} + { +#ifndef __CUDA_ARCH__ + CUDF_EXPECTS(dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64, + "Unexpected offsets type"); +#else + cudf_assert((dtype.id() == type_id::INT32 || dtype.id() == type_id::INT64) && + "Unexpected offsets type"); +#endif + } + + protected: + char* p_; /// pointer to the integer data in device memory +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/detail/offsets_iterator_factory.cuh b/cpp/include/cudf/detail/offsets_iterator_factory.cuh new file mode 100644 index 00000000000..5b4c6b825d2 --- /dev/null +++ b/cpp/include/cudf/detail/offsets_iterator_factory.cuh @@ -0,0 +1,47 @@ +/* + * Copyright (c) 2023, NVIDIA 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. + */ + +#pragma once + +#include +#include + +namespace cudf { +namespace detail { + +/** + * @brief Use this class to create an offsetalator instance. + */ +struct offsetalator_factory { + /** + * @brief Create an input offsetalator instance from an offsets column + */ + static input_offsetalator make_input_iterator(column_view const& offsets) + { + return input_offsetalator(offsets.head(), offsets.type()); + } + + /** + * @brief Create an output offsetalator instance from an offsets column + */ + static output_offsetalator make_output_iterator(mutable_column_view const& offsets) + { + return output_offsetalator(offsets.head(), offsets.type()); + } +}; + +} // namespace detail +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b0382d15807..7b628649051 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -393,6 +393,7 @@ set_tests_properties( ConfigureTest( ITERATOR_TEST iterator/indexalator_test.cu + iterator/offsetalator_test.cu iterator/optional_iterator_test_chrono.cu iterator/optional_iterator_test_numeric.cu iterator/pair_iterator_test_chrono.cu diff --git a/cpp/tests/iterator/indexalator_test.cu b/cpp/tests/iterator/indexalator_test.cu index 3e8bcd5cb0d..0c10853ec02 100644 --- a/cpp/tests/iterator/indexalator_test.cu +++ b/cpp/tests/iterator/indexalator_test.cu @@ -157,40 +157,3 @@ TYPED_TEST(IndexalatorTest, output_iterator) expected = cudf::test::fixed_width_column_wrapper({0, 1, 1, 2, 3, 4, 5, 5, 7}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); } - -/** - * For testing creating and using the indexalator in device code. - */ -struct device_functor_fn { - cudf::column_device_view const d_col; - __device__ cudf::size_type operator()(cudf::size_type idx) - { - auto itr = cudf::detail::input_indexalator(d_col.head(), d_col.type()); - return itr[idx] * 3; - } -}; - -TYPED_TEST(IndexalatorTest, device_indexalator) -{ - using T = TypeParam; - - auto d_col1 = - cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); - auto d_col2 = - cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); - auto input = cudf::column_view(d_col1); - auto output = cudf::mutable_column_view(d_col2); - auto stream = cudf::get_default_stream(); - - auto d_input = cudf::column_device_view::create(input, stream); - - thrust::transform(rmm::exec_policy(stream), - thrust::counting_iterator(0), - thrust::counting_iterator(input.size()), - output.begin(), - device_functor_fn{*d_input}); - - auto expected = - cudf::test::fixed_width_column_wrapper({0, 18, 21, 42, 69, 99, 129, 135, 189}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); -} diff --git a/cpp/tests/iterator/offsetalator_test.cu b/cpp/tests/iterator/offsetalator_test.cu new file mode 100644 index 00000000000..e569e58f42a --- /dev/null +++ b/cpp/tests/iterator/offsetalator_test.cu @@ -0,0 +1,140 @@ +/* + * Copyright (c) 2023, NVIDIA 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 + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +using TestingTypes = cudf::test::Types; + +template +struct OffsetalatorTest : public IteratorTest {}; + +TYPED_TEST_SUITE(OffsetalatorTest, TestingTypes); + +TYPED_TEST(OffsetalatorTest, input_iterator) +{ + using T = TypeParam; + + auto host_values = cudf::test::make_type_param_vector({0, 6, 0, -14, 13, 64, -13, -20, 45}); + + auto d_col = cudf::test::fixed_width_column_wrapper(host_values.begin(), host_values.end()); + + auto expected_values = thrust::host_vector(host_values.size()); + std::transform(host_values.begin(), host_values.end(), expected_values.begin(), [](auto v) { + return static_cast(v); + }); + + auto it_dev = cudf::detail::offsetalator_factory::make_input_iterator(d_col); + this->iterator_test_thrust(expected_values, it_dev, host_values.size()); +} + +TYPED_TEST(OffsetalatorTest, output_iterator) +{ + using T = TypeParam; + + auto d_col1 = cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); + auto d_col2 = cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); + auto itr = cudf::detail::offsetalator_factory::make_output_iterator(d_col2); + auto input = cudf::column_view(d_col1); + auto stream = cudf::get_default_stream(); + + auto map = cudf::test::fixed_width_column_wrapper({0, 2, 4, 6, 8, 1, 3, 5, 7}); + auto d_map = cudf::column_view(map); + thrust::gather(rmm::exec_policy_nosync(stream), + d_map.begin(), + d_map.end(), + input.begin(), + itr); + auto expected = cudf::test::fixed_width_column_wrapper({0, 7, 23, 43, 63, 6, 14, 33, 45}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::scatter(rmm::exec_policy_nosync(stream), + input.begin(), + input.end(), + d_map.begin(), + itr); + expected = cudf::test::fixed_width_column_wrapper({0, 33, 6, 43, 7, 45, 14, 63, 23}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::fill(rmm::exec_policy(stream), itr, itr + input.size(), 77); + expected = cudf::test::fixed_width_column_wrapper({77, 77, 77, 77, 77, 77, 77, 77, 77}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + thrust::sequence(rmm::exec_policy(stream), itr, itr + input.size()); + expected = cudf::test::fixed_width_column_wrapper({0, 1, 2, 3, 4, 5, 6, 7, 8}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); + + auto offsets = + cudf::test::fixed_width_column_wrapper({0, 10, 20, 30, 40, 50, 60, 70, 80}); + auto d_offsets = cudf::column_view(offsets); + thrust::lower_bound(rmm::exec_policy(stream), + d_offsets.begin(), + d_offsets.end(), + input.begin(), + input.end(), + itr); + expected = cudf::test::fixed_width_column_wrapper({0, 1, 1, 2, 3, 4, 5, 5, 7}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); +} + +namespace { +/** + * For testing creating and using the offsetalator in device code. + */ +struct device_functor_fn { + cudf::column_device_view const d_col; + __device__ int32_t operator()(int idx) + { + auto const itr = cudf::detail::input_offsetalator(d_col.head(), d_col.type()); + return static_cast(itr[idx] * 3); + } +}; +} // namespace + +TYPED_TEST(OffsetalatorTest, device_offsetalator) +{ + using T = TypeParam; + + auto d_col1 = cudf::test::fixed_width_column_wrapper({0, 6, 7, 14, 23, 33, 43, 45, 63}); + auto d_col2 = cudf::test::fixed_width_column_wrapper({0, 0, 0, 0, 0, 0, 0, 0, 0}); + auto input = cudf::column_view(d_col1); + auto output = cudf::mutable_column_view(d_col2); + auto stream = cudf::get_default_stream(); + + auto d_input = cudf::column_device_view::create(input, stream); + + thrust::transform(rmm::exec_policy(stream), + thrust::counting_iterator(0), + thrust::counting_iterator(input.size()), + output.begin(), + device_functor_fn{*d_input}); + + auto expected = + cudf::test::fixed_width_column_wrapper({0, 18, 21, 42, 69, 99, 129, 135, 189}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(d_col2, expected); +} From 4313cfa9b3fcff41f67b48ac8797dc015d441ecc Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Mon, 13 Nov 2023 11:40:36 -0800 Subject: [PATCH 04/44] Use new rapids-dask-dependency metapackage for managing dask versions (#14364) * Update dependency lists * Update wheel building to stop needing manual installations * Update wheel dependency with alpha spec * Rename the package * Update update-version.sh * Update conda/recipes/dask-cudf/meta.yaml Co-authored-by: GALI PREM SAGAR * Make pip/conda dependencies consistent and fix recipe * dfg * Apply suggestions from code review --------- Co-authored-by: GALI PREM SAGAR --- ci/build_wheel.sh | 2 ++ ci/release/update-version.sh | 1 + ci/test_wheel_dask_cudf.sh | 3 --- conda/environments/all_cuda-118_arch-x86_64.yaml | 4 +--- conda/environments/all_cuda-120_arch-x86_64.yaml | 4 +--- conda/recipes/dask-cudf/meta.yaml | 8 +------- dependencies.yaml | 4 +--- python/dask_cudf/pyproject.toml | 3 +-- 8 files changed, 8 insertions(+), 21 deletions(-) diff --git a/ci/build_wheel.sh b/ci/build_wheel.sh index 08716cdb3d9..ae1d9c3fb1a 100755 --- a/ci/build_wheel.sh +++ b/ci/build_wheel.sh @@ -36,6 +36,8 @@ fi if [[ ${package_name} == "dask_cudf" ]]; then sed -r -i "s/cudf==(.*)\"/cudf${PACKAGE_CUDA_SUFFIX}==\1${alpha_spec}\"/g" ${pyproject_file} + sed -r -i "s/dask-cuda==(.*)\"/dask-cuda==\1${alpha_spec}\"/g" ${pyproject_file} + sed -r -i "s/rapids-dask-dependency==(.*)\"/rapids-dask-dependency==\1${alpha_spec}\"/g" ${pyproject_file} else sed -r -i "s/rmm(.*)\"/rmm${PACKAGE_CUDA_SUFFIX}\1${alpha_spec}\"/g" ${pyproject_file} # ptxcompiler and cubinlinker aren't version constrained diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 7574b4174e9..843abd3c3c1 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -80,6 +80,7 @@ DEPENDENCIES=( kvikio libkvikio librmm + rapids-dask-dependency rmm ) for DEP in "${DEPENDENCIES[@]}"; do diff --git a/ci/test_wheel_dask_cudf.sh b/ci/test_wheel_dask_cudf.sh index 118bea753d0..e9162b816aa 100755 --- a/ci/test_wheel_dask_cudf.sh +++ b/ci/test_wheel_dask_cudf.sh @@ -23,9 +23,6 @@ manylinux="manylinux_${manylinux_version}" RAPIDS_PY_WHEEL_NAME="cudf_${manylinux}_${RAPIDS_PY_CUDA_SUFFIX}" rapids-download-wheels-from-s3 ./local-cudf-dep python -m pip install --no-deps ./local-cudf-dep/cudf*.whl -# Always install latest dask for testing -python -m pip install git+https://github.com/dask/dask.git@main git+https://github.com/dask/distributed.git@main git+https://github.com/rapidsai/dask-cuda.git@branch-23.12 - # echo to expand wildcard before adding `[extra]` requires for pip python -m pip install $(echo ./dist/dask_cudf*.whl)[test] diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index adf4fcad32d..9b85888a7b3 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -25,10 +25,7 @@ dependencies: - cupy>=12.0.0 - cxx-compiler - cython>=3.0.3 -- dask-core>=2023.9.2 - dask-cuda==23.12.* -- dask>=2023.9.2 -- distributed>=2023.9.2 - dlpack>=0.5,<0.6.0a0 - doxygen=1.9.1 - fastavro>=0.22.9 @@ -80,6 +77,7 @@ dependencies: - python-snappy>=0.6.0 - python>=3.9,<3.11 - pytorch<1.12.0 +- rapids-dask-dependency==23.12.* - rich - rmm==23.12.* - s3fs>=2022.3.0 diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index a69ef587570..da2b4e109b3 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -27,10 +27,7 @@ dependencies: - cupy>=12.0.0 - cxx-compiler - cython>=3.0.3 -- dask-core>=2023.9.2 - dask-cuda==23.12.* -- dask>=2023.9.2 -- distributed>=2023.9.2 - dlpack>=0.5,<0.6.0a0 - doxygen=1.9.1 - fastavro>=0.22.9 @@ -78,6 +75,7 @@ dependencies: - python-snappy>=0.6.0 - python>=3.9,<3.11 - pytorch<1.12.0 +- rapids-dask-dependency==23.12.* - rich - rmm==23.12.* - s3fs>=2022.3.0 diff --git a/conda/recipes/dask-cudf/meta.yaml b/conda/recipes/dask-cudf/meta.yaml index 9dc9f76d9f5..16638926492 100644 --- a/conda/recipes/dask-cudf/meta.yaml +++ b/conda/recipes/dask-cudf/meta.yaml @@ -37,17 +37,11 @@ build: requirements: host: - python - - cudf ={{ version }} - - dask >=2023.9.2 - - dask-core >=2023.9.2 - - distributed >=2023.9.2 - cuda-version ={{ cuda_version }} run: - python - cudf ={{ version }} - - dask >=2023.9.2 - - dask-core >=2023.9.2 - - distributed >=2023.9.2 + - rapids-dask-dependency ={{ minor_version }} - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} test: diff --git a/dependencies.yaml b/dependencies.yaml index a16b51f4483..b971a682571 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -500,12 +500,10 @@ dependencies: common: - output_types: [conda, requirements, pyproject] packages: - - dask>=2023.9.2 - - distributed>=2023.9.2 + - rapids-dask-dependency==23.12.* - output_types: conda packages: - cupy>=12.0.0 - - dask-core>=2023.9.2 # dask-core in conda is the actual package & dask is the meta package - output_types: pyproject packages: - &cudf cudf==23.12.* diff --git a/python/dask_cudf/pyproject.toml b/python/dask_cudf/pyproject.toml index 32c7bb9fd15..0306da3de46 100644 --- a/python/dask_cudf/pyproject.toml +++ b/python/dask_cudf/pyproject.toml @@ -20,11 +20,10 @@ requires-python = ">=3.9" dependencies = [ "cudf==23.12.*", "cupy-cuda11x>=12.0.0", - "dask>=2023.9.2", - "distributed>=2023.9.2", "fsspec>=0.6.0", "numpy>=1.21,<1.25", "pandas>=1.3,<1.6.0dev0", + "rapids-dask-dependency==23.12.*", ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. classifiers = [ "Intended Audience :: Developers", From 5d09d38bc8ea44e1bdf1fa29e11a820c7417bac5 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Tue, 14 Nov 2023 00:51:42 -0500 Subject: [PATCH 05/44] Always build nvbench statically so we don't need to package it (#14399) Corrects failures seen in C++ CI where libnvbench.so can't be found Authors: - Robert Maynard (https://github.com/robertmaynard) - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14399 --- cpp/cmake/thirdparty/get_nvbench.cmake | 2 +- cpp/cmake/thirdparty/patches/nvbench_override.json | 5 ----- 2 files changed, 1 insertion(+), 6 deletions(-) diff --git a/cpp/cmake/thirdparty/get_nvbench.cmake b/cpp/cmake/thirdparty/get_nvbench.cmake index f0642145fa0..bbd22693ba4 100644 --- a/cpp/cmake/thirdparty/get_nvbench.cmake +++ b/cpp/cmake/thirdparty/get_nvbench.cmake @@ -21,7 +21,7 @@ function(find_and_configure_nvbench) set(cudf_patch_dir "${CMAKE_CURRENT_FUNCTION_LIST_DIR}/patches") rapids_cpm_package_override("${cudf_patch_dir}/nvbench_override.json") - rapids_cpm_nvbench() + rapids_cpm_nvbench(BUILD_STATIC) endfunction() diff --git a/cpp/cmake/thirdparty/patches/nvbench_override.json b/cpp/cmake/thirdparty/patches/nvbench_override.json index 7be868081b6..ad9b19c29c1 100644 --- a/cpp/cmake/thirdparty/patches/nvbench_override.json +++ b/cpp/cmake/thirdparty/patches/nvbench_override.json @@ -7,11 +7,6 @@ "file" : "${current_json_dir}/nvbench_global_setup.diff", "issue" : "Fix add support for global setup to initialize RMM in nvbench [https://github.com/NVIDIA/nvbench/pull/123]", "fixed_in" : "" - }, - { - "file" : "nvbench/use_existing_fmt.diff", - "issue" : "Fix add support for using an existing fmt [https://github.com/NVIDIA/nvbench/pull/125]", - "fixed_in" : "" } ] } From e982d3736f095e680298af85bde732d9b5a73122 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath <3190405+shwina@users.noreply.github.com> Date: Tue, 14 Nov 2023 09:51:02 -0500 Subject: [PATCH 06/44] cudf.pandas: cuDF subpath checking in module `__getattr__` (#14388) Closes https://github.com/rapidsai/cudf/issues/14384. `x.startswith(y)` is not a good enough check for if `x` is a subdirectory of `y`. It causes `pandasai` to be reported as a sub-package of `pandas`. Authors: - Ashwin Srinath (https://github.com/shwina) Approvers: - https://github.com/brandon-b-miller URL: https://github.com/rapidsai/cudf/pull/14388 --- python/cudf/cudf/pandas/module_accelerator.py | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/python/cudf/cudf/pandas/module_accelerator.py b/python/cudf/cudf/pandas/module_accelerator.py index eb35c4adaaf..180d75d96e8 100644 --- a/python/cudf/cudf/pandas/module_accelerator.py +++ b/python/cudf/cudf/pandas/module_accelerator.py @@ -10,6 +10,7 @@ import importlib.abc import importlib.machinery import os +import pathlib import sys import threading import warnings @@ -554,9 +555,10 @@ def getattr_real_or_wrapped( frame = sys._getframe() # We cannot possibly be at the top level. assert frame.f_back - calling_module = frame.f_back.f_code.co_filename + calling_module = pathlib.PurePath(frame.f_back.f_code.co_filename) use_real = any( - calling_module.startswith(path) for path in loader._denylist + calling_module.is_relative_to(path) + for path in loader._denylist ) try: if use_real: From 7f3fba164c4dd28c701ea2941d0525fc782a639c Mon Sep 17 00:00:00 2001 From: Jeremy Dyer Date: Tue, 14 Nov 2023 12:02:10 -0500 Subject: [PATCH 07/44] Refactor cudf_kafka to use skbuild (#14292) Refactor the currently outdated cudf_kafka build setup to use skbuild instead. Authors: - Jeremy Dyer (https://github.com/jdye64) - Bradley Dice (https://github.com/bdice) - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Bradley Dice (https://github.com/bdice) - AJ Schmidt (https://github.com/ajschmidt8) URL: https://github.com/rapidsai/cudf/pull/14292 --- build.sh | 2 +- ci/release/update-version.sh | 1 + .../all_cuda-120_arch-x86_64.yaml | 1 - conda/recipes/cudf_kafka/build.sh | 13 --- .../cudf_kafka/conda_build_config.yaml | 6 ++ conda/recipes/cudf_kafka/meta.yaml | 21 ++-- cpp/libcudf_kafka/CMakeLists.txt | 8 +- .../cmake/thirdparty/get_cudf.cmake | 16 +-- cpp/libcudf_kafka/tests/CMakeLists.txt | 2 +- dependencies.yaml | 13 +-- python/cudf/cudf/_lib/CMakeLists.txt | 6 -- python/cudf_kafka/CMakeLists.txt | 47 +++++++++ python/cudf_kafka/LICENSE | 1 + python/cudf_kafka/README.md | 1 + .../cudf_kafka/cudf_kafka/_lib/CMakeLists.txt | 62 ++++++++++++ python/cudf_kafka/cudf_kafka/_lib/kafka.pxd | 4 +- python/cudf_kafka/pyproject.toml | 1 + python/cudf_kafka/setup.py | 97 ++----------------- 18 files changed, 160 insertions(+), 142 deletions(-) create mode 100644 python/cudf_kafka/CMakeLists.txt create mode 120000 python/cudf_kafka/LICENSE create mode 120000 python/cudf_kafka/README.md create mode 100644 python/cudf_kafka/cudf_kafka/_lib/CMakeLists.txt diff --git a/build.sh b/build.sh index 2ad69712e5d..e5beb51dedf 100755 --- a/build.sh +++ b/build.sh @@ -369,7 +369,7 @@ fi # build cudf_kafka Python package if hasArg cudf_kafka; then cd ${REPODIR}/python/cudf_kafka - SKBUILD_CONFIGURE_OPTIONS="-DCMAKE_LIBRARY_PATH=${LIBCUDF_BUILD_DIR}" \ + SKBUILD_CONFIGURE_OPTIONS="-DCMAKE_PREFIX_PATH=${INSTALL_PREFIX} -DCMAKE_LIBRARY_PATH=${LIBCUDF_BUILD_DIR} ${EXTRA_CMAKE_ARGS}" \ SKBUILD_BUILD_OPTIONS="-j${PARALLEL_LEVEL:-1}" \ python -m pip install --no-build-isolation --no-deps . fi diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 843abd3c3c1..4f1cbc47d1d 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -43,6 +43,7 @@ sed_runner 's/'"VERSION ${CURRENT_SHORT_TAG}.*"'/'"VERSION ${NEXT_FULL_TAG}"'/g' # Python CMakeLists updates sed_runner 's/'"cudf_version .*)"'/'"cudf_version ${NEXT_FULL_TAG})"'/g' python/cudf/CMakeLists.txt +sed_runner 's/'"cudf_kafka_version .*)"'/'"cudf_kafka_version ${NEXT_FULL_TAG})"'/g' python/cudf_kafka/CMakeLists.txt # cpp libcudf_kafka update sed_runner 's/'"VERSION ${CURRENT_SHORT_TAG}.*"'/'"VERSION ${NEXT_FULL_TAG}"'/g' cpp/libcudf_kafka/CMakeLists.txt diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index da2b4e109b3..a3eeb3dd99f 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -17,7 +17,6 @@ dependencies: - cmake>=3.26.4 - cramjam - cuda-cudart-dev -- cuda-gdb - cuda-nvcc - cuda-nvrtc-dev - cuda-nvtx-dev diff --git a/conda/recipes/cudf_kafka/build.sh b/conda/recipes/cudf_kafka/build.sh index f4bb6e1bc91..9458349d101 100644 --- a/conda/recipes/cudf_kafka/build.sh +++ b/conda/recipes/cudf_kafka/build.sh @@ -1,16 +1,3 @@ # Copyright (c) 2020-2023, NVIDIA CORPORATION. -# This assumes the script is executed from the root of the repo directory -# Need to set CUDA_HOME inside conda environments because the hacked together -# setup.py for cudf-kafka searches that way. -# TODO: Remove after https://github.com/rapidsai/cudf/pull/14292 updates -# cudf_kafka to use scikit-build -CUDA_MAJOR=${RAPIDS_CUDA_VERSION%%.*} -if [[ ${CUDA_MAJOR} == "12" ]]; then - target_name="x86_64-linux" - if [[ ! $(arch) == "x86_64" ]]; then - target_name="sbsa-linux" - fi - export CUDA_HOME="${PREFIX}/targets/${target_name}/" -fi ./build.sh -v cudf_kafka diff --git a/conda/recipes/cudf_kafka/conda_build_config.yaml b/conda/recipes/cudf_kafka/conda_build_config.yaml index b63a136ad2d..c98c2701653 100644 --- a/conda/recipes/cudf_kafka/conda_build_config.yaml +++ b/conda/recipes/cudf_kafka/conda_build_config.yaml @@ -9,3 +9,9 @@ sysroot_version: cmake_version: - ">=3.26.4" + +cuda_compiler: + - cuda-nvcc + +cuda11_compiler: + - nvcc diff --git a/conda/recipes/cudf_kafka/meta.yaml b/conda/recipes/cudf_kafka/meta.yaml index 9440f8bf124..343ec2519f1 100644 --- a/conda/recipes/cudf_kafka/meta.yaml +++ b/conda/recipes/cudf_kafka/meta.yaml @@ -33,28 +33,31 @@ build: - SCCACHE_S3_KEY_PREFIX=cudf-kafka-linux64 # [linux64] - SCCACHE_S3_USE_SSL - SCCACHE_S3_NO_CREDENTIALS - # TODO: Remove after https://github.com/rapidsai/cudf/pull/14292 updates - # cudf_kafka to use scikit-build - - RAPIDS_CUDA_VERSION + ignore_run_exports_from: + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} + {% endif %} requirements: build: - cmake {{ cmake_version }} + - ninja - {{ compiler('c') }} - {{ compiler('cxx') }} - - ninja - - sysroot_{{ target_platform }} {{ sysroot_version }} - # TODO: Remove after https://github.com/rapidsai/cudf/pull/14292 updates - # cudf_kafka to use scikit-build - {% if cuda_major == "12" %} - - cuda-gdb + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} ={{ cuda_version }} + {% else %} + - {{ compiler('cuda') }} {% endif %} + - cuda-version ={{ cuda_version }} + - sysroot_{{ target_platform }} {{ sysroot_version }} host: - python - cython >=3.0.3 - cuda-version ={{ cuda_version }} - cudf ={{ version }} - libcudf_kafka ={{ version }} + - scikit-build >=0.13.1 - setuptools {% if cuda_major == "12" %} - cuda-cudart-dev diff --git a/cpp/libcudf_kafka/CMakeLists.txt b/cpp/libcudf_kafka/CMakeLists.txt index 1a15a3ec2cd..4128afa3935 100644 --- a/cpp/libcudf_kafka/CMakeLists.txt +++ b/cpp/libcudf_kafka/CMakeLists.txt @@ -21,7 +21,7 @@ include(rapids-export) include(rapids-find) project( - CUDA_KAFKA + CUDF_KAFKA VERSION 23.12.00 LANGUAGES CXX ) @@ -64,7 +64,7 @@ add_library(cudf_kafka SHARED src/kafka_consumer.cpp src/kafka_callback.cpp) # ################################################################################################## # * include paths --------------------------------------------------------------------------------- target_include_directories( - cudf_kafka PUBLIC "$" + cudf_kafka PUBLIC "$" "$" ) @@ -85,6 +85,8 @@ set_target_properties( CXX_STANDARD_REQUIRED ON ) +add_library(cudf_kafka::cudf_kafka ALIAS cudf_kafka) + # ################################################################################################## # * cudf_kafka Install ---------------------------------------------------------------------------- rapids_cmake_install_lib_dir(lib_dir) @@ -94,7 +96,7 @@ install( EXPORT cudf_kafka-exports ) -install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include DESTINATION include) +install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) rapids_export( INSTALL cudf_kafka diff --git a/cpp/libcudf_kafka/cmake/thirdparty/get_cudf.cmake b/cpp/libcudf_kafka/cmake/thirdparty/get_cudf.cmake index aa4c5b60e7a..20aa9873f43 100644 --- a/cpp/libcudf_kafka/cmake/thirdparty/get_cudf.cmake +++ b/cpp/libcudf_kafka/cmake/thirdparty/get_cudf.cmake @@ -1,5 +1,5 @@ # ============================================================================= -# Copyright (c) 2021-2022, NVIDIA CORPORATION. +# Copyright (c) 2021-2023, NVIDIA 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 @@ -35,21 +35,21 @@ function(find_and_configure_cudf VERSION) endif() endfunction() -set(CUDA_KAFKA_MIN_VERSION_cudf - "${CUDA_KAFKA_VERSION_MAJOR}.${CUDA_KAFKA_VERSION_MINOR}.${CUDA_KAFKA_VERSION_PATCH}" +set(CUDF_KAFKA_MIN_VERSION + "${CUDF_KAFKA_VERSION_MAJOR}.${CUDF_KAFKA_VERSION_MINOR}.${CUDF_KAFKA_VERSION_PATCH}" ) -find_and_configure_cudf(${CUDA_KAFKA_MIN_VERSION_cudf}) +find_and_configure_cudf(${CUDF_KAFKA_MIN_VERSION}) if(cudf_REQUIRES_CUDA) - rapids_cuda_init_architectures(CUDA_KAFKA) + rapids_cuda_init_architectures(CUDF_KAFKA) # Since we are building cudf as part of ourselves we need to enable the CUDA language in the # top-most scope enable_language(CUDA) - # Since CUDA_KAFKA only enables CUDA optionally we need to manually include the file that + # Since CUDF_KAFKA only enables CUDA optionally we need to manually include the file that # rapids_cuda_init_architectures relies on `project` calling - if(DEFINED CMAKE_PROJECT_CUDA_KAFKA_INCLUDE) - include("${CMAKE_PROJECT_CUDA_KAFKA_INCLUDE}") + if(DEFINED CMAKE_PROJECT_CUDF_KAFKA_INCLUDE) + include("${CMAKE_PROJECT_CUDF_KAFKA_INCLUDE}") endif() endif() diff --git a/cpp/libcudf_kafka/tests/CMakeLists.txt b/cpp/libcudf_kafka/tests/CMakeLists.txt index 68a5327b455..b819cb6fc3b 100644 --- a/cpp/libcudf_kafka/tests/CMakeLists.txt +++ b/cpp/libcudf_kafka/tests/CMakeLists.txt @@ -26,7 +26,7 @@ function(ConfigureTest test_name) add_executable(${test_name} ${ARGN}) set_target_properties( ${test_name} - PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" + PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" INSTALL_RPATH "\$ORIGIN/../../../lib" ) target_link_libraries( diff --git a/dependencies.yaml b/dependencies.yaml index b971a682571..97149a5e2ba 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -9,8 +9,8 @@ files: - build_all - build_cpp - build_wheels - - build_python - build_python_common + - build_python_cudf - cudatoolkit - develop - docs @@ -71,8 +71,8 @@ files: table: build-system includes: - build_all - - build_python - build_python_common + - build_python_cudf - build_wheels py_run_cudf: output: pyproject @@ -138,8 +138,8 @@ files: extras: table: build-system includes: - - build_wheels - build_python_common + - build_wheels py_run_cudf_kafka: output: pyproject pyproject_dir: python/cudf_kafka @@ -259,16 +259,16 @@ dependencies: - cython>=3.0.3 # TODO: Pin to numpy<1.25 until cudf requires pandas 2 - &numpy numpy>=1.21,<1.25 + - scikit-build>=0.13.1 - output_types: [conda, requirements, pyproject] packages: # Hard pin the patch version used during the build. This must be kept # in sync with the version pinned in get_arrow.cmake. - pyarrow==14.0.1.* - build_python: + build_python_cudf: common: - output_types: [conda, requirements, pyproject] packages: - - scikit-build>=0.13.1 - rmm==23.12.* - output_types: conda packages: @@ -302,9 +302,6 @@ dependencies: - cuda-nvrtc-dev - cuda-nvtx-dev - libcurand-dev - # TODO: Remove after https://github.com/rapidsai/cudf/pull/14292 updates - # cudf_kafka to use scikit-build - - cuda-gdb - matrix: cuda: "11.8" packages: diff --git a/python/cudf/cudf/_lib/CMakeLists.txt b/python/cudf/cudf/_lib/CMakeLists.txt index 1b543b94589..c041c7f4842 100644 --- a/python/cudf/cudf/_lib/CMakeLists.txt +++ b/python/cudf/cudf/_lib/CMakeLists.txt @@ -81,12 +81,6 @@ target_link_libraries(strings_udf cudf_strings_udf) # necessary. The relevant command is tar -xf /opt/_internal/static-libs-for-embedding-only.tar.xz -C # /opt/_internal" find_package(NumPy REQUIRED) -set(targets_using_numpy interop avro csv orc json parquet) -foreach(target IN LISTS targets_using_numpy) - target_include_directories(${target} PRIVATE "${NumPy_INCLUDE_DIRS}") - # Switch to the line below when we switch back to FindPython.cmake in CMake 3.24. - # target_include_directories(${target} PRIVATE "${Python_NumPy_INCLUDE_DIRS}") -endforeach() set(targets_using_dlpack interop) foreach(target IN LISTS targets_using_dlpack) diff --git a/python/cudf_kafka/CMakeLists.txt b/python/cudf_kafka/CMakeLists.txt new file mode 100644 index 00000000000..d55c3fdc076 --- /dev/null +++ b/python/cudf_kafka/CMakeLists.txt @@ -0,0 +1,47 @@ +# ============================================================================= +# Copyright (c) 2022-2023, NVIDIA 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. +# ============================================================================= + +cmake_minimum_required(VERSION 3.26.4 FATAL_ERROR) + +set(cudf_kafka_version 23.12.00) + +include(../../fetch_rapids.cmake) + +project( + cudf-kafka-python + VERSION ${cudf_kafka_version} + LANGUAGES # TODO: Building Python extension modules via the python_extension_module requires the C + # language to be enabled here. The test project that is built in scikit-build to verify + # various linking options for the python library is hardcoded to build with C, so until + # that is fixed we need to keep C. + C CXX +) + +find_package(cudf_kafka ${cudf_kafka_version} REQUIRED) + +if(NOT cudf_kafka_FOUND) + message( + FATAL_ERROR + "cudf_kafka package not found. cudf_kafka C++ is required to build this Python package." + ) +endif() + +include(rapids-cython) +rapids_cython_init() + +add_subdirectory(cudf_kafka/_lib) + +if(DEFINED cython_lib_dir) + rapids_cython_add_rpath_entries(TARGET cudf_kafka PATHS "${cython_lib_dir}") +endif() diff --git a/python/cudf_kafka/LICENSE b/python/cudf_kafka/LICENSE new file mode 120000 index 00000000000..30cff7403da --- /dev/null +++ b/python/cudf_kafka/LICENSE @@ -0,0 +1 @@ +../../LICENSE \ No newline at end of file diff --git a/python/cudf_kafka/README.md b/python/cudf_kafka/README.md new file mode 120000 index 00000000000..fe840054137 --- /dev/null +++ b/python/cudf_kafka/README.md @@ -0,0 +1 @@ +../../README.md \ No newline at end of file diff --git a/python/cudf_kafka/cudf_kafka/_lib/CMakeLists.txt b/python/cudf_kafka/cudf_kafka/_lib/CMakeLists.txt new file mode 100644 index 00000000000..3262b7d5ebe --- /dev/null +++ b/python/cudf_kafka/cudf_kafka/_lib/CMakeLists.txt @@ -0,0 +1,62 @@ +# ============================================================================= +# Copyright (c) 2022-2023, NVIDIA 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. +# ============================================================================= + +set(cython_sources kafka.pyx) +set(linked_libraries cudf_kafka::cudf_kafka) + +rapids_cython_create_modules( + CXX ASSOCIATED_TARGETS cudf_kafka + SOURCE_FILES "${cython_sources}" + LINKED_LIBRARIES "${linked_libraries}" +) + +# TODO: Finding NumPy currently requires finding Development due to a bug in CMake. This bug was +# fixed in https://gitlab.kitware.com/cmake/cmake/-/merge_requests/7410 and will be available in +# CMake 3.24, so we can remove the Development component once we upgrade to CMake 3.24. +# find_package(Python REQUIRED COMPONENTS Development NumPy) + +# Note: The bug noted above prevents us from finding NumPy successfully using FindPython.cmake +# inside the manylinux images used to build wheels because manylinux images do not contain +# libpython.so and therefore Development cannot be found. Until we upgrade to CMake 3.24, we should +# use FindNumpy.cmake instead (provided by scikit-build). When we switch to 3.24 we can try +# switching back, but it may not work if that implicitly still requires Python libraries. In that +# case we'll need to follow up with the CMake team to remove that dependency. The stopgap solution +# is to unpack the static lib tarballs in the wheel building jobs so that there are at least static +# libs to be found, but that should be a last resort since it implies a dependency that isn't really +# necessary. The relevant command is tar -xf /opt/_internal/static-libs-for-embedding-only.tar.xz -C +# /opt/_internal" +find_package(NumPy REQUIRED) + +find_package(Python 3.9 REQUIRED COMPONENTS Interpreter) + +execute_process( + COMMAND "${Python_EXECUTABLE}" -c "import pyarrow; print(pyarrow.get_include())" + OUTPUT_VARIABLE PYARROW_INCLUDE_DIR + ERROR_VARIABLE PYARROW_ERROR + RESULT_VARIABLE PYARROW_RESULT + OUTPUT_STRIP_TRAILING_WHITESPACE +) + +if(${PYARROW_RESULT}) + message(FATAL_ERROR "Error while trying to obtain pyarrow include directory:\n${PYARROW_ERROR}") +endif() + +# TODO: Due to cudf's scalar.pyx needing to cimport pylibcudf's scalar.pyx (because there are parts +# of cudf Cython that need to directly access the c_obj underlying the pylibcudf Scalar) the +# requirement for arrow headers infects all of cudf. That in turn requires including numpy headers. +# These requirements will go away once all scalar-related Cython code is removed from cudf. +foreach(target IN LISTS RAPIDS_CYTHON_CREATED_TARGETS) + target_include_directories(${target} PRIVATE "${NumPy_INCLUDE_DIRS}") + target_include_directories(${target} PRIVATE "${PYARROW_INCLUDE_DIR}") +endforeach() diff --git a/python/cudf_kafka/cudf_kafka/_lib/kafka.pxd b/python/cudf_kafka/cudf_kafka/_lib/kafka.pxd index ca729c62512..068837d04ee 100644 --- a/python/cudf_kafka/cudf_kafka/_lib/kafka.pxd +++ b/python/cudf_kafka/cudf_kafka/_lib/kafka.pxd @@ -11,12 +11,12 @@ from cudf._lib.cpp.io.datasource cimport datasource from cudf._lib.io.datasource cimport Datasource -cdef extern from "kafka_callback.hpp" \ +cdef extern from "cudf_kafka/kafka_callback.hpp" \ namespace "cudf::io::external::kafka" nogil: ctypedef object (*python_callable_type)() -cdef extern from "kafka_consumer.hpp" \ +cdef extern from "cudf_kafka/kafka_consumer.hpp" \ namespace "cudf::io::external::kafka" nogil: cpdef cppclass kafka_consumer: diff --git a/python/cudf_kafka/pyproject.toml b/python/cudf_kafka/pyproject.toml index 4829f06ab09..15431161d75 100644 --- a/python/cudf_kafka/pyproject.toml +++ b/python/cudf_kafka/pyproject.toml @@ -6,6 +6,7 @@ requires = [ "cython>=3.0.3", "numpy>=1.21,<1.25", "pyarrow==14.0.1.*", + "scikit-build>=0.13.1", "setuptools", "wheel", ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. diff --git a/python/cudf_kafka/setup.py b/python/cudf_kafka/setup.py index 6f3909d4528..6a99e9ed968 100644 --- a/python/cudf_kafka/setup.py +++ b/python/cudf_kafka/setup.py @@ -1,96 +1,13 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. -import os -import shutil -import sysconfig -from distutils.sysconfig import get_python_lib - -import numpy as np -import pyarrow as pa -from Cython.Build import cythonize -from setuptools import find_packages, setup -from setuptools.extension import Extension - -cython_files = ["cudf_kafka/_lib/*.pyx"] - -CUDA_HOME = os.environ.get("CUDA_HOME", False) -if not CUDA_HOME: - path_to_cuda_gdb = shutil.which("cuda-gdb") - if path_to_cuda_gdb is None: - raise OSError( - "Could not locate CUDA. " - "Please set the environment variable " - "CUDA_HOME to the path to the CUDA installation " - "and try again." - ) - CUDA_HOME = os.path.dirname(os.path.dirname(path_to_cuda_gdb)) - -if not os.path.isdir(CUDA_HOME): - raise OSError(f"Invalid CUDA_HOME: directory does not exist: {CUDA_HOME}") - -cuda_include_dir = os.path.join(CUDA_HOME, "include") - -CUDF_ROOT = os.environ.get( - "CUDF_ROOT", - os.path.abspath( - os.path.join( - os.path.dirname(os.path.abspath(__file__)), "../../cpp/build/" - ) - ), -) -CUDF_KAFKA_ROOT = os.environ.get( - "CUDF_KAFKA_ROOT", "../../cpp/libcudf_kafka/build" -) - -try: - nthreads = int(os.environ.get("PARALLEL_LEVEL", "0") or "0") -except Exception: - nthreads = 0 - -extensions = [ - Extension( - "*", - sources=cython_files, - include_dirs=[ - os.path.abspath(os.path.join(CUDF_ROOT, "../include/cudf")), - os.path.abspath(os.path.join(CUDF_ROOT, "../include")), - os.path.abspath( - os.path.join(CUDF_ROOT, "../libcudf_kafka/include/cudf_kafka") - ), - os.path.join(CUDF_ROOT, "include"), - os.path.join(CUDF_ROOT, "_deps/libcudacxx-src/include"), - os.path.join( - os.path.dirname(sysconfig.get_path("include")), - "rapids/libcudacxx", - ), - os.path.dirname(sysconfig.get_path("include")), - np.get_include(), - pa.get_include(), - cuda_include_dir, - ], - library_dirs=( - [ - get_python_lib(), - os.path.join(os.sys.prefix, "lib"), - CUDF_KAFKA_ROOT, - ] - ), - libraries=["cudf", "cudf_kafka"], - language="c++", - extra_compile_args=["-std=c++17", "-DFMT_HEADER_ONLY=1"], - ) -] +# Copyright (c) 2018-2023, NVIDIA CORPORATION. +from setuptools import find_packages +from skbuild import setup packages = find_packages(include=["cudf_kafka*"]) + setup( - # Include the separately-compiled shared library - ext_modules=cythonize( - extensions, - nthreads=nthreads, - compiler_directives=dict( - profile=False, language_level=3, embedsignature=True - ), - ), packages=packages, - package_data={key: ["VERSION", "*.pxd"] for key in packages}, + package_data={ + key: ["VERSION", "*.pxd", "*.hpp", "*.cuh"] for key in packages + }, zip_safe=False, ) From b0c1b7b82ccdf1a7e4159cb3bffa1984092440d4 Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Tue, 14 Nov 2023 12:48:32 -0500 Subject: [PATCH 08/44] Add BytePairEncoder class to cuDF (#13891) Adds a new BytePairEncoding class to cuDF ``` >>> import cudf >>> from cudf.core.byte_pair_encoding import BytePairEncoder >>> mps = cudf.read_text('merges.txt', delimiter='\n', strip_delimiters=True) >>> bpe = BytePairEncoder(mps) >>> str_series = cudf.Series(['This is a sentence', 'thisisit']) >>> bpe(str_series) 0 This is a sent ence 1 this is it dtype: object ``` This class wraps the existing `nvtext::byte_pair_encoding` APIs to load the merge-pairs data and encode a column of strings. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/13891 --- .../cudf/_lib/cpp/nvtext/byte_pair_encode.pxd | 24 ++++++++ python/cudf/cudf/_lib/nvtext/CMakeLists.txt | 4 +- .../cudf/_lib/nvtext/byte_pair_encode.pyx | 50 ++++++++++++++++ python/cudf/cudf/core/byte_pair_encoding.py | 59 +++++++++++++++++++ .../cudf/cudf/tests/text/test_text_methods.py | 41 +++++++++++++ 5 files changed, 176 insertions(+), 2 deletions(-) create mode 100644 python/cudf/cudf/_lib/cpp/nvtext/byte_pair_encode.pxd create mode 100644 python/cudf/cudf/_lib/nvtext/byte_pair_encode.pyx create mode 100644 python/cudf/cudf/core/byte_pair_encoding.py diff --git a/python/cudf/cudf/_lib/cpp/nvtext/byte_pair_encode.pxd b/python/cudf/cudf/_lib/cpp/nvtext/byte_pair_encode.pxd new file mode 100644 index 00000000000..e678e4e84db --- /dev/null +++ b/python/cudf/cudf/_lib/cpp/nvtext/byte_pair_encode.pxd @@ -0,0 +1,24 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. + +from libcpp.memory cimport unique_ptr +from libcpp.string cimport string + +from cudf._lib.cpp.column.column cimport column +from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.cpp.scalar.scalar cimport string_scalar + + +cdef extern from "nvtext/byte_pair_encoding.hpp" namespace "nvtext" nogil: + + cdef struct bpe_merge_pairs "nvtext::bpe_merge_pairs": + pass + + cdef unique_ptr[bpe_merge_pairs] load_merge_pairs( + const column_view &merge_pairs + ) except + + + cdef unique_ptr[column] byte_pair_encoding( + const column_view &strings, + const bpe_merge_pairs &merge_pairs, + const string_scalar &separator + ) except + diff --git a/python/cudf/cudf/_lib/nvtext/CMakeLists.txt b/python/cudf/cudf/_lib/nvtext/CMakeLists.txt index d4e2392ee04..d7cbdeb5bda 100644 --- a/python/cudf/cudf/_lib/nvtext/CMakeLists.txt +++ b/python/cudf/cudf/_lib/nvtext/CMakeLists.txt @@ -13,8 +13,8 @@ # ============================================================================= set(cython_sources - edit_distance.pyx generate_ngrams.pyx jaccard.pyx minhash.pyx ngrams_tokenize.pyx normalize.pyx - replace.pyx stemmer.pyx subword_tokenize.pyx tokenize.pyx + byte_pair_encode.pyx edit_distance.pyx generate_ngrams.pyx jaccard.pyx minhash.pyx + ngrams_tokenize.pyx normalize.pyx replace.pyx stemmer.pyx subword_tokenize.pyx tokenize.pyx ) set(linked_libraries cudf::cudf) rapids_cython_create_modules( diff --git a/python/cudf/cudf/_lib/nvtext/byte_pair_encode.pyx b/python/cudf/cudf/_lib/nvtext/byte_pair_encode.pyx new file mode 100644 index 00000000000..cfc76afa8a5 --- /dev/null +++ b/python/cudf/cudf/_lib/nvtext/byte_pair_encode.pyx @@ -0,0 +1,50 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. + + +from cudf.core.buffer import acquire_spill_lock + +from libcpp.memory cimport unique_ptr +from libcpp.utility cimport move + +from cudf._lib.column cimport Column +from cudf._lib.cpp.column.column cimport column +from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.cpp.nvtext.byte_pair_encode cimport ( + bpe_merge_pairs as cpp_bpe_merge_pairs, + byte_pair_encoding as cpp_byte_pair_encoding, + load_merge_pairs as cpp_load_merge_pairs, +) +from cudf._lib.cpp.scalar.scalar cimport string_scalar +from cudf._lib.scalar cimport DeviceScalar + + +cdef class BPEMergePairs: + cdef unique_ptr[cpp_bpe_merge_pairs] c_obj + + def __cinit__(self, Column merge_pairs): + cdef column_view c_pairs = merge_pairs.view() + with nogil: + self.c_obj = move(cpp_load_merge_pairs(c_pairs)) + + +@acquire_spill_lock() +def byte_pair_encoding( + Column strings, + BPEMergePairs merge_pairs, + object separator +): + cdef column_view c_strings = strings.view() + cdef DeviceScalar d_separator = separator.device_value + cdef const string_scalar* c_separator = d_separator\ + .get_raw_ptr() + cdef unique_ptr[column] c_result + with nogil: + c_result = move( + cpp_byte_pair_encoding( + c_strings, + merge_pairs.c_obj.get()[0], + c_separator[0] + ) + ) + + return Column.from_unique_ptr(move(c_result)) diff --git a/python/cudf/cudf/core/byte_pair_encoding.py b/python/cudf/cudf/core/byte_pair_encoding.py new file mode 100644 index 00000000000..4c881022ecf --- /dev/null +++ b/python/cudf/cudf/core/byte_pair_encoding.py @@ -0,0 +1,59 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. + +from __future__ import annotations + +import cudf +from cudf._lib.nvtext.byte_pair_encode import ( + BPEMergePairs as cpp_merge_pairs, + byte_pair_encoding as cpp_byte_pair_encoding, +) + + +class BytePairEncoder: + """ + Given a merge pairs strings series, performs byte pair encoding on + a strings series using the provided separator. + + Parameters + ---------- + merges_pairs : str + Strings column of merge pairs + + Returns + ------- + BytePairEncoder + """ + + def __init__(self, merges_pair: "cudf.Series"): + self.merge_pairs = cpp_merge_pairs(merges_pair._column) + + def __call__(self, text, separator: str = " "): + """ + + Parameters + ---------- + text : cudf string series + The strings to be encoded. + + Returns + ------- + Encoded strings + + Examples + -------- + >>> import cudf + >>> from cudf.core.byte_pair_encoding import BytePairEncoder + >>> mps = cudf.Series(["e n", "i t", "i s", "e s", "en t", + ... "c e", "es t", "en ce", "T h", "Th is", + ... "t est", "s ent", "t h", "th is"]) + >>> bpe = BytePairEncoder(mps) + >>> str_series = cudf.Series(['This is the sentence', 'thisisit']) + >>> bpe(str_series) + 0 This is a sent ence + 1 this is it + dtype: object + """ + sep = cudf.Scalar(separator, dtype="str") + result = cpp_byte_pair_encoding(text._column, self.merge_pairs, sep) + + return cudf.Series(result) diff --git a/python/cudf/cudf/tests/text/test_text_methods.py b/python/cudf/cudf/tests/text/test_text_methods.py index e565df8f3da..2dccd583b23 100644 --- a/python/cudf/cudf/tests/text/test_text_methods.py +++ b/python/cudf/cudf/tests/text/test_text_methods.py @@ -7,6 +7,7 @@ import pytest import cudf +from cudf.core.byte_pair_encoding import BytePairEncoder from cudf.core.tokenize_vocabulary import TokenizeVocabulary from cudf.testing._utils import assert_eq @@ -1024,3 +1025,43 @@ def test_jaccard_index_random_strings(): actual = str1.str.jaccard_index(str2, jaccard_width) assert_eq(expected, actual) + + +@pytest.mark.parametrize( + "separator, input, results", + [ + (" ", "thetestsentence", "the test sent ence"), + ("_", "sentenceistest", "sent_ence_is_test"), + ("$", "istestsentencehere", "is$test$sent$ence$he$r$e"), + ], +) +def test_byte_pair_encoding(separator, input, results): + pairs_table = cudf.Series( + [ + "t he", + "h e", + "e n", + "i t", + "i s", + "e s", + "en t", + "c e", + "es t", + "en ce", + "t h", + "h i", + "th is", + "t est", + "s i", + "s ent", + ] + ) + encoder = BytePairEncoder(pairs_table) + + strings = cudf.Series([input, None, "", input]) + + expected = cudf.Series([results, None, "", results]) + + actual = encoder(strings, separator) + assert type(expected) == type(actual) + assert_eq(expected, actual) From b446a6f187241e765c925da1053ece2679313a06 Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Tue, 14 Nov 2023 12:49:19 -0500 Subject: [PATCH 09/44] Fix token-count logic in nvtext::tokenize_with_vocabulary (#14393) Fixes a bug introduced in #14336 when trying to simplify the token-counting logic as per this discussion https://github.com/rapidsai/cudf/pull/14336#discussion_r1378173552 The simplification caused an error which was found when running the nvtext benchmarks. The appropriate gtest has been updated to cover this case now. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Bradley Dice (https://github.com/bdice) - Karthikeyan (https://github.com/karthikeyann) URL: https://github.com/rapidsai/cudf/pull/14393 --- cpp/benchmarks/text/vocab.cpp | 2 +- cpp/src/text/vocabulary_tokenize.cu | 8 ++++++-- cpp/tests/text/tokenize_tests.cpp | 12 ++++++------ 3 files changed, 13 insertions(+), 9 deletions(-) diff --git a/cpp/benchmarks/text/vocab.cpp b/cpp/benchmarks/text/vocab.cpp index 6922b7214ff..80942e2697d 100644 --- a/cpp/benchmarks/text/vocab.cpp +++ b/cpp/benchmarks/text/vocab.cpp @@ -53,7 +53,7 @@ static void bench_vocab_tokenize(nvbench::state& state) auto const vocab_col = [] { data_profile const profile = data_profile_builder().no_validity().distribution( - cudf::type_id::STRING, distribution_id::NORMAL, 0, 5); + cudf::type_id::STRING, distribution_id::NORMAL, 0, 15); auto const col = create_random_column(cudf::type_id::STRING, row_count{100}, profile); return cudf::strings::filter_characters_of_type( cudf::strings_column_view(col->view()), diff --git a/cpp/src/text/vocabulary_tokenize.cu b/cpp/src/text/vocabulary_tokenize.cu index 41f8c0a8731..511f1995374 100644 --- a/cpp/src/text/vocabulary_tokenize.cu +++ b/cpp/src/text/vocabulary_tokenize.cu @@ -276,8 +276,12 @@ __global__ void token_counts_fn(cudf::column_device_view const d_strings, __syncwarp(); for (auto itr = d_output + lane_idx + 1; itr < d_output_end; itr += cudf::detail::warp_size) { - // add one if at the edge of a token or at the string's end - count += ((*itr && !(*(itr - 1))) || (itr + 1 == d_output_end)); + // add one if at the edge of a token or if at the string's end + if (*itr) { + count += !(*(itr - 1)); + } else { + count += (itr + 1 == d_output_end); + } } __syncwarp(); diff --git a/cpp/tests/text/tokenize_tests.cpp b/cpp/tests/text/tokenize_tests.cpp index 8118183a458..ea36e13de6f 100644 --- a/cpp/tests/text/tokenize_tests.cpp +++ b/cpp/tests/text/tokenize_tests.cpp @@ -246,14 +246,14 @@ TEST_F(TextTokenizeTest, Vocabulary) TEST_F(TextTokenizeTest, VocabularyLongStrings) { - cudf::test::strings_column_wrapper vocabulary( // leaving out 'cat' on purpose + cudf::test::strings_column_wrapper vocabulary( {"ate", "chased", "cheese", "dog", "fox", "jumped", "mouse", "mousé", "over", "the"}); auto vocab = nvtext::load_vocabulary(cudf::strings_column_view(vocabulary)); std::vector h_strings( 4, "the fox jumped chased the dog cheese mouse at the over there dog mouse cat plus the horse " - "jumped over the mouse house with the dog"); + "jumped over the mousé house with the dog "); cudf::test::strings_column_wrapper input(h_strings.begin(), h_strings.end()); auto input_view = cudf::strings_column_view(input); auto delimiter = cudf::string_scalar(" "); @@ -262,10 +262,10 @@ TEST_F(TextTokenizeTest, VocabularyLongStrings) using LCW = cudf::test::lists_column_wrapper; // clang-format off - LCW expected({LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 6, -1, -1, 9, 3}, - LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 6, -1, -1, 9, 3}, - LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 6, -1, -1, 9, 3}, - LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 6, -1, -1, 9, 3}}); + LCW expected({LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 7, -1, -1, 9, 3}, + LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 7, -1, -1, 9, 3}, + LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 7, -1, -1, 9, 3}, + LCW{ 9, 4, 5, 1, 9, 3, 2, 6, -1, 9, 8, -1, 3, 6, -1, -1, 9, -1, 5, 8, 9, 7, -1, -1, 9, 3}}); // clang-format on CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); From 8106a0c3d2050786f42152a280bd9315b897379e Mon Sep 17 00:00:00 2001 From: GALI PREM SAGAR Date: Tue, 14 Nov 2023 16:03:54 -0600 Subject: [PATCH 10/44] Cleanup remaining usages of dask dependencies (#14407) This PR switches remaining usages of `dask` dependencies to use `rapids-dask-dependency` Authors: - GALI PREM SAGAR (https://github.com/galipremsagar) Approvers: - Bradley Dice (https://github.com/bdice) - Jake Awe (https://github.com/AyodeAwe) - Vyas Ramasubramani (https://github.com/vyasr) URL: https://github.com/rapidsai/cudf/pull/14407 --- conda/recipes/custreamz/meta.yaml | 4 +--- conda/recipes/dask-cudf/run_test.sh | 36 ----------------------------- 2 files changed, 1 insertion(+), 39 deletions(-) delete mode 100644 conda/recipes/dask-cudf/run_test.sh diff --git a/conda/recipes/custreamz/meta.yaml b/conda/recipes/custreamz/meta.yaml index fb6efabffd4..b8c5918ea60 100644 --- a/conda/recipes/custreamz/meta.yaml +++ b/conda/recipes/custreamz/meta.yaml @@ -45,9 +45,7 @@ requirements: - streamz - cudf ={{ version }} - cudf_kafka ={{ version }} - - dask >=2023.9.2 - - dask-core >=2023.9.2 - - distributed >=2023.9.2 + - rapids-dask-dependency ={{ version }} - python-confluent-kafka >=1.9.0,<1.10.0a0 - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} diff --git a/conda/recipes/dask-cudf/run_test.sh b/conda/recipes/dask-cudf/run_test.sh deleted file mode 100644 index e7238d00f2b..00000000000 --- a/conda/recipes/dask-cudf/run_test.sh +++ /dev/null @@ -1,36 +0,0 @@ -#!/bin/bash -# Copyright (c) 2020-2023, NVIDIA CORPORATION. - -set -e - -# Logger function for build status output -function logger() { - echo -e "\n>>>> $@\n" -} - -# Importing cudf on arm64 CPU only nodes is currently not working due to a -# difference in reported gpu devices between arm64 and amd64 -ARCH=$(arch) - -if [ "${ARCH}" = "aarch64" ]; then - logger "Skipping tests on arm64" - exit 0 -fi - -# Dask & Distributed option to install main(nightly) or `conda-forge` packages. -export INSTALL_DASK_MAIN=1 - -# Dask version to install when `INSTALL_DASK_MAIN=0` -export DASK_STABLE_VERSION="2023.9.2" - -# Install the conda-forge or nightly version of dask and distributed -if [[ "${INSTALL_DASK_MAIN}" == 1 ]]; then - rapids-logger "rapids-mamba-retry install -c dask/label/dev 'dask/label/dev::dask' 'dask/label/dev::distributed'" - rapids-mamba-retry install -c dask/label/dev "dask/label/dev::dask" "dask/label/dev::distributed" -else - rapids-logger "rapids-mamba-retry install conda-forge::dask=={$DASK_STABLE_VERSION} conda-forge::distributed=={$DASK_STABLE_VERSION} conda-forge::dask-core=={$DASK_STABLE_VERSION} --force-reinstall" - rapids-mamba-retry install conda-forge::dask=={$DASK_STABLE_VERSION} conda-forge::distributed=={$DASK_STABLE_VERSION} conda-forge::dask-core=={$DASK_STABLE_VERSION} --force-reinstall -fi - -logger "python -c 'import dask_cudf'" -python -c "import dask_cudf" From 27b052d01ebdfd3690b90588971817423614acc0 Mon Sep 17 00:00:00 2001 From: shrshi Date: Tue, 14 Nov 2023 14:39:07 -0800 Subject: [PATCH 11/44] Added streams to CSV reader and writer api (#14340) This PR contributes to https://github.com/rapidsai/cudf/issues/13744. -Added stream parameters to public APIs `cudf::io::read_csv` `cudf::io::write_csv` -Added stream gtests Authors: - https://github.com/shrshi - Karthikeyan (https://github.com/karthikeyann) Approvers: - Karthikeyan (https://github.com/karthikeyann) - Vukasin Milovanovic (https://github.com/vuule) - Yunsong Wang (https://github.com/PointKernel) URL: https://github.com/rapidsai/cudf/pull/14340 --- cpp/include/cudf/io/csv.hpp | 4 + cpp/include/cudf/io/detail/csv.hpp | 1 - cpp/include/cudf_test/column_wrapper.hpp | 16 ++-- cpp/src/io/csv/writer_impl.cu | 38 +++++---- cpp/src/io/functions.cpp | 12 ++- cpp/tests/CMakeLists.txt | 1 + cpp/tests/streams/io/csv_test.cpp | 102 +++++++++++++++++++++++ 7 files changed, 150 insertions(+), 24 deletions(-) create mode 100644 cpp/tests/streams/io/csv_test.cpp diff --git a/cpp/include/cudf/io/csv.hpp b/cpp/include/cudf/io/csv.hpp index ac885c54356..435583e805d 100644 --- a/cpp/include/cudf/io/csv.hpp +++ b/cpp/include/cudf/io/csv.hpp @@ -1307,6 +1307,7 @@ class csv_reader_options_builder { * @endcode * * @param options Settings for controlling reading behavior + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate device memory of the table in the returned * table_with_metadata * @@ -1314,6 +1315,7 @@ class csv_reader_options_builder { */ table_with_metadata read_csv( csv_reader_options options, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group @@ -1715,9 +1717,11 @@ class csv_writer_options_builder { * @endcode * * @param options Settings for controlling writing behavior + * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource to use for device memory allocation */ void write_csv(csv_writer_options const& options, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group diff --git a/cpp/include/cudf/io/detail/csv.hpp b/cpp/include/cudf/io/detail/csv.hpp index 9fdc7a47fb9..40ddcf385b0 100644 --- a/cpp/include/cudf/io/detail/csv.hpp +++ b/cpp/include/cudf/io/detail/csv.hpp @@ -17,7 +17,6 @@ #pragma once #include -#include #include diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp index e94dfea9dcf..b9f2e0d9868 100644 --- a/cpp/include/cudf_test/column_wrapper.hpp +++ b/cpp/include/cudf_test/column_wrapper.hpp @@ -803,7 +803,8 @@ class strings_column_wrapper : public detail::column_wrapper { offsets, cudf::test::get_default_stream(), rmm::mr::get_current_device_resource()); auto d_bitmask = cudf::detail::make_device_uvector_sync( null_mask, cudf::test::get_default_stream(), rmm::mr::get_current_device_resource()); - wrapped = cudf::make_strings_column(d_chars, d_offsets, d_bitmask, null_count); + wrapped = cudf::make_strings_column( + d_chars, d_offsets, d_bitmask, null_count, cudf::test::get_default_stream()); } /** @@ -1846,7 +1847,8 @@ class structs_column_wrapper : public detail::column_wrapper { child_column_wrappers.end(), std::back_inserter(child_columns), [&](auto const& column_wrapper) { - return std::make_unique(column_wrapper.get()); + return std::make_unique(column_wrapper.get(), + cudf::test::get_default_stream()); }); init(std::move(child_columns), validity); } @@ -1882,7 +1884,8 @@ class structs_column_wrapper : public detail::column_wrapper { child_column_wrappers.end(), std::back_inserter(child_columns), [&](auto const& column_wrapper) { - return std::make_unique(column_wrapper.get()); + return std::make_unique(column_wrapper.get(), + cudf::test::get_default_stream()); }); init(std::move(child_columns), validity_iter); } @@ -1906,8 +1909,11 @@ class structs_column_wrapper : public detail::column_wrapper { return cudf::test::detail::make_null_mask(validity.begin(), validity.end()); }(); - wrapped = cudf::make_structs_column( - num_rows, std::move(child_columns), null_count, std::move(null_mask)); + wrapped = cudf::make_structs_column(num_rows, + std::move(child_columns), + null_count, + std::move(null_mask), + cudf::test::get_default_stream()); } template diff --git a/cpp/src/io/csv/writer_impl.cu b/cpp/src/io/csv/writer_impl.cu index 8c586306ad5..6e9c634804c 100644 --- a/cpp/src/io/csv/writer_impl.cu +++ b/cpp/src/io/csv/writer_impl.cu @@ -146,6 +146,12 @@ struct column_to_strings_fn { { } + ~column_to_strings_fn() = default; + column_to_strings_fn(column_to_strings_fn const&) = delete; + column_to_strings_fn& operator=(column_to_strings_fn const&) = delete; + column_to_strings_fn(column_to_strings_fn&&) = delete; + column_to_strings_fn& operator=(column_to_strings_fn&&) = delete; + // Note: `null` replacement with `na_rep` deferred to `concatenate()` // instead of column-wise; might be faster // @@ -160,8 +166,9 @@ struct column_to_strings_fn { std::enable_if_t, std::unique_ptr> operator()( column_view const& column) const { - return cudf::strings::detail::from_booleans( - column, options_.get_true_value(), options_.get_false_value(), stream_, mr_); + string_scalar true_string{options_.get_true_value(), true, stream_}; + string_scalar false_string{options_.get_false_value(), true, stream_}; + return cudf::strings::detail::from_booleans(column, true_string, false_string, stream_, mr_); } // strings: @@ -367,10 +374,10 @@ void write_chunked(data_sink* out_sink, CUDF_EXPECTS(str_column_view.size() > 0, "Unexpected empty strings column."); - cudf::string_scalar newline{options.get_line_terminator()}; + cudf::string_scalar newline{options.get_line_terminator(), true, stream}; auto p_str_col_w_nl = cudf::strings::detail::join_strings(str_column_view, newline, - string_scalar("", false), + string_scalar{"", false, stream}, stream, rmm::mr::get_current_device_resource()); strings_column_view strings_column{p_str_col_w_nl->view()}; @@ -455,12 +462,14 @@ void write_csv(data_sink* out_sink, // populate vector of string-converted columns: // - std::transform(sub_view.begin(), - sub_view.end(), - std::back_inserter(str_column_vec), - [converter](auto const& current_col) { - return cudf::type_dispatcher(current_col.type(), converter, current_col); - }); + std::transform( + sub_view.begin(), + sub_view.end(), + std::back_inserter(str_column_vec), + [&converter = std::as_const(converter)](auto const& current_col) { + return cudf::type_dispatcher( + current_col.type(), converter, current_col); + }); // create string table view from str_column_vec: // @@ -470,18 +479,19 @@ void write_csv(data_sink* out_sink, // concatenate columns in each row into one big string column // (using null representation and delimiter): // - std::string delimiter_str{options.get_inter_column_delimiter()}; auto str_concat_col = [&] { + cudf::string_scalar delimiter_str{ + std::string{options.get_inter_column_delimiter()}, true, stream}; + cudf::string_scalar options_narep{options.get_na_rep(), true, stream}; if (str_table_view.num_columns() > 1) return cudf::strings::detail::concatenate(str_table_view, delimiter_str, - options.get_na_rep(), + options_narep, strings::separator_on_nulls::YES, stream, rmm::mr::get_current_device_resource()); - cudf::string_scalar narep{options.get_na_rep()}; return cudf::strings::detail::replace_nulls( - str_table_view.column(0), narep, stream, rmm::mr::get_current_device_resource()); + str_table_view.column(0), options_narep, stream, rmm::mr::get_current_device_resource()); }(); write_chunked(out_sink, str_concat_col->view(), options, stream, mr); diff --git a/cpp/src/io/functions.cpp b/cpp/src/io/functions.cpp index 00d56008611..964e40e36cd 100644 --- a/cpp/src/io/functions.cpp +++ b/cpp/src/io/functions.cpp @@ -230,7 +230,9 @@ void write_json(json_writer_options const& options, mr); } -table_with_metadata read_csv(csv_reader_options options, rmm::mr::device_memory_resource* mr) +table_with_metadata read_csv(csv_reader_options options, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); @@ -245,12 +247,14 @@ table_with_metadata read_csv(csv_reader_options options, rmm::mr::device_memory_ return cudf::io::detail::csv::read_csv( // std::move(datasources[0]), options, - cudf::get_default_stream(), + stream, mr); } // Freeform API wraps the detail writer class API -void write_csv(csv_writer_options const& options, rmm::mr::device_memory_resource* mr) +void write_csv(csv_writer_options const& options, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { using namespace cudf::io::detail; @@ -262,7 +266,7 @@ void write_csv(csv_writer_options const& options, rmm::mr::device_memory_resourc options.get_table(), options.get_names(), options, - cudf::get_default_stream(), + stream, mr); } diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 7b628649051..1be8566fb0f 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -635,6 +635,7 @@ ConfigureTest( ConfigureTest(STREAM_BINARYOP_TEST streams/binaryop_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_CONCATENATE_TEST streams/concatenate_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_COPYING_TEST streams/copying_test.cpp STREAM_MODE testing) +ConfigureTest(STREAM_CSVIO_TEST streams/io/csv_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_DICTIONARY_TEST streams/dictionary_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_FILLING_TEST streams/filling_test.cpp STREAM_MODE testing) ConfigureTest(STREAM_GROUPBY_TEST streams/groupby_test.cpp STREAM_MODE testing) diff --git a/cpp/tests/streams/io/csv_test.cpp b/cpp/tests/streams/io/csv_test.cpp new file mode 100644 index 00000000000..88514fa412c --- /dev/null +++ b/cpp/tests/streams/io/csv_test.cpp @@ -0,0 +1,102 @@ +/* + * Copyright (c) 2023, NVIDIA 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 +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include + +auto const temp_env = static_cast( + ::testing::AddGlobalTestEnvironment(new cudf::test::TempDirTestEnvironment)); + +class CSVTest : public cudf::test::BaseFixture {}; + +TEST_F(CSVTest, CSVWriter) +{ + constexpr auto num_rows = 10; + + std::vector zeros(num_rows, 0); + std::vector ones(num_rows, 1); + auto col6_data = cudf::detail::make_counting_transform_iterator(0, [&](auto i) { + return numeric::decimal128{ones[i], numeric::scale_type{12}}; + }); + auto col7_data = cudf::detail::make_counting_transform_iterator(0, [&](auto i) { + return numeric::decimal128{ones[i], numeric::scale_type{-12}}; + }); + + cudf::test::fixed_width_column_wrapper col0(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col1(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col2(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col3(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col4(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col5(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col6(col6_data, col6_data + num_rows); + cudf::test::fixed_width_column_wrapper col7(col7_data, col7_data + num_rows); + + std::vector col8_data(num_rows, "rapids"); + cudf::test::strings_column_wrapper col8(col8_data.begin(), col8_data.end()); + + cudf::table_view tab({col0, col1, col2, col3, col4, col5, col6, col7, col8}); + + auto const filepath = temp_env->get_temp_dir() + "multicolumn.csv"; + auto w_options = cudf::io::csv_writer_options::builder(cudf::io::sink_info{filepath}, tab) + .include_header(false) + .inter_column_delimiter(','); + cudf::io::write_csv(w_options.build(), cudf::test::get_default_stream()); +} + +TEST_F(CSVTest, CSVReader) +{ + constexpr auto num_rows = 10; + + std::vector zeros(num_rows, 0); + std::vector ones(num_rows, 1); + auto col6_data = cudf::detail::make_counting_transform_iterator(0, [&](auto i) { + return numeric::decimal128{ones[i], numeric::scale_type{12}}; + }); + auto col7_data = cudf::detail::make_counting_transform_iterator(0, [&](auto i) { + return numeric::decimal128{ones[i], numeric::scale_type{-12}}; + }); + + cudf::test::fixed_width_column_wrapper col0(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col1(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col2(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col3(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col4(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col5(zeros.begin(), zeros.end()); + cudf::test::fixed_width_column_wrapper col6(col6_data, col6_data + num_rows); + cudf::test::fixed_width_column_wrapper col7(col7_data, col7_data + num_rows); + + std::vector col8_data(num_rows, "rapids"); + cudf::test::strings_column_wrapper col8(col8_data.begin(), col8_data.end()); + + cudf::table_view tab({col0, col1, col2, col3, col4, col5, col6, col7, col8}); + + auto const filepath = temp_env->get_temp_dir() + "multicolumn.csv"; + auto w_options = cudf::io::csv_writer_options::builder(cudf::io::sink_info{filepath}, tab) + .include_header(false) + .inter_column_delimiter(','); + cudf::io::write_csv(w_options.build(), cudf::test::get_default_stream()); +} From 330d389b26a05676d9f079503a3d96b571762337 Mon Sep 17 00:00:00 2001 From: Robert Maynard Date: Tue, 14 Nov 2023 17:56:55 -0500 Subject: [PATCH 12/44] Ensure nvbench initializes nvml context when built statically (#14411) Port https://github.com/NVIDIA/nvbench/pull/148 to cudf so that nvbench benchmarks work now that we always use a static version of nvbench. Authors: - Robert Maynard (https://github.com/robertmaynard) Approvers: - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14411 --- cpp/cmake/thirdparty/patches/nvbench_override.json | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/cpp/cmake/thirdparty/patches/nvbench_override.json b/cpp/cmake/thirdparty/patches/nvbench_override.json index ad9b19c29c1..f85bdb9486c 100644 --- a/cpp/cmake/thirdparty/patches/nvbench_override.json +++ b/cpp/cmake/thirdparty/patches/nvbench_override.json @@ -7,6 +7,11 @@ "file" : "${current_json_dir}/nvbench_global_setup.diff", "issue" : "Fix add support for global setup to initialize RMM in nvbench [https://github.com/NVIDIA/nvbench/pull/123]", "fixed_in" : "" + }, + { + "file" : "nvbench/nvml_with_static_builds.diff", + "issue" : "Add support for nvml with static nvbench [https://github.com/NVIDIA/nvbench/pull/148]", + "fixed_in" : "" } ] } From 8a0a08f34ff804a7329ea640aa1e0a9b188d2162 Mon Sep 17 00:00:00 2001 From: Matthew Roeschke <10647082+mroeschke@users.noreply.github.com> Date: Tue, 14 Nov 2023 17:55:16 -1000 Subject: [PATCH 13/44] Fix as_column(pd.Timestamp/Timedelta, length=) not respecting length (#14390) Noticed this while trying to clean up `as_column` Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14390 --- python/cudf/cudf/core/column/column.py | 5 ++++- python/cudf/cudf/tests/test_column.py | 13 ++++++++----- 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/python/cudf/cudf/core/column/column.py b/python/cudf/cudf/core/column/column.py index a5e99abd79e..b4f65693d85 100644 --- a/python/cudf/cudf/core/column/column.py +++ b/python/cudf/cudf/core/column/column.py @@ -2102,7 +2102,10 @@ def as_column( elif isinstance(arbitrary, (pd.Timestamp, pd.Timedelta)): # This will always treat NaTs as nulls since it's not technically a # discrete value like NaN - data = as_column(pa.array(pd.Series([arbitrary]), from_pandas=True)) + length = length or 1 + data = as_column( + pa.array(pd.Series([arbitrary] * length), from_pandas=True) + ) if dtype is not None: data = data.astype(dtype) diff --git a/python/cudf/cudf/tests/test_column.py b/python/cudf/cudf/tests/test_column.py index db0446d506c..0546638f388 100644 --- a/python/cudf/cudf/tests/test_column.py +++ b/python/cudf/cudf/tests/test_column.py @@ -193,12 +193,15 @@ def test_column_mixed_dtype(data, error): @pytest.mark.parametrize("nan_as_null", [True, False]) -def test_as_column_scalar_with_nan(nan_as_null): - size = 10 - scalar = np.nan - +@pytest.mark.parametrize( + "scalar", + [np.nan, pd.Timedelta(days=1), pd.Timestamp(2020, 1, 1)], + ids=repr, +) +@pytest.mark.parametrize("size", [1, 10]) +def test_as_column_scalar_with_nan(nan_as_null, scalar, size): expected = ( - cudf.Series([np.nan] * size, nan_as_null=nan_as_null) + cudf.Series([scalar] * size, nan_as_null=nan_as_null) .dropna() .to_numpy() ) From ab2248ea8e693143823d02bb8b806c65bfc3bf32 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Tue, 14 Nov 2023 23:30:51 -0800 Subject: [PATCH 14/44] Fix and disable encoding for nanosecond statistics in ORC writer (#14367) Issue https://github.com/rapidsai/cudf/issues/14325 Use uint when reading/writing nano stats because nanoseconds have int32 encoding (different from both unit32 and sint32, _obviously_), which does not use zigzag. sint32 uses zigzag, and unit32 does not allow negative numbers, so we can use uint since we'll never have negative nanoseconds. Also disabled the nanoseconds because it should only be written after ORC-135; we don't write the version so readers get confused if nanoseconds are there. Planning to re-enable once we start writing the version. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/14367 --- cpp/include/cudf/io/orc_metadata.hpp | 8 +++---- cpp/src/io/orc/orc.cpp | 13 +++++++++++ cpp/src/io/orc/orc.hpp | 6 +++++ cpp/src/io/orc/stats_enc.cu | 35 +++++++++++++++++++++------- cpp/tests/io/orc_test.cpp | 16 +++++++++---- 5 files changed, 61 insertions(+), 17 deletions(-) diff --git a/cpp/include/cudf/io/orc_metadata.hpp b/cpp/include/cudf/io/orc_metadata.hpp index 82d59803c25..9531a012e49 100644 --- a/cpp/include/cudf/io/orc_metadata.hpp +++ b/cpp/include/cudf/io/orc_metadata.hpp @@ -141,10 +141,10 @@ using binary_statistics = sum_statistics; * the UNIX epoch. The `minimum_utc` and `maximum_utc` are the same values adjusted to UTC. */ struct timestamp_statistics : minmax_statistics { - std::optional minimum_utc; ///< minimum in milliseconds - std::optional maximum_utc; ///< maximum in milliseconds - std::optional minimum_nanos; ///< nanoseconds part of the minimum - std::optional maximum_nanos; ///< nanoseconds part of the maximum + std::optional minimum_utc; ///< minimum in milliseconds + std::optional maximum_utc; ///< maximum in milliseconds + std::optional minimum_nanos; ///< nanoseconds part of the minimum + std::optional maximum_nanos; ///< nanoseconds part of the maximum }; namespace orc { diff --git a/cpp/src/io/orc/orc.cpp b/cpp/src/io/orc/orc.cpp index bc399b75ef9..ee5fa4e8b5a 100644 --- a/cpp/src/io/orc/orc.cpp +++ b/cpp/src/io/orc/orc.cpp @@ -182,6 +182,19 @@ void ProtobufReader::read(timestamp_statistics& s, size_t maxlen) field_reader(5, s.minimum_nanos), field_reader(6, s.maximum_nanos)); function_builder(s, maxlen, op); + + // Adjust nanoseconds because they are encoded as (value + 1) + // Range [1, 1000'000] is translated here to [0, 999'999] + if (s.minimum_nanos.has_value()) { + auto& min_nanos = s.minimum_nanos.value(); + CUDF_EXPECTS(min_nanos >= 1 and min_nanos <= 1000'000, "Invalid minimum nanoseconds"); + --min_nanos; + } + if (s.maximum_nanos.has_value()) { + auto& max_nanos = s.maximum_nanos.value(); + CUDF_EXPECTS(max_nanos >= 1 and max_nanos <= 1000'000, "Invalid maximum nanoseconds"); + --max_nanos; + } } void ProtobufReader::read(column_statistics& s, size_t maxlen) diff --git a/cpp/src/io/orc/orc.hpp b/cpp/src/io/orc/orc.hpp index 6f65e384d2d..783ed4206b6 100644 --- a/cpp/src/io/orc/orc.hpp +++ b/cpp/src/io/orc/orc.hpp @@ -41,6 +41,12 @@ static constexpr uint32_t block_header_size = 3; // Seconds from January 1st, 1970 to January 1st, 2015 static constexpr int64_t orc_utc_epoch = 1420070400; +// Used for the nanosecond remainder in timestamp statistics when the actual nanoseconds of min/max +// are not included. As the timestamp statistics are stored as milliseconds + nanosecond remainder, +// the maximum nanosecond remainder is 999,999 (nanoseconds in a millisecond - 1). +static constexpr int32_t DEFAULT_MIN_NANOS = 0; +static constexpr int32_t DEFAULT_MAX_NANOS = 999'999; + struct PostScript { uint64_t footerLength = 0; // the length of the footer section in bytes CompressionKind compression = NONE; // the kind of generic compression used diff --git a/cpp/src/io/orc/stats_enc.cu b/cpp/src/io/orc/stats_enc.cu index 479a2dfada3..429fd5b929d 100644 --- a/cpp/src/io/orc/stats_enc.cu +++ b/cpp/src/io/orc/stats_enc.cu @@ -27,6 +27,10 @@ namespace cudf::io::orc::gpu { using strings::detail::fixed_point_string_size; +// Nanosecond statistics should not be enabled until the spec version is set correctly in the output +// files. See https://github.com/rapidsai/cudf/issues/14325 for more details +constexpr bool enable_nanosecond_statistics = false; + constexpr unsigned int init_threads_per_group = 32; constexpr unsigned int init_groups_per_block = 4; constexpr unsigned int init_threads_per_block = init_threads_per_group * init_groups_per_block; @@ -96,8 +100,10 @@ __global__ void __launch_bounds__(block_size, 1) stats_len = pb_fldlen_common + pb_fld_hdrlen + 2 * (pb_fld_hdrlen + pb_fldlen_int64); break; case dtype_timestamp64: - stats_len = pb_fldlen_common + pb_fld_hdrlen + 4 * (pb_fld_hdrlen + pb_fldlen_int64) + - 2 * (pb_fld_hdrlen + pb_fldlen_int32); + stats_len = pb_fldlen_common + pb_fld_hdrlen + 4 * (pb_fld_hdrlen + pb_fldlen_int64); + if constexpr (enable_nanosecond_statistics) { + stats_len += 2 * (pb_fld_hdrlen + pb_fldlen_int32); + } break; case dtype_float32: case dtype_float64: @@ -405,7 +411,8 @@ __global__ void __launch_bounds__(encode_threads_per_block) // optional sint64 minimumUtc = 3; // min,max values saved as milliseconds since UNIX epoch // optional sint64 maximumUtc = 4; // optional int32 minimumNanos = 5; // lower 6 TS digits for min/max to achieve nanosecond - // precision optional int32 maximumNanos = 6; + // precision + // optional int32 maximumNanos = 6; // } if (s->chunk.has_minmax) { cur[0] = 9 * 8 + ProtofType::FIXEDLEN; @@ -416,12 +423,22 @@ __global__ void __launch_bounds__(encode_threads_per_block) split_nanosecond_timestamp(s->chunk.max_value.i_val); // minimum/maximum are the same as minimumUtc/maximumUtc as we always write files in UTC - cur = pb_put_int(cur, 1, min_ms); // minimum - cur = pb_put_int(cur, 2, max_ms); // maximum - cur = pb_put_int(cur, 3, min_ms); // minimumUtc - cur = pb_put_int(cur, 4, max_ms); // maximumUtc - cur = pb_put_int(cur, 5, min_ns_remainder); // minimumNanos - cur = pb_put_int(cur, 6, max_ns_remainder); // maximumNanos + cur = pb_put_int(cur, 1, min_ms); // minimum + cur = pb_put_int(cur, 2, max_ms); // maximum + cur = pb_put_int(cur, 3, min_ms); // minimumUtc + cur = pb_put_int(cur, 4, max_ms); // maximumUtc + + if constexpr (enable_nanosecond_statistics) { + if (min_ns_remainder != DEFAULT_MIN_NANOS) { + // using uint because positive values are not zigzag encoded + cur = pb_put_uint(cur, 5, min_ns_remainder + 1); // minimumNanos + } + if (max_ns_remainder != DEFAULT_MAX_NANOS) { + // using uint because positive values are not zigzag encoded + cur = pb_put_uint(cur, 6, max_ns_remainder + 1); // maximumNanos + } + } + fld_start[1] = cur - (fld_start + 2); } break; diff --git a/cpp/tests/io/orc_test.cpp b/cpp/tests/io/orc_test.cpp index 234716749ff..dca3886db14 100644 --- a/cpp/tests/io/orc_test.cpp +++ b/cpp/tests/io/orc_test.cpp @@ -1054,8 +1054,12 @@ TEST_F(OrcStatisticsTest, Basic) EXPECT_EQ(*ts4.maximum, 3); EXPECT_EQ(*ts4.minimum_utc, -4); EXPECT_EQ(*ts4.maximum_utc, 3); - EXPECT_EQ(*ts4.minimum_nanos, 999994); - EXPECT_EQ(*ts4.maximum_nanos, 6); + // nanosecond precision can't be included until we write a writer version that includes ORC-135 + // see https://github.com/rapidsai/cudf/issues/14325 + // EXPECT_EQ(*ts4.minimum_nanos, 999994); + EXPECT_FALSE(ts4.minimum_nanos.has_value()); + // EXPECT_EQ(*ts4.maximum_nanos, 6); + EXPECT_FALSE(ts4.maximum_nanos.has_value()); auto& s5 = stats[5]; EXPECT_EQ(*s5.number_of_values, 4ul); @@ -1065,8 +1069,12 @@ TEST_F(OrcStatisticsTest, Basic) EXPECT_EQ(*ts5.maximum, 3000); EXPECT_EQ(*ts5.minimum_utc, -3001); EXPECT_EQ(*ts5.maximum_utc, 3000); - EXPECT_EQ(*ts5.minimum_nanos, 994000); - EXPECT_EQ(*ts5.maximum_nanos, 6000); + // nanosecond precision can't be included until we write a writer version that includes ORC-135 + // see https://github.com/rapidsai/cudf/issues/14325 + // EXPECT_EQ(*ts5.minimum_nanos, 994000); + EXPECT_FALSE(ts5.minimum_nanos.has_value()); + // EXPECT_EQ(*ts5.maximum_nanos, 6000); + EXPECT_FALSE(ts5.maximum_nanos.has_value()); auto& s6 = stats[6]; EXPECT_EQ(*s6.number_of_values, 4ul); From 8deb3dd7573000e7d87f18a9e2bbe39cf2932e10 Mon Sep 17 00:00:00 2001 From: GALI PREM SAGAR Date: Wed, 15 Nov 2023 07:56:37 -0600 Subject: [PATCH 15/44] Raise error in `reindex` when `index` is not unique (#14400) Fixes: #14398 This PR raises an error in `reindex` API when reindexing is performed on a non-unique index column. Authors: - GALI PREM SAGAR (https://github.com/galipremsagar) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14400 --- python/cudf/cudf/core/indexed_frame.py | 4 ++++ python/cudf/cudf/tests/test_dataframe.py | 12 ++++++++++++ python/cudf/cudf/tests/test_series.py | 12 ++++++++++++ python/dask_cudf/dask_cudf/backends.py | 13 ++++--------- 4 files changed, 32 insertions(+), 9 deletions(-) diff --git a/python/cudf/cudf/core/indexed_frame.py b/python/cudf/cudf/core/indexed_frame.py index 376bef6d0b2..4211a8c24bf 100644 --- a/python/cudf/cudf/core/indexed_frame.py +++ b/python/cudf/cudf/core/indexed_frame.py @@ -2607,6 +2607,10 @@ def _reindex( df = self if index is not None: + if not df._index.is_unique: + raise ValueError( + "cannot reindex on an axis with duplicate labels" + ) index = cudf.core.index.as_index( index, name=getattr(index, "name", self._index.name) ) diff --git a/python/cudf/cudf/tests/test_dataframe.py b/python/cudf/cudf/tests/test_dataframe.py index d44cf594e8b..5677f97408a 100644 --- a/python/cudf/cudf/tests/test_dataframe.py +++ b/python/cudf/cudf/tests/test_dataframe.py @@ -10723,3 +10723,15 @@ def test_dataframe_series_dot(): expected = gser @ [12, 13] assert_eq(expected, actual) + + +def test_dataframe_duplicate_index_reindex(): + gdf = cudf.DataFrame({"a": [0, 1, 2, 3]}, index=[0, 0, 1, 1]) + pdf = gdf.to_pandas() + + assert_exceptions_equal( + gdf.reindex, + pdf.reindex, + lfunc_args_and_kwargs=([10, 11, 12, 13], {}), + rfunc_args_and_kwargs=([10, 11, 12, 13], {}), + ) diff --git a/python/cudf/cudf/tests/test_series.py b/python/cudf/cudf/tests/test_series.py index 8f8f87c20e0..c15a797713f 100644 --- a/python/cudf/cudf/tests/test_series.py +++ b/python/cudf/cudf/tests/test_series.py @@ -2638,3 +2638,15 @@ def test_series_setitem_mixed_bool_dtype(): s = cudf.Series([True, False, True]) with pytest.raises(TypeError): s[0] = 10 + + +def test_series_duplicate_index_reindex(): + gs = cudf.Series([0, 1, 2, 3], index=[0, 0, 1, 1]) + ps = gs.to_pandas() + + assert_exceptions_equal( + gs.reindex, + ps.reindex, + lfunc_args_and_kwargs=([10, 11, 12, 13], {}), + rfunc_args_and_kwargs=([10, 11, 12, 13], {}), + ) diff --git a/python/dask_cudf/dask_cudf/backends.py b/python/dask_cudf/dask_cudf/backends.py index b0da82eaeee..387643587d1 100644 --- a/python/dask_cudf/dask_cudf/backends.py +++ b/python/dask_cudf/dask_cudf/backends.py @@ -427,17 +427,12 @@ def union_categoricals_cudf( ) -@_dask_cudf_nvtx_annotate -def safe_hash(frame): - return cudf.Series(frame.hash_values(), index=frame.index) - - @hash_object_dispatch.register((cudf.DataFrame, cudf.Series)) @_dask_cudf_nvtx_annotate def hash_object_cudf(frame, index=True): if index: - return safe_hash(frame.reset_index()) - return safe_hash(frame) + frame = frame.reset_index() + return frame.hash_values() @hash_object_dispatch.register(cudf.BaseIndex) @@ -445,10 +440,10 @@ def hash_object_cudf(frame, index=True): def hash_object_cudf_index(ind, index=None): if isinstance(ind, cudf.MultiIndex): - return safe_hash(ind.to_frame(index=False)) + return ind.to_frame(index=False).hash_values() col = cudf.core.column.as_column(ind) - return safe_hash(cudf.Series(col)) + return cudf.Series(col).hash_values() @group_split_dispatch.register((cudf.Series, cudf.DataFrame)) From 9e7f8a5fdd03d6a24630687621d0ee14c2db26d7 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Wed, 15 Nov 2023 15:27:28 -0800 Subject: [PATCH 16/44] Fix dask dependency in custreamz (#14420) #14407 added a dask dependency to custreamz, but it added too tight of a pinning by requiring the exact same version. This is not valid because rapids-dask-dependency won't release a new version corresponding to each new cudf release, so pinning to the exact same version up to the alpha creates an unsatisfiable constraint. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Ray Douglass (https://github.com/raydouglass) - Bradley Dice (https://github.com/bdice) - GALI PREM SAGAR (https://github.com/galipremsagar) --- conda/recipes/custreamz/meta.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/conda/recipes/custreamz/meta.yaml b/conda/recipes/custreamz/meta.yaml index b8c5918ea60..755394e3936 100644 --- a/conda/recipes/custreamz/meta.yaml +++ b/conda/recipes/custreamz/meta.yaml @@ -45,7 +45,7 @@ requirements: - streamz - cudf ={{ version }} - cudf_kafka ={{ version }} - - rapids-dask-dependency ={{ version }} + - rapids-dask-dependency ={{ minor_version }} - python-confluent-kafka >=1.9.0,<1.10.0a0 - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} From e4e69757b340ce45e8ceca53047f079c8b3eb648 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Wed, 15 Nov 2023 16:31:03 -0800 Subject: [PATCH 17/44] Update cudf_kafka_version. --- python/cudf_kafka/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/python/cudf_kafka/CMakeLists.txt b/python/cudf_kafka/CMakeLists.txt index d55c3fdc076..1e21c873585 100644 --- a/python/cudf_kafka/CMakeLists.txt +++ b/python/cudf_kafka/CMakeLists.txt @@ -14,7 +14,7 @@ cmake_minimum_required(VERSION 3.26.4 FATAL_ERROR) -set(cudf_kafka_version 23.12.00) +set(cudf_kafka_version 24.02.00) include(../../fetch_rapids.cmake) From f9c586d48aa2a879b2267318088d3cc38f398662 Mon Sep 17 00:00:00 2001 From: Ferdinand Xu Date: Thu, 16 Nov 2023 10:14:19 +0800 Subject: [PATCH 18/44] Support java AST String literal with desired encoding (#14402) Authors: - Ferdinand Xu (https://github.com/winningsix) - Jason Lowe (https://github.com/jlowe) Approvers: - Robert (Bobby) Evans (https://github.com/revans2) - Jason Lowe (https://github.com/jlowe) URL: https://github.com/rapidsai/cudf/pull/14402 --- java/src/main/java/ai/rapids/cudf/ast/Literal.java | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/java/src/main/java/ai/rapids/cudf/ast/Literal.java b/java/src/main/java/ai/rapids/cudf/ast/Literal.java index 427dd286b0c..4e1e886c282 100644 --- a/java/src/main/java/ai/rapids/cudf/ast/Literal.java +++ b/java/src/main/java/ai/rapids/cudf/ast/Literal.java @@ -20,6 +20,7 @@ import java.nio.ByteBuffer; import java.nio.ByteOrder; +import java.nio.charset.StandardCharsets; /** A literal value in an AST expression. */ public final class Literal extends AstExpression { @@ -205,7 +206,14 @@ public static Literal ofString(String value) { if (value == null) { return ofNull(DType.STRING); } - byte[] stringBytes = value.getBytes(); + return ofUTF8String(value.getBytes(StandardCharsets.UTF_8)); + } + + /** Construct a string literal directly with byte array to skip transcoding. */ + public static Literal ofUTF8String(byte[] stringBytes) { + if (stringBytes == null) { + return ofNull(DType.STRING); + } byte[] serializedValue = new byte[stringBytes.length + Integer.BYTES]; ByteBuffer.wrap(serializedValue).order(ByteOrder.nativeOrder()).putInt(stringBytes.length); System.arraycopy(stringBytes, 0, serializedValue, Integer.BYTES, stringBytes.length); From afd7d189b83cbcccba783877f42bb153b5cf315e Mon Sep 17 00:00:00 2001 From: Divye Gala Date: Wed, 15 Nov 2023 23:33:28 -0500 Subject: [PATCH 19/44] Example code for blog on new row comparators (#13795) Example code using a few libcudf APIs to demonstrate nested-type usage. Authors: - Divye Gala (https://github.com/divyegala) - Karthikeyan (https://github.com/karthikeyann) - Vyas Ramasubramani (https://github.com/vyasr) - Bradley Dice (https://github.com/bdice) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Bradley Dice (https://github.com/bdice) - Karthikeyan (https://github.com/karthikeyann) - Yunsong Wang (https://github.com/PointKernel) - Nghia Truong (https://github.com/ttnghia) - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cudf/pull/13795 --- ci/release/update-version.sh | 3 +- cpp/examples/README.md | 1 + cpp/examples/basic/CMakeLists.txt | 18 +- cpp/examples/build.sh | 31 ++- cpp/examples/fetch_dependencies.cmake | 30 +++ cpp/examples/nested_types/CMakeLists.txt | 16 ++ cpp/examples/nested_types/deduplication.cpp | 209 ++++++++++++++++++++ cpp/examples/nested_types/example.json | 5 + cpp/examples/strings/CMakeLists.txt | 18 +- 9 files changed, 279 insertions(+), 52 deletions(-) create mode 100644 cpp/examples/fetch_dependencies.cmake create mode 100644 cpp/examples/nested_types/CMakeLists.txt create mode 100644 cpp/examples/nested_types/deduplication.cpp create mode 100644 cpp/examples/nested_types/example.json diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 4f1cbc47d1d..16742465c32 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -101,8 +101,7 @@ sed_runner "s/version == ${CURRENT_SHORT_TAG}/version == ${NEXT_SHORT_TAG}/g" RE sed_runner "s/cudf=${CURRENT_SHORT_TAG}/cudf=${NEXT_SHORT_TAG}/g" README.md # Libcudf examples update -sed_runner "s/CUDF_TAG branch-${CURRENT_SHORT_TAG}/CUDF_TAG branch-${NEXT_SHORT_TAG}/" cpp/examples/basic/CMakeLists.txt -sed_runner "s/CUDF_TAG branch-${CURRENT_SHORT_TAG}/CUDF_TAG branch-${NEXT_SHORT_TAG}/" cpp/examples/strings/CMakeLists.txt +sed_runner "s/CUDF_TAG branch-${CURRENT_SHORT_TAG}/CUDF_TAG branch-${NEXT_SHORT_TAG}/" cpp/examples/fetch_dependencies.cmake # CI files for FILE in .github/workflows/*.yaml; do diff --git a/cpp/examples/README.md b/cpp/examples/README.md index b2e8dd399d0..7f2b769f4a5 100644 --- a/cpp/examples/README.md +++ b/cpp/examples/README.md @@ -7,3 +7,4 @@ Current examples: - Basic: demonstrates a basic use case with libcudf and building a custom application with libcudf - Strings: demonstrates using libcudf for accessing and creating strings columns and for building custom kernels for strings +- Nested Types: demonstrates using libcudf for some operations on nested types diff --git a/cpp/examples/basic/CMakeLists.txt b/cpp/examples/basic/CMakeLists.txt index 9ff716f41e4..759a43b5627 100644 --- a/cpp/examples/basic/CMakeLists.txt +++ b/cpp/examples/basic/CMakeLists.txt @@ -8,23 +8,7 @@ project( LANGUAGES CXX CUDA ) -set(CPM_DOWNLOAD_VERSION v0.35.3) -file( - DOWNLOAD - https://github.com/cpm-cmake/CPM.cmake/releases/download/${CPM_DOWNLOAD_VERSION}/get_cpm.cmake - ${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake -) -include(${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake) - -set(CUDF_TAG branch-23.12) -CPMFindPackage( - NAME cudf GIT_REPOSITORY https://github.com/rapidsai/cudf - GIT_TAG ${CUDF_TAG} - GIT_SHALLOW - TRUE - SOURCE_SUBDIR - cpp -) +include(../fetch_dependencies.cmake) # Configure your project here add_executable(basic_example src/process_csv.cpp) diff --git a/cpp/examples/build.sh b/cpp/examples/build.sh index 7d389cd318d..001cdeec694 100755 --- a/cpp/examples/build.sh +++ b/cpp/examples/build.sh @@ -1,6 +1,6 @@ #!/bin/bash -# Copyright (c) 2021-2022, NVIDIA CORPORATION. +# Copyright (c) 2021-2023, NVIDIA CORPORATION. # libcudf examples build script @@ -14,18 +14,17 @@ LIB_BUILD_DIR=${LIB_BUILD_DIR:-$(readlink -f "${EXAMPLES_DIR}/../build")} ################################################################################ # Add individual libcudf examples build scripts down below -# Basic example -BASIC_EXAMPLE_DIR=${EXAMPLES_DIR}/basic -BASIC_EXAMPLE_BUILD_DIR=${BASIC_EXAMPLE_DIR}/build -# Configure -cmake -S ${BASIC_EXAMPLE_DIR} -B ${BASIC_EXAMPLE_BUILD_DIR} -Dcudf_ROOT="${LIB_BUILD_DIR}" -# Build -cmake --build ${BASIC_EXAMPLE_BUILD_DIR} -j${PARALLEL_LEVEL} - -# Strings example -STRINGS_EXAMPLE_DIR=${EXAMPLES_DIR}/strings -STRINGS_EXAMPLE_BUILD_DIR=${STRINGS_EXAMPLE_DIR}/build -# Configure -cmake -S ${STRINGS_EXAMPLE_DIR} -B ${STRINGS_EXAMPLE_BUILD_DIR} -Dcudf_ROOT="${LIB_BUILD_DIR}" -# Build -cmake --build ${STRINGS_EXAMPLE_BUILD_DIR} -j${PARALLEL_LEVEL} +build_example() { + example_dir=${1} + example_dir="${EXAMPLES_DIR}/${example_dir}" + build_dir="${example_dir}/build" + + # Configure + cmake -S ${example_dir} -B ${build_dir} -Dcudf_ROOT="${LIB_BUILD_DIR}" + # Build + cmake --build ${build_dir} -j${PARALLEL_LEVEL} +} + +build_example basic +build_example strings +build_example nested_types diff --git a/cpp/examples/fetch_dependencies.cmake b/cpp/examples/fetch_dependencies.cmake new file mode 100644 index 00000000000..dc86c6a9aa5 --- /dev/null +++ b/cpp/examples/fetch_dependencies.cmake @@ -0,0 +1,30 @@ +# ============================================================================= +# Copyright (c) 2023, NVIDIA 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. +# ============================================================================= +set(CPM_DOWNLOAD_VERSION v0.35.3) +file( + DOWNLOAD + https://github.com/cpm-cmake/CPM.cmake/releases/download/${CPM_DOWNLOAD_VERSION}/get_cpm.cmake + ${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake +) +include(${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake) + +set(CUDF_TAG branch-23.12) +CPMFindPackage( + NAME cudf GIT_REPOSITORY https://github.com/rapidsai/cudf + GIT_TAG ${CUDF_TAG} + GIT_SHALLOW + TRUE + SOURCE_SUBDIR + cpp +) diff --git a/cpp/examples/nested_types/CMakeLists.txt b/cpp/examples/nested_types/CMakeLists.txt new file mode 100644 index 00000000000..cb9430db237 --- /dev/null +++ b/cpp/examples/nested_types/CMakeLists.txt @@ -0,0 +1,16 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. + +cmake_minimum_required(VERSION 3.26.4) + +project( + nested_types + VERSION 0.0.1 + LANGUAGES CXX CUDA +) + +include(../fetch_dependencies.cmake) + +# Configure your project here +add_executable(deduplication deduplication.cpp) +target_link_libraries(deduplication PRIVATE cudf::cudf) +target_compile_features(deduplication PRIVATE cxx_std_17) diff --git a/cpp/examples/nested_types/deduplication.cpp b/cpp/examples/nested_types/deduplication.cpp new file mode 100644 index 00000000000..5969985cc72 --- /dev/null +++ b/cpp/examples/nested_types/deduplication.cpp @@ -0,0 +1,209 @@ +/* + * Copyright (c) 2023, NVIDIA 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 +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include + +/** + * @file deduplication.cpp + * @brief Demonstrates usage of the libcudf APIs to perform operations on nested-type tables. + * + * The algorithms chosen to be demonstrated are to showcase nested-type row operators of three + * kinds: + * 1. hashing: Used by functions `count_aggregate` and `join_count` to hash inputs of any type + * 2. equality: Used by functions `count_aggregate` and `join_count` in conjunction with hashing + * to determine equality for nested types + * 3. lexicographic: Used by function `sort_keys` to create a lexicographical order for nested-types + * so as to enable sorting + * + */ + +/** + * @brief Create memory resource for libcudf functions + * + * @param pool Whether to use a pool memory resource. + * @return Memory resource instance + */ +std::shared_ptr create_memory_resource(bool pool) +{ + auto cuda_mr = std::make_shared(); + if (pool) { return rmm::mr::make_owning_wrapper(cuda_mr); } + return cuda_mr; +} + +/** + * @brief Read JSON input from file + * + * @param filepath path to input JSON file + * @return cudf::io::table_with_metadata + */ +cudf::io::table_with_metadata read_json(std::string filepath) +{ + auto source_info = cudf::io::source_info(filepath); + auto builder = cudf::io::json_reader_options::builder(source_info).lines(true); + auto options = builder.build(); + return cudf::io::read_json(options); +} + +/** + * @brief Write JSON output to file + * + * @param input table to write + * @param metadata metadata of input table read by JSON reader + * @param filepath path to output JSON file + */ +void write_json(cudf::table_view input, cudf::io::table_metadata metadata, std::string filepath) +{ + // write the data for inspection + auto sink_info = cudf::io::sink_info(filepath); + auto builder = cudf::io::json_writer_options::builder(sink_info, input).lines(true); + builder.metadata(metadata); + auto options = builder.build(); + cudf::io::write_json(options); +} + +/** + * @brief Aggregate count of duplicate rows in nested-type column + * + * @param input table to aggregate + * @return std::unique_ptr + */ +std::unique_ptr count_aggregate(cudf::table_view input) +{ + // Get count for each key + auto keys = cudf::table_view{{input.column(0)}}; + auto val = cudf::make_numeric_column(cudf::data_type{cudf::type_id::INT32}, keys.num_rows()); + + cudf::groupby::groupby grpby_obj(keys); + std::vector requests; + requests.emplace_back(cudf::groupby::aggregation_request()); + auto agg = cudf::make_count_aggregation(); + requests[0].aggregations.push_back(std::move(agg)); + requests[0].values = *val; + auto agg_results = grpby_obj.aggregate(requests); + auto result_key = std::move(agg_results.first); + auto result_val = std::move(agg_results.second[0].results[0]); + + auto left_cols = result_key->release(); + left_cols.push_back(std::move(result_val)); + + return std::make_unique(std::move(left_cols)); +} + +/** + * @brief Join each row with its duplicate counts + * + * @param left left table + * @param right right table + * @return std::unique_ptr + */ +std::unique_ptr join_count(cudf::table_view left, cudf::table_view right) +{ + auto [left_indices, right_indices] = + cudf::inner_join(cudf::table_view{{left.column(0)}}, cudf::table_view{{right.column(0)}}); + auto new_left = cudf::gather(left, cudf::device_span{*left_indices}); + auto new_right = cudf::gather(right, cudf::device_span{*right_indices}); + + auto left_cols = new_left->release(); + auto right_cols = new_right->release(); + left_cols.push_back(std::move(right_cols[1])); + + return std::make_unique(std::move(left_cols)); +} + +/** + * @brief Sort nested-type column + * + * @param input table to sort + * @return std::unique_ptr + * + * @note if stability is desired, use `cudf::stable_sorted_order` + */ +std::unique_ptr sort_keys(cudf::table_view input) +{ + auto sort_order = cudf::sorted_order(cudf::table_view{{input.column(0)}}); + return cudf::gather(input, *sort_order); +} + +/** + * @brief Main for nested_types examples + * + * Command line parameters: + * 1. JSON input file name/path (default: "example.json") + * 2. JSON output file name/path (default: "output.json") + * 3. Memory resource (optional): "pool" or "cuda" (default: "pool") + * + * Example invocation from directory `cudf/cpp/examples/nested_types`: + * ./build/deduplication example.json output.json pool + * + */ +int main(int argc, char const** argv) +{ + std::string input_filepath; + std::string output_filepath; + std::string mr_name; + if (argc != 4 && argc != 1) { + std::cout << "Either provide all command-line arguments, or none to use defaults" << std::endl; + return 1; + } + if (argc == 1) { + input_filepath = "example.json"; + output_filepath = "output.json"; + mr_name = "pool"; + } else { + input_filepath = argv[1]; + output_filepath = argv[2]; + mr_name = argv[3]; + } + + auto pool = mr_name == "pool"; + auto resource = create_memory_resource(pool); + rmm::mr::set_current_device_resource(resource.get()); + + std::cout << "Reading " << input_filepath << "..." << std::endl; + // read input file + auto [input, metadata] = read_json(input_filepath); + + auto count = count_aggregate(input->view()); + + auto combined = join_count(input->view(), count->view()); + + auto sorted = sort_keys(combined->view()); + + metadata.schema_info.emplace_back("count"); + + std::cout << "Writing " << output_filepath << "..." << std::endl; + write_json(sorted->view(), metadata, output_filepath); + + return 0; +} diff --git a/cpp/examples/nested_types/example.json b/cpp/examples/nested_types/example.json new file mode 100644 index 00000000000..efaa37817d6 --- /dev/null +++ b/cpp/examples/nested_types/example.json @@ -0,0 +1,5 @@ +{"features": {"key": "a1", "values": [{"info": "message_1", "type": "device_a", "dt": 1688750001}]}, "source": "network_a", "quality": 0.7} +{"features": {"key": "a2", "values": [{"info": "message_2", "type": "device_a", "dt": 1688750002}]}, "source": "network_a", "quality": 0.7} +{"features": {"key": "a3", "values": [{"info": "message_3", "type": "device_a", "dt": 1688750003}]}, "source": "network_b", "quality": 0.8} +{"features": {"key": "a1", "values": [{"info": "message_1", "type": "device_a", "dt": 1688750001}]}, "source": "network_b", "quality": 0.9} +{"features": {"key": "a4", "values": [{"info": "message_4", "type": "device_a", "dt": 1688750004}]}, "source": "network_b", "quality": 0.9} diff --git a/cpp/examples/strings/CMakeLists.txt b/cpp/examples/strings/CMakeLists.txt index 4b500d3a92e..c90fa9dde16 100644 --- a/cpp/examples/strings/CMakeLists.txt +++ b/cpp/examples/strings/CMakeLists.txt @@ -8,23 +8,7 @@ project( LANGUAGES CXX CUDA ) -set(CPM_DOWNLOAD_VERSION v0.35.3) -file( - DOWNLOAD - https://github.com/cpm-cmake/CPM.cmake/releases/download/${CPM_DOWNLOAD_VERSION}/get_cpm.cmake - ${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake -) -include(${CMAKE_BINARY_DIR}/cmake/get_cpm.cmake) - -set(CUDF_TAG branch-23.12) -CPMFindPackage( - NAME cudf GIT_REPOSITORY https://github.com/rapidsai/cudf - GIT_TAG ${CUDF_TAG} - GIT_SHALLOW - TRUE - SOURCE_SUBDIR - cpp -) +include(../fetch_dependencies.cmake) list(APPEND CUDF_CUDA_FLAGS --expt-extended-lambda --expt-relaxed-constexpr) From 53127de4d9e06f9fa172ac34952f85104eb7bac9 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Thu, 16 Nov 2023 08:28:44 -0600 Subject: [PATCH 20/44] Remove needs: wheel-build-cudf. (#14427) This PR fixes a nightly test failure due to an extraneous `needs:` entry in `test.yaml`. ``` Invalid workflow file: .github/workflows/test.yaml#L100 The workflow is not valid. .github/workflows/test.yaml (Line: 100, Col: 12): Job 'unit-tests-cudf-pandas' depends on unknown job 'wheel-build-cudf'. ``` Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cudf/pull/14427 --- .github/workflows/test.yaml | 1 - 1 file changed, 1 deletion(-) diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 0b6b55069f6..0d4401160e1 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -97,7 +97,6 @@ jobs: sha: ${{ inputs.sha }} script: ci/test_wheel_dask_cudf.sh unit-tests-cudf-pandas: - needs: wheel-build-cudf secrets: inherit uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-23.12 with: From 8e1ef05b2b96775ce7e1a2f22894ec7a8ebb65a4 Mon Sep 17 00:00:00 2001 From: Divye Gala Date: Thu, 16 Nov 2023 16:43:29 -0500 Subject: [PATCH 21/44] Change `nullable()` to `has_nulls()` in `cudf::detail::gather` (#14363) In https://github.com/rapidsai/cudf/pull/13795, we found out that `nullable()` causes severe perf degradation for the nested-type case when the input is read from file via `cudf::io::read_json`. This is because the JSON reader adds a null mask for columns that don't have NULLs. This change is a no-overhead replacement that checks the actual null count instead of checking if a null mask is present. This PR also solves a bug in quantile/median groupby where NULLs were being [set](https://github.com/rapidsai/cudf/blob/8deb3dd7573000e7d87f18a9e2bbe39cf2932e10/cpp/src/groupby/sort/group_quantiles.cu#L73) but the null count was not updated. Authors: - Divye Gala (https://github.com/divyegala) - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Bradley Dice (https://github.com/bdice) - Vyas Ramasubramani (https://github.com/vyasr) - David Wendt (https://github.com/davidwendt) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/14363 --- cpp/include/cudf/detail/gather.cuh | 22 ++++++---- cpp/include/cudf/detail/null_mask.hpp | 17 ++++++++ cpp/include/cudf/table/table_view.hpp | 17 ++++++++ cpp/src/bitmask/null_mask.cu | 15 +++++++ cpp/src/groupby/sort/group_quantiles.cu | 17 +++++--- cpp/tests/join/join_tests.cpp | 56 ------------------------- 6 files changed, 75 insertions(+), 69 deletions(-) diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index 955f9914632..c9975ef2199 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -673,14 +673,20 @@ std::unique_ptr gather(table_view const& source_table, mr)); } - auto const nullable = bounds_policy == out_of_bounds_policy::NULLIFY || - std::any_of(source_table.begin(), source_table.end(), [](auto const& col) { - return col.nullable(); - }); - if (nullable) { - auto const op = bounds_policy == out_of_bounds_policy::NULLIFY ? gather_bitmask_op::NULLIFY - : gather_bitmask_op::DONT_CHECK; - gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr); + auto needs_new_bitmask = bounds_policy == out_of_bounds_policy::NULLIFY || + cudf::has_nested_nullable_columns(source_table); + if (needs_new_bitmask) { + needs_new_bitmask = needs_new_bitmask || cudf::has_nested_nulls(source_table); + if (needs_new_bitmask) { + auto const op = bounds_policy == out_of_bounds_policy::NULLIFY + ? gather_bitmask_op::NULLIFY + : gather_bitmask_op::DONT_CHECK; + gather_bitmask(source_table, gather_map_begin, destination_columns, op, stream, mr); + } else { + for (size_type i = 0; i < source_table.num_columns(); ++i) { + set_all_valid_null_masks(source_table.column(i), *destination_columns[i], stream, mr); + } + } } return std::make_unique
(std::move(destination_columns)); diff --git a/cpp/include/cudf/detail/null_mask.hpp b/cpp/include/cudf/detail/null_mask.hpp index 8c10bbe416f..74e2ccd2ea1 100644 --- a/cpp/include/cudf/detail/null_mask.hpp +++ b/cpp/include/cudf/detail/null_mask.hpp @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include #include @@ -259,6 +260,22 @@ cudf::size_type inplace_bitmask_and(device_span dest_mask, size_type mask_size_bits, rmm::cuda_stream_view stream); +/** + * @brief Recursively set valid null masks for all children. + * + * This function applies all valid null masks to the output column if input column satisfies + * `nullable() == true` condition + * + * @param input input column to check for nullability + * @param output output column to mirror nullability of input + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + */ +void set_all_valid_null_masks(column_view const& input, + column& output, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr); + } // namespace detail } // namespace cudf diff --git a/cpp/include/cudf/table/table_view.hpp b/cpp/include/cudf/table/table_view.hpp index b90b2dac012..5d9c930d137 100644 --- a/cpp/include/cudf/table/table_view.hpp +++ b/cpp/include/cudf/table/table_view.hpp @@ -336,6 +336,23 @@ inline bool has_nested_nulls(table_view const& input) }); } +/** + * @brief Returns True if the table has a nullable column at any level of the column hierarchy + * + * @param input The table to check for nullable columns + * @return True if the table has nullable columns at any level of the column hierarchy, false + * otherwise + */ +inline bool has_nested_nullable_columns(table_view const& input) +{ + return std::any_of(input.begin(), input.end(), [](auto const& col) { + return col.nullable() || + std::any_of(col.child_begin(), col.child_end(), [](auto const& child_col) { + return has_nested_nullable_columns(table_view{{child_col}}); + }); + }); +} + /** * @brief The function to collect all nullable columns at all nested levels in a given table. * diff --git a/cpp/src/bitmask/null_mask.cu b/cpp/src/bitmask/null_mask.cu index 3ff56eabe1e..1a1cbb17d15 100644 --- a/cpp/src/bitmask/null_mask.cu +++ b/cpp/src/bitmask/null_mask.cu @@ -510,6 +510,21 @@ std::pair bitmask_or(table_view const& view, return std::pair(std::move(null_mask), 0); } +void set_all_valid_null_masks(column_view const& input, + column& output, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + if (input.nullable()) { + auto mask = detail::create_null_mask(output.size(), mask_state::ALL_VALID, stream, mr); + output.set_null_mask(std::move(mask), 0); + + for (size_type i = 0; i < input.num_children(); ++i) { + set_all_valid_null_masks(input.child(i), output.child(i), stream, mr); + } + } +} + } // namespace detail // Create a bitmask from a specific range diff --git a/cpp/src/groupby/sort/group_quantiles.cu b/cpp/src/groupby/sort/group_quantiles.cu index a9edcfecbf7..a456d4b5964 100644 --- a/cpp/src/groupby/sort/group_quantiles.cu +++ b/cpp/src/groupby/sort/group_quantiles.cu @@ -49,6 +49,7 @@ struct calculate_quantile_fn { double const* d_quantiles; size_type num_quantiles; interpolation interpolation; + size_type* null_count; __device__ void operator()(size_type i) { @@ -68,11 +69,13 @@ struct calculate_quantile_fn { thrust::for_each_n(thrust::seq, thrust::make_counting_iterator(0), num_quantiles, - [d_result = d_result, segment_size, offset](size_type j) { - if (segment_size == 0) + [d_result = d_result, segment_size, offset, this](size_type j) { + if (segment_size == 0) { d_result.set_null(offset + j); - else + atomicAdd(this->null_count, 1); + } else { d_result.set_valid(offset + j); + } }); } }; @@ -104,6 +107,7 @@ struct quantiles_functor { auto values_view = column_device_view::create(values, stream); auto group_size_view = column_device_view::create(group_sizes, stream); auto result_view = mutable_column_device_view::create(result->mutable_view(), stream); + auto null_count = rmm::device_scalar(0, stream, mr); // For each group, calculate quantile if (!cudf::is_dictionary(values.type())) { @@ -118,7 +122,8 @@ struct quantiles_functor { group_offsets.data(), quantile.data(), static_cast(quantile.size()), - interpolation}); + interpolation, + null_count.data()}); } else { auto values_iter = cudf::dictionary::detail::make_dictionary_iterator(*values_view); thrust::for_each_n(rmm::exec_policy(stream), @@ -131,9 +136,11 @@ struct quantiles_functor { group_offsets.data(), quantile.data(), static_cast(quantile.size()), - interpolation}); + interpolation, + null_count.data()}); } + result->set_null_count(null_count.value(stream)); return result; } diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index 089db315748..a416df0c7c3 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -1941,62 +1941,6 @@ TEST_F(JoinTest, FullJoinWithStructsAndNulls) CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); } -TEST_F(JoinTest, Repro_StructsWithoutNullsPushedDown) -{ - // When joining on a STRUCT column, if the parent nulls are not reflected in - // the children, the join might produce incorrect results. - // - // In this test, a fact table of structs is joined against a dimension table. - // Both tables must match (only) on the NULL row. This will fail if the fact table's - // nulls are not pushed down into its children. - using ints = column_wrapper; - using structs = cudf::test::structs_column_wrapper; - using namespace cudf::test::iterators; - - auto make_table = [](auto&& col) { - auto columns = CVector{}; - columns.push_back(std::move(col)); - return cudf::table{std::move(columns)}; - }; - - auto const fact_table = [make_table] { - auto fact_ints = ints{0, 1, 2, 3, 4}; - auto fact_structs = structs{{fact_ints}, no_nulls()}.release(); - // Now set struct validity to invalidate index#3. - cudf::detail::set_null_mask( - fact_structs->mutable_view().null_mask(), 3, 4, false, cudf::get_default_stream()); - // Struct row#3 is null, but Struct.child has a non-null value. - return make_table(std::move(fact_structs)); - }(); - - auto const dimension_table = [make_table] { - auto dim_ints = ints{999}; - auto dim_structs = structs{{dim_ints}, null_at(0)}; - return make_table(dim_structs.release()); - }(); - - auto const result = inner_join(fact_table.view(), dimension_table.view(), {0}, {0}); - EXPECT_EQ(result->num_rows(), 1); // The null STRUCT rows should match. - - // Note: Join result might not have nulls pushed down, since it's an output of gather(). - // Must superimpose parent nulls before comparisons. - auto [superimposed_results, _] = cudf::structs::detail::push_down_nulls( - *result, cudf::get_default_stream(), rmm::mr::get_current_device_resource()); - - auto const expected = [] { - auto fact_ints = ints{0}; - auto fact_structs = structs{{fact_ints}, null_at(0)}; - auto dim_ints = ints{0}; - auto dim_structs = structs{{dim_ints}, null_at(0)}; - auto columns = CVector{}; - columns.push_back(fact_structs.release()); - columns.push_back(dim_structs.release()); - return cudf::table{std::move(columns)}; - }(); - - CUDF_TEST_EXPECT_TABLES_EQUIVALENT(superimposed_results, expected); -} - using lcw = cudf::test::lists_column_wrapper; using cudf::test::iterators::null_at; From bf63d1049db70c28ea961b677ad5f207aa648860 Mon Sep 17 00:00:00 2001 From: Ed Seidl Date: Thu, 16 Nov 2023 14:47:36 -0800 Subject: [PATCH 22/44] Add decoder for DELTA_BYTE_ARRAY to Parquet reader (#14101) Part of #13501. Adds ability to decode DELTA_BYTE_ARRAY encoded pages. Authors: - Ed Seidl (https://github.com/etseidl) - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - https://github.com/nvdbaranec - GALI PREM SAGAR (https://github.com/galipremsagar) URL: https://github.com/rapidsai/cudf/pull/14101 --- cpp/src/io/parquet/delta_binary.cuh | 58 ++- cpp/src/io/parquet/page_data.cu | 12 +- cpp/src/io/parquet/page_decode.cuh | 12 +- cpp/src/io/parquet/page_delta_decode.cu | 490 +++++++++++++++++- cpp/src/io/parquet/page_hdr.cu | 17 +- cpp/src/io/parquet/page_string_decode.cu | 344 ++++++++++-- cpp/src/io/parquet/parquet_gpu.hpp | 79 ++- cpp/src/io/parquet/reader_impl.cpp | 38 +- cpp/src/io/parquet/reader_impl_preprocess.cu | 4 +- .../tests/data/parquet/delta_byte_arr.parquet | Bin 0 -> 5783 bytes python/cudf/cudf/tests/test_parquet.py | 104 ++++ 11 files changed, 1044 insertions(+), 114 deletions(-) create mode 100644 python/cudf/cudf/tests/data/parquet/delta_byte_arr.parquet diff --git a/cpp/src/io/parquet/delta_binary.cuh b/cpp/src/io/parquet/delta_binary.cuh index e3b23f4c0a0..ccc28791071 100644 --- a/cpp/src/io/parquet/delta_binary.cuh +++ b/cpp/src/io/parquet/delta_binary.cuh @@ -39,15 +39,15 @@ namespace cudf::io::parquet::detail { // per mini-block. While encoding, the lowest delta value is subtracted from all the deltas in the // block to ensure that all encoded values are positive. The deltas for each mini-block are bit // packed using the same encoding as the RLE/Bit-Packing Hybrid encoder. -// -// DELTA_BYTE_ARRAY encoding (incremental encoding or front compression), is used for BYTE_ARRAY -// columns. For each element in a sequence of strings, a prefix length from the preceding string -// and a suffix is stored. The prefix lengths are DELTA_BINARY_PACKED encoded. The suffixes are -// encoded with DELTA_LENGTH_BYTE_ARRAY encoding, which is a DELTA_BINARY_PACKED list of suffix -// lengths, followed by the concatenated suffix data. -// we decode one mini-block at a time. max mini-block size seen is 64. -constexpr int delta_rolling_buf_size = 128; +// The largest mini-block size we can currently support. +constexpr int max_delta_mini_block_size = 64; + +// The first pass decodes `values_per_mb` values, and then the second pass does another +// batch of size `values_per_mb`. The largest value for values_per_miniblock among the +// major writers seems to be 64, so 2 * 64 should be good. We save the first value separately +// since it is not encoded in the first mini-block. +constexpr int delta_rolling_buf_size = 2 * max_delta_mini_block_size; /** * @brief Read a ULEB128 varint integer @@ -90,7 +90,8 @@ struct delta_binary_decoder { uleb128_t mini_block_count; // usually 4, chosen such that block_size/mini_block_count is a // multiple of 32 uleb128_t value_count; // total values encoded in the block - zigzag128_t last_value; // last value decoded, initialized to first_value from header + zigzag128_t first_value; // initial value, stored in the header + zigzag128_t last_value; // last value decoded uint32_t values_per_mb; // block_size / mini_block_count, must be multiple of 32 uint32_t current_value_idx; // current value index, initialized to 0 at start of block @@ -102,6 +103,13 @@ struct delta_binary_decoder { uleb128_t value[delta_rolling_buf_size]; // circular buffer of delta values + // returns the value stored in the `value` array at index + // `rolling_index(idx)`. If `idx` is `0`, then return `first_value`. + constexpr zigzag128_t value_at(size_type idx) + { + return idx == 0 ? first_value : value[rolling_index(idx)]; + } + // returns the number of values encoded in the block data. when all_values is true, // account for the first value in the header. otherwise just count the values encoded // in the mini-block data. @@ -145,7 +153,8 @@ struct delta_binary_decoder { block_size = get_uleb128(d_start, d_end); mini_block_count = get_uleb128(d_start, d_end); value_count = get_uleb128(d_start, d_end); - last_value = get_zz128(d_start, d_end); + first_value = get_zz128(d_start, d_end); + last_value = first_value; current_value_idx = 0; values_per_mb = block_size / mini_block_count; @@ -179,6 +188,28 @@ struct delta_binary_decoder { } } + // given start/end pointers in the data, find the end of the binary encoded block. when done, + // `this` will be initialized with the correct start and end positions. returns the end, which is + // start of data/next block. should only be called from thread 0. + inline __device__ uint8_t const* find_end_of_block(uint8_t const* start, uint8_t const* end) + { + // read block header + init_binary_block(start, end); + + // test for no encoded values. a single value will be in the block header. + if (value_count <= 1) { return block_start; } + + // read mini-block headers and skip over data + while (current_value_idx < num_encoded_values(false)) { + setup_next_mini_block(false); + } + // calculate the correct end of the block + auto const* const new_end = cur_mb == 0 ? block_start : cur_mb_start; + // re-init block with correct end + init_binary_block(start, new_end); + return new_end; + } + // decode the current mini-batch of deltas, and convert to values. // called by all threads in a warp, currently only one warp supported. inline __device__ void calc_mini_block_values(int lane_id) @@ -186,12 +217,9 @@ struct delta_binary_decoder { using cudf::detail::warp_size; if (current_value_idx >= value_count) { return; } - // need to save first value from header on first pass + // need to account for the first value from header on first pass if (current_value_idx == 0) { - if (lane_id == 0) { - current_value_idx++; - value[0] = last_value; - } + if (lane_id == 0) { current_value_idx++; } __syncwarp(); if (current_value_idx >= value_count) { return; } } diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index a783b489c02..0c53877f7c7 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -449,8 +449,13 @@ __global__ void __launch_bounds__(decode_block_size) int out_thread0; [[maybe_unused]] null_count_back_copier _{s, t}; - if (!setupLocalPageInfo( - s, &pages[page_idx], chunks, min_row, num_rows, mask_filter{KERNEL_MASK_GENERAL}, true)) { + if (!setupLocalPageInfo(s, + &pages[page_idx], + chunks, + min_row, + num_rows, + mask_filter{decode_kernel_mask::GENERAL}, + true)) { return; } @@ -486,6 +491,7 @@ __global__ void __launch_bounds__(decode_block_size) target_pos = min(s->nz_count, src_pos + decode_block_size - out_thread0); if (out_thread0 > 32) { target_pos = min(target_pos, s->dict_pos); } } + // TODO(ets): see if this sync can be removed __syncthreads(); if (t < 32) { // decode repetition and definition levels. @@ -603,7 +609,7 @@ __global__ void __launch_bounds__(decode_block_size) } struct mask_tform { - __device__ uint32_t operator()(PageInfo const& p) { return p.kernel_mask; } + __device__ uint32_t operator()(PageInfo const& p) { return static_cast(p.kernel_mask); } }; } // anonymous namespace diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index ab1cc68923d..4db9bd3904b 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -991,8 +991,15 @@ struct all_types_filter { * @brief Functor for setupLocalPageInfo that takes a mask of allowed types. */ struct mask_filter { - int mask; - __device__ inline bool operator()(PageInfo const& page) { return (page.kernel_mask & mask) != 0; } + uint32_t mask; + + __device__ mask_filter(uint32_t m) : mask(m) {} + __device__ mask_filter(decode_kernel_mask m) : mask(static_cast(m)) {} + + __device__ inline bool operator()(PageInfo const& page) + { + return BitAnd(mask, page.kernel_mask) != 0; + } }; /** @@ -1306,6 +1313,7 @@ inline __device__ bool setupLocalPageInfo(page_state_s* const s, s->dict_run = 0; } break; case Encoding::DELTA_BINARY_PACKED: + case Encoding::DELTA_BYTE_ARRAY: // nothing to do, just don't error break; default: { diff --git a/cpp/src/io/parquet/page_delta_decode.cu b/cpp/src/io/parquet/page_delta_decode.cu index bb5e5066b69..bc025c6fc3e 100644 --- a/cpp/src/io/parquet/page_delta_decode.cu +++ b/cpp/src/io/parquet/page_delta_decode.cu @@ -27,6 +27,277 @@ namespace cudf::io::parquet::detail { namespace { +constexpr int decode_block_size = 128; + +// DELTA_BYTE_ARRAY encoding (incremental encoding or front compression), is used for BYTE_ARRAY +// columns. For each element in a sequence of strings, a prefix length from the preceding string +// and a suffix is stored. The prefix lengths are DELTA_BINARY_PACKED encoded. The suffixes are +// encoded with DELTA_LENGTH_BYTE_ARRAY encoding, which is a DELTA_BINARY_PACKED list of suffix +// lengths, followed by the concatenated suffix data. +struct delta_byte_array_decoder { + uint8_t const* last_string; // pointer to last decoded string...needed for its prefix + uint8_t const* suffix_char_data; // pointer to the start of character data + + uint8_t* temp_buf; // buffer used when skipping values + uint32_t start_val; // decoded strings up to this index will be dumped to temp_buf + uint32_t last_string_len; // length of the last decoded string + + delta_binary_decoder prefixes; // state of decoder for prefix lengths + delta_binary_decoder suffixes; // state of decoder for suffix lengths + + // initialize the prefixes and suffixes blocks + __device__ void init(uint8_t const* start, uint8_t const* end, uint32_t start_idx, uint8_t* temp) + { + auto const* suffix_start = prefixes.find_end_of_block(start, end); + suffix_char_data = suffixes.find_end_of_block(suffix_start, end); + last_string = nullptr; + temp_buf = temp; + start_val = start_idx; + } + + // kind of like an inclusive scan for strings. takes prefix_len bytes from preceding + // string and prepends to the suffix we've already copied into place. called from + // within loop over values_in_mb, so this only needs to handle a single warp worth of data + // at a time. + __device__ void string_scan(uint8_t* strings_out, + uint8_t const* last_string, + uint32_t start_idx, + uint32_t end_idx, + uint32_t offset, + uint32_t lane_id) + { + using cudf::detail::warp_size; + + // let p(n) === length(prefix(string_n)) + // + // if p(n-1) > p(n), then string_n can be completed when string_n-2 is completed. likewise if + // p(m) > p(n), then string_n can be completed with string_m-1. however, if p(m) < p(n), then m + // is a "blocker" for string_n; string_n can be completed only after string_m is. + // + // we will calculate the nearest blocking position for each lane, and then fill in string_0. we + // then iterate, finding all lanes that have had their "blocker" filled in and completing them. + // when all lanes are filled in, we return. this will still hit the worst case if p(n-1) < p(n) + // for all n + __shared__ __align__(8) int64_t prefix_lens[warp_size]; + __shared__ __align__(8) uint8_t const* offsets[warp_size]; + + uint32_t const ln_idx = start_idx + lane_id; + uint64_t prefix_len = ln_idx < end_idx ? prefixes.value_at(ln_idx) : 0; + uint8_t* const lane_out = ln_idx < end_idx ? strings_out + offset : nullptr; + + prefix_lens[lane_id] = prefix_len; + offsets[lane_id] = lane_out; + + // if all prefix_len's are zero, then there's nothing to do + if (__all_sync(0xffff'ffff, prefix_len == 0)) { return; } + + // find a neighbor to the left that has a prefix length less than this lane. once that + // neighbor is complete, this lane can be completed. + int blocker = lane_id - 1; + while (blocker > 0 && prefix_lens[blocker] != 0 && prefix_len <= prefix_lens[blocker]) { + blocker--; + } + + // fill in lane 0 (if necessary) + if (lane_id == 0 && prefix_len > 0) { + memcpy(lane_out, last_string, prefix_len); + prefix_lens[0] = prefix_len = 0; + } + __syncwarp(); + + // now fill in blockers until done + for (uint32_t i = 1; i < warp_size && i + start_idx < end_idx; i++) { + if (prefix_len != 0 && prefix_lens[blocker] == 0 && lane_out != nullptr) { + memcpy(lane_out, offsets[blocker], prefix_len); + prefix_lens[lane_id] = prefix_len = 0; + } + + // check for finished + if (__all_sync(0xffff'ffff, prefix_len == 0)) { return; } + } + } + + // calculate a mini-batch of string values, writing the results to + // `strings_out`. starting at global index `start_idx` and decoding + // up to `num_values` strings. + // called by all threads in a warp. used for strings <= 32 chars. + // returns number of bytes written + __device__ size_t calculate_string_values(uint8_t* strings_out, + uint32_t start_idx, + uint32_t num_values, + uint32_t lane_id) + { + using cudf::detail::warp_size; + using WarpScan = cub::WarpScan; + __shared__ WarpScan::TempStorage scan_temp; + + if (start_idx >= suffixes.value_count) { return 0; } + auto end_idx = start_idx + min(suffixes.values_per_mb, num_values); + end_idx = min(end_idx, static_cast(suffixes.value_count)); + + auto p_strings_out = strings_out; + auto p_temp_out = temp_buf; + + auto copy_batch = [&](uint8_t* out, uint32_t idx, uint32_t end) { + uint32_t const ln_idx = idx + lane_id; + + // calculate offsets into suffix data + uint64_t const suffix_len = ln_idx < end ? suffixes.value_at(ln_idx) : 0; + uint64_t suffix_off = 0; + WarpScan(scan_temp).ExclusiveSum(suffix_len, suffix_off); + + // calculate offsets into string data + uint64_t const prefix_len = ln_idx < end ? prefixes.value_at(ln_idx) : 0; + uint64_t const string_len = prefix_len + suffix_len; + + // get offset into output for each lane + uint64_t string_off, warp_total; + WarpScan(scan_temp).ExclusiveSum(string_len, string_off, warp_total); + auto const so_ptr = out + string_off; + + // copy suffixes into string data + if (ln_idx < end) { memcpy(so_ptr + prefix_len, suffix_char_data + suffix_off, suffix_len); } + __syncwarp(); + + // copy prefixes into string data. + string_scan(out, last_string, idx, end, string_off, lane_id); + + // save the position of the last computed string. this will be used in + // the next iteration to reconstruct the string in lane 0. + if (ln_idx == end - 1 || (ln_idx < end && lane_id == 31)) { + // set last_string to this lane's string + last_string = out + string_off; + last_string_len = string_len; + // and consume used suffix_char_data + suffix_char_data += suffix_off + suffix_len; + } + + return warp_total; + }; + + uint64_t string_total = 0; + for (int idx = start_idx; idx < end_idx; idx += warp_size) { + auto const n_in_batch = min(warp_size, end_idx - idx); + // account for the case where start_val occurs in the middle of this batch + if (idx < start_val && idx + n_in_batch > start_val) { + // dump idx...start_val into temp_buf + copy_batch(p_temp_out, idx, start_val); + __syncwarp(); + + // start_val...idx + n_in_batch into strings_out + auto nbytes = copy_batch(p_strings_out, start_val, idx + n_in_batch); + p_strings_out += nbytes; + string_total = nbytes; + } else { + if (idx < start_val) { + p_temp_out += copy_batch(p_temp_out, idx, end_idx); + } else { + auto nbytes = copy_batch(p_strings_out, idx, end_idx); + p_strings_out += nbytes; + string_total += nbytes; + } + } + __syncwarp(); + } + + return string_total; + } + + // character parallel version of CalculateStringValues(). This is faster for strings longer than + // 32 chars. + __device__ size_t calculate_string_values_cp(uint8_t* strings_out, + uint32_t start_idx, + uint32_t num_values, + uint32_t lane_id) + { + using cudf::detail::warp_size; + __shared__ __align__(8) uint8_t* so_ptr; + + if (start_idx >= suffixes.value_count) { return; } + auto end_idx = start_idx + min(suffixes.values_per_mb, num_values); + end_idx = min(end_idx, static_cast(suffixes.value_count)); + + if (lane_id == 0) { so_ptr = start_idx < start_val ? temp_buf : strings_out; } + __syncwarp(); + + uint64_t string_total = 0; + for (int idx = start_idx; idx < end_idx; idx++) { + uint64_t const suffix_len = suffixes.value_at(idx); + uint64_t const prefix_len = prefixes.value_at(idx); + uint64_t const string_len = prefix_len + suffix_len; + + // copy prefix and suffix data into current strings_out position + // for longer strings use a 4-byte version stolen from gather_chars_fn_string_parallel. + if (string_len > 64) { + if (prefix_len > 0) { wideStrcpy(so_ptr, last_string, prefix_len, lane_id); } + if (suffix_len > 0) { + wideStrcpy(so_ptr + prefix_len, suffix_char_data, suffix_len, lane_id); + } + } else { + for (int i = lane_id; i < string_len; i += warp_size) { + so_ptr[i] = i < prefix_len ? last_string[i] : suffix_char_data[i - prefix_len]; + } + } + __syncwarp(); + + if (idx >= start_val) { string_total += string_len; } + + if (lane_id == 0) { + last_string = so_ptr; + last_string_len = string_len; + suffix_char_data += suffix_len; + if (idx == start_val - 1) { + so_ptr = strings_out; + } else { + so_ptr += string_len; + } + } + __syncwarp(); + } + + return string_total; + } + + // dump strings before start_val to temp buf + __device__ void skip(bool use_char_ll) + { + using cudf::detail::warp_size; + int const t = threadIdx.x; + int const lane_id = t % warp_size; + + // is this even necessary? return if asking to skip the whole block. + if (start_val >= prefixes.num_encoded_values(true)) { return; } + + // prefixes and suffixes will have the same parameters (it's checked earlier) + auto const batch_size = prefixes.values_per_mb; + + uint32_t skip_pos = 0; + while (prefixes.current_value_idx < start_val) { + // warp 0 gets prefixes and warp 1 gets suffixes + auto* const db = t < 32 ? &prefixes : &suffixes; + + // this will potentially decode past start_val, but that's ok + if (t < 64) { db->decode_batch(); } + __syncthreads(); + + // warp 0 decodes the batch. + if (t < 32) { + auto const num_to_decode = min(batch_size, start_val - skip_pos); + auto const bytes_written = + use_char_ll ? calculate_string_values_cp(temp_buf, skip_pos, num_to_decode, lane_id) + : calculate_string_values(temp_buf, skip_pos, num_to_decode, lane_id); + // store last_string someplace safe in temp buffer + if (t == 0) { + memcpy(temp_buf + bytes_written, last_string, last_string_len); + last_string = temp_buf + bytes_written; + } + } + skip_pos += prefixes.values_per_mb; + __syncthreads(); + } + } +}; + // Decode page data that is DELTA_BINARY_PACKED encoded. This encoding is // only used for int32 and int64 physical types (and appears to only be used // with V2 page headers; see https://www.mail-archive.com/dev@parquet.apache.org/msg11826.html). @@ -52,13 +323,9 @@ __global__ void __launch_bounds__(96) auto* const db = &db_state; [[maybe_unused]] null_count_back_copier _{s, t}; - if (!setupLocalPageInfo(s, - &pages[page_idx], - chunks, - min_row, - num_rows, - mask_filter{KERNEL_MASK_DELTA_BINARY}, - true)) { + auto const mask = decode_kernel_mask::DELTA_BINARY; + if (!setupLocalPageInfo( + s, &pages[page_idx], chunks, min_row, num_rows, mask_filter{mask}, true)) { return; } @@ -78,6 +345,10 @@ __global__ void __launch_bounds__(96) __syncthreads(); auto const batch_size = db->values_per_mb; + if (batch_size > max_delta_mini_block_size) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), error_code); + return; + } // if skipped_leaf_values is non-zero, then we need to decode up to the first mini-block // that has a value we need. @@ -93,6 +364,7 @@ __global__ void __launch_bounds__(96) } else { // warp2 target_pos = min(s->nz_count, src_pos + batch_size); } + // TODO(ets): see if this sync can be removed __syncthreads(); // warp0 will decode the rep/def levels, warp1 will unpack a mini-batch of deltas. @@ -125,23 +397,12 @@ __global__ void __launch_bounds__(96) // place value for this thread if (dst_pos >= 0 && sp < target_pos) { void* const dst = nesting_info_base[leaf_level_index].data_out + dst_pos * s->dtype_len; + auto const val = db->value_at(sp + skipped_leaf_values); switch (s->dtype_len) { - case 1: - *static_cast(dst) = - db->value[rolling_index(sp + skipped_leaf_values)]; - break; - case 2: - *static_cast(dst) = - db->value[rolling_index(sp + skipped_leaf_values)]; - break; - case 4: - *static_cast(dst) = - db->value[rolling_index(sp + skipped_leaf_values)]; - break; - case 8: - *static_cast(dst) = - db->value[rolling_index(sp + skipped_leaf_values)]; - break; + case 1: *static_cast(dst) = val; break; + case 2: *static_cast(dst) = val; break; + case 4: *static_cast(dst) = val; break; + case 8: *static_cast(dst) = val; break; } } } @@ -154,6 +415,164 @@ __global__ void __launch_bounds__(96) if (t == 0 and s->error != 0) { set_error(s->error, error_code); } } +// Decode page data that is DELTA_BYTE_ARRAY packed. This encoding consists of a DELTA_BINARY_PACKED +// array of prefix lengths, followed by a DELTA_BINARY_PACKED array of suffix lengths, followed by +// the suffixes (technically the suffixes are DELTA_LENGTH_BYTE_ARRAY encoded). The latter two can +// be used to create an offsets array for the suffix data, but then this needs to be combined with +// the prefix lengths to do the final decode for each value. Because the lengths of the prefixes and +// suffixes are not encoded in the header, we're going to have to first do a quick pass through them +// to find the start/end of each structure. +template +__global__ void __launch_bounds__(decode_block_size) + gpuDecodeDeltaByteArray(PageInfo* pages, + device_span chunks, + size_t min_row, + size_t num_rows, + int32_t* error_code) +{ + using cudf::detail::warp_size; + __shared__ __align__(16) delta_byte_array_decoder db_state; + __shared__ __align__(16) page_state_s state_g; + __shared__ __align__(16) page_state_buffers_s state_buffers; + + page_state_s* const s = &state_g; + auto* const sb = &state_buffers; + int const page_idx = blockIdx.x; + int const t = threadIdx.x; + int const lane_id = t % warp_size; + auto* const prefix_db = &db_state.prefixes; + auto* const suffix_db = &db_state.suffixes; + auto* const dba = &db_state; + [[maybe_unused]] null_count_back_copier _{s, t}; + + auto const mask = decode_kernel_mask::DELTA_BYTE_ARRAY; + if (!setupLocalPageInfo( + s, &pages[page_idx], chunks, min_row, num_rows, mask_filter{mask}, true)) { + return; + } + + bool const has_repetition = s->col.max_level[level_type::REPETITION] > 0; + + // choose a character parallel string copy when the average string is longer than a warp + auto const use_char_ll = (s->page.str_bytes / s->page.num_valids) > cudf::detail::warp_size; + + // copying logic from gpuDecodePageData. + PageNestingDecodeInfo const* nesting_info_base = s->nesting_info; + + __shared__ level_t rep[delta_rolling_buf_size]; // circular buffer of repetition level values + __shared__ level_t def[delta_rolling_buf_size]; // circular buffer of definition level values + + // skipped_leaf_values will always be 0 for flat hierarchies. + uint32_t const skipped_leaf_values = s->page.skipped_leaf_values; + + if (t == 0) { + // initialize the prefixes and suffixes blocks + dba->init(s->data_start, s->data_end, s->page.start_val, s->page.temp_string_buf); + } + __syncthreads(); + + // assert that prefix and suffix have same mini-block size + if (prefix_db->values_per_mb != suffix_db->values_per_mb or + prefix_db->block_size != suffix_db->block_size or + prefix_db->value_count != suffix_db->value_count) { + set_error(static_cast(decode_error::DELTA_PARAM_MISMATCH), error_code); + return; + } + + // pointer to location to output final strings + int const leaf_level_index = s->col.max_nesting_depth - 1; + auto strings_data = nesting_info_base[leaf_level_index].string_out; + + auto const batch_size = prefix_db->values_per_mb; + if (batch_size > max_delta_mini_block_size) { + set_error(static_cast(decode_error::DELTA_PARAMS_UNSUPPORTED), error_code); + return; + } + + // if this is a bounds page and nested, then we need to skip up front. non-nested will work + // its way through the page. + int string_pos = has_repetition ? s->page.start_val : 0; + auto const is_bounds_pg = is_bounds_page(s, min_row, num_rows, has_repetition); + if (is_bounds_pg && string_pos > 0) { dba->skip(use_char_ll); } + + while (!s->error && (s->input_value_count < s->num_input_values || s->src_pos < s->nz_count)) { + uint32_t target_pos; + uint32_t const src_pos = s->src_pos; + + if (t < 3 * warp_size) { // warp 0..2 + target_pos = min(src_pos + 2 * batch_size, s->nz_count + s->first_row + batch_size); + } else { // warp 3 + target_pos = min(s->nz_count, src_pos + batch_size); + } + // TODO(ets): see if this sync can be removed + __syncthreads(); + + // warp0 will decode the rep/def levels, warp1 will unpack a mini-batch of prefixes, warp 2 will + // unpack a mini-batch of suffixes. warp3 waits one cycle for warps 0-2 to produce a batch, and + // then stuffs values into the proper location in the output. + if (t < warp_size) { + // decode repetition and definition levels. + // - update validity vectors + // - updates offsets (for nested columns) + // - produces non-NULL value indices in s->nz_idx for subsequent decoding + gpuDecodeLevels(s, sb, target_pos, rep, def, t); + + } else if (t < 2 * warp_size) { + // warp 1 + prefix_db->decode_batch(); + + } else if (t < 3 * warp_size) { + // warp 2 + suffix_db->decode_batch(); + + } else if (src_pos < target_pos) { + // warp 3 + + int const nproc = min(batch_size, s->page.end_val - string_pos); + strings_data += use_char_ll + ? dba->calculate_string_values_cp(strings_data, string_pos, nproc, lane_id) + : dba->calculate_string_values(strings_data, string_pos, nproc, lane_id); + string_pos += nproc; + + // process the mini-block in batches of 32 + for (uint32_t sp = src_pos + lane_id; sp < src_pos + batch_size; sp += 32) { + // the position in the output column/buffer + int dst_pos = sb->nz_idx[rolling_index(sp)]; + + // handle skip_rows here. flat hierarchies can just skip up to first_row. + if (!has_repetition) { dst_pos -= s->first_row; } + + if (dst_pos >= 0 && sp < target_pos) { + auto const offptr = + reinterpret_cast(nesting_info_base[leaf_level_index].data_out) + dst_pos; + auto const src_idx = sp + skipped_leaf_values; + *offptr = prefix_db->value_at(src_idx) + suffix_db->value_at(src_idx); + } + __syncwarp(); + } + + if (lane_id == 0) { s->src_pos = src_pos + batch_size; } + } + + __syncthreads(); + } + + // now turn array of lengths into offsets + int value_count = nesting_info_base[leaf_level_index].value_count; + + // if no repetition we haven't calculated start/end bounds and instead just skipped + // values until we reach first_row. account for that here. + if (!has_repetition) { value_count -= s->first_row; } + + auto const offptr = reinterpret_cast(nesting_info_base[leaf_level_index].data_out); + block_excl_sum(offptr, value_count, s->page.str_offset); + + if (t == 0 and s->error != 0) { + cuda::atomic_ref ref{*error_code}; + ref.fetch_or(s->error, cuda::std::memory_order_relaxed); + } +} + } // anonymous namespace /** @@ -181,4 +600,29 @@ void __host__ DecodeDeltaBinary(cudf::detail::hostdevice_vector& pages } } +/** + * @copydoc cudf::io::parquet::gpu::DecodeDeltaByteArray + */ +void __host__ DecodeDeltaByteArray(cudf::detail::hostdevice_vector& pages, + cudf::detail::hostdevice_vector const& chunks, + size_t num_rows, + size_t min_row, + int level_type_size, + int32_t* error_code, + rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(pages.size() > 0, "There is no page to decode"); + + dim3 const dim_block(decode_block_size, 1); + dim3 const dim_grid(pages.size(), 1); // 1 threadblock per page + + if (level_type_size == 1) { + gpuDecodeDeltaByteArray<<>>( + pages.device_ptr(), chunks, min_row, num_rows, error_code); + } else { + gpuDecodeDeltaByteArray<<>>( + pages.device_ptr(), chunks, min_row, num_rows, error_code); + } +} + } // namespace cudf::io::parquet::detail diff --git a/cpp/src/io/parquet/page_hdr.cu b/cpp/src/io/parquet/page_hdr.cu index 22add2fffc6..595dd40cdc2 100644 --- a/cpp/src/io/parquet/page_hdr.cu +++ b/cpp/src/io/parquet/page_hdr.cu @@ -146,18 +146,21 @@ __device__ void skip_struct_field(byte_stream_s* bs, int field_type) * @param chunk Column chunk the page belongs to * @return `kernel_mask_bits` value for the given page */ -__device__ uint32_t kernel_mask_for_page(PageInfo const& page, ColumnChunkDesc const& chunk) +__device__ decode_kernel_mask kernel_mask_for_page(PageInfo const& page, + ColumnChunkDesc const& chunk) { - if (page.flags & PAGEINFO_FLAGS_DICTIONARY) { return 0; } + if (page.flags & PAGEINFO_FLAGS_DICTIONARY) { return decode_kernel_mask::NONE; } if (page.encoding == Encoding::DELTA_BINARY_PACKED) { - return KERNEL_MASK_DELTA_BINARY; + return decode_kernel_mask::DELTA_BINARY; + } else if (page.encoding == Encoding::DELTA_BYTE_ARRAY) { + return decode_kernel_mask::DELTA_BYTE_ARRAY; } else if (is_string_col(chunk)) { - return KERNEL_MASK_STRING; + return decode_kernel_mask::STRING; } // non-string, non-delta - return KERNEL_MASK_GENERAL; + return decode_kernel_mask::GENERAL; } /** @@ -380,7 +383,9 @@ __global__ void __launch_bounds__(128) bs->page.skipped_values = -1; bs->page.skipped_leaf_values = 0; bs->page.str_bytes = 0; - bs->page.kernel_mask = 0; + bs->page.temp_string_size = 0; + bs->page.temp_string_buf = nullptr; + bs->page.kernel_mask = decode_kernel_mask::NONE; } num_values = bs->ck.num_values; page_info = bs->ck.page_info; diff --git a/cpp/src/io/parquet/page_string_decode.cu b/cpp/src/io/parquet/page_string_decode.cu index 4c7d8e3c20a..e29db042401 100644 --- a/cpp/src/io/parquet/page_string_decode.cu +++ b/cpp/src/io/parquet/page_string_decode.cu @@ -14,20 +14,28 @@ * limitations under the License. */ +#include "delta_binary.cuh" #include "page_decode.cuh" #include "page_string_utils.cuh" #include +#include #include +#include +#include + +#include + namespace cudf::io::parquet::detail { namespace { -constexpr int preprocess_block_size = 512; -constexpr int decode_block_size = 128; -constexpr int rolling_buf_size = decode_block_size * 2; -constexpr int preproc_buf_size = LEVEL_DECODE_BUF_SIZE; +constexpr int preprocess_block_size = 512; +constexpr int decode_block_size = 128; +constexpr int delta_preproc_block_size = 64; +constexpr int rolling_buf_size = decode_block_size * 2; +constexpr int preproc_buf_size = LEVEL_DECODE_BUF_SIZE; /** * @brief Compute the start and end page value bounds for this page @@ -450,12 +458,107 @@ __device__ size_t totalPlainEntriesSize(uint8_t const* data, } /** - * @brief Kernel for computing string page output size information. + * @brief Compute string size information for DELTA_BYTE_ARRAY encoded strings. + * + * This traverses the packed prefix and suffix lengths, summing them to obtain the total + * number of bytes needed for the decoded string data. It also calculates an upper bound + * for the largest string length to obtain an upper bound on temporary space needed if + * rows will be skipped. + * + * Called with 64 threads. + * + * @param data Pointer to the start of the page data stream + * @param end Pointer to the end of the page data stream + * @param start_value Do not count values that occur before this index + * @param end_value Do not count values that occur after this index + * @return A pair of `size_t` values representing the total string size and temp buffer size + * required for decoding + */ +__device__ thrust::pair totalDeltaByteArraySize(uint8_t const* data, + uint8_t const* end, + int start_value, + int end_value) +{ + using cudf::detail::warp_size; + using WarpReduce = cub::WarpReduce; + __shared__ typename WarpReduce::TempStorage temp_storage[2]; + + __shared__ __align__(16) delta_binary_decoder prefixes; + __shared__ __align__(16) delta_binary_decoder suffixes; + + int const t = threadIdx.x; + int const lane_id = t % warp_size; + int const warp_id = t / warp_size; + + if (t == 0) { + auto const* suffix_start = prefixes.find_end_of_block(data, end); + suffixes.init_binary_block(suffix_start, end); + } + __syncthreads(); + + // two warps will traverse the prefixes and suffixes and sum them up + auto const db = t < warp_size ? &prefixes : t < 2 * warp_size ? &suffixes : nullptr; + + size_t total_bytes = 0; + uleb128_t max_len = 0; + + if (db != nullptr) { + // initialize with first value (which is stored in last_value) + if (lane_id == 0 && start_value == 0) { total_bytes = db->last_value; } + + uleb128_t lane_sum = 0; + uleb128_t lane_max = 0; + while (db->current_value_idx < end_value && + db->current_value_idx < db->num_encoded_values(true)) { + // calculate values for current mini-block + db->calc_mini_block_values(lane_id); + + // get per lane sum for mini-block + for (uint32_t i = 0; i < db->values_per_mb; i += 32) { + uint32_t const idx = db->current_value_idx + i + lane_id; + if (idx >= start_value && idx < end_value && idx < db->value_count) { + lane_sum += db->value[rolling_index(idx)]; + lane_max = max(lane_max, db->value[rolling_index(idx)]); + } + } + + if (lane_id == 0) { db->setup_next_mini_block(true); } + __syncwarp(); + } + + // get sum for warp. + // note: warp_sum will only be valid on lane 0. + auto const warp_sum = WarpReduce(temp_storage[warp_id]).Sum(lane_sum); + auto const warp_max = WarpReduce(temp_storage[warp_id]).Reduce(lane_max, cub::Max()); + + if (lane_id == 0) { + total_bytes += warp_sum; + max_len = warp_max; + } + } + __syncthreads(); + + // now sum up total_bytes from the two warps + auto const final_bytes = + cudf::detail::single_lane_block_sum_reduce(total_bytes); + + // Sum up prefix and suffix max lengths to get a max possible string length. Multiply that + // by the number of strings in a mini-block, plus one to save the last string. + auto const temp_bytes = + cudf::detail::single_lane_block_sum_reduce(max_len) * + (db->values_per_mb + 1); + + return {final_bytes, temp_bytes}; +} + +/** + * @brief Kernel for computing string page bounds information. * - * String columns need accurate data size information to preallocate memory in the column buffer to - * store the char data. This calls a kernel to calculate information needed by the string decoding - * kernel. On exit, the `str_bytes`, `num_nulls`, and `num_valids` fields of the PageInfo struct - * are updated. This call ignores non-string columns. + * This kernel traverses the repetition and definition level data to determine start and end values + * for pages with string-like data. Also calculates the number of null and valid values in the + * page. Does nothing if the page mask is neither `STRING` nor `DELTA_BYTE_ARRAY`. On exit the + * `num_nulls`, `num_valids`, `start_val` and `end_val` fields of the `PageInfo` struct will be + * populated. * * @param pages All pages to be decoded * @param chunks All chunks to be decoded @@ -464,7 +567,7 @@ __device__ size_t totalPlainEntriesSize(uint8_t const* data, * @tparam level_t Type used to store decoded repetition and definition levels */ template -__global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSizes( +__global__ void __launch_bounds__(preprocess_block_size) gpuComputeStringPageBounds( PageInfo* pages, device_span chunks, size_t min_row, size_t num_rows) { __shared__ __align__(16) page_state_s state_g; @@ -474,8 +577,13 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz int const t = threadIdx.x; PageInfo* const pp = &pages[page_idx]; - // reset str_bytes to 0 in case it's already been calculated - if (t == 0) { pp->str_bytes = 0; } + if (t == 0) { + s->page.num_nulls = 0; + s->page.num_valids = 0; + // reset str_bytes to 0 in case it's already been calculated (esp needed for chunked reads). + // TODO: need to rethink this once str_bytes is in the statistics + pp->str_bytes = 0; + } // whether or not we have repetition levels (lists) bool const has_repetition = chunks[pp->chunk_idx].max_level[level_type::REPETITION] > 0; @@ -491,23 +599,11 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz {rep_runs}}; // setup page info - if (!setupLocalPageInfo( - s, pp, chunks, min_row, num_rows, mask_filter{KERNEL_MASK_STRING}, false)) { - return; - } - - if (!t) { - s->page.num_nulls = 0; - s->page.num_valids = 0; - s->page.str_bytes = 0; - } - __syncthreads(); + auto const mask = BitOr(decode_kernel_mask::STRING, decode_kernel_mask::DELTA_BYTE_ARRAY); + if (!setupLocalPageInfo(s, pp, chunks, min_row, num_rows, mask_filter{mask}, true)) { return; } bool const is_bounds_pg = is_bounds_page(s, min_row, num_rows, has_repetition); - // if we're skipping this page anyway, no need to count it - if (!is_bounds_pg && !is_page_contained(s, min_row, num_rows)) { return; } - // find start/end value indices auto const [start_value, end_value] = page_bounds(s, min_row, num_rows, is_bounds_pg, has_repetition, decoders); @@ -516,7 +612,106 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz if (t == 0) { pp->num_nulls = s->page.num_nulls; pp->num_valids = s->page.num_valids; + pp->start_val = start_value; + pp->end_val = end_value; } +} + +/** + * @brief Kernel for computing string page output size information for delta_byte_array encoding. + * + * This call ignores columns that are not DELTA_BYTE_ARRAY encoded. On exit the `str_bytes` field + * of the `PageInfo` struct will be populated. Also fills in the `temp_string_size` field if rows + * are to be skipped. + * + * @param pages All pages to be decoded + * @param chunks All chunks to be decoded + * @param min_rows crop all rows below min_row + * @param num_rows Maximum number of rows to read + */ +__global__ void __launch_bounds__(delta_preproc_block_size) gpuComputeDeltaPageStringSizes( + PageInfo* pages, device_span chunks, size_t min_row, size_t num_rows) +{ + __shared__ __align__(16) page_state_s state_g; + + page_state_s* const s = &state_g; + int const page_idx = blockIdx.x; + int const t = threadIdx.x; + PageInfo* const pp = &pages[page_idx]; + + // whether or not we have repetition levels (lists) + bool const has_repetition = chunks[pp->chunk_idx].max_level[level_type::REPETITION] > 0; + + // setup page info + auto const mask = decode_kernel_mask::DELTA_BYTE_ARRAY; + if (!setupLocalPageInfo(s, pp, chunks, min_row, num_rows, mask_filter{mask}, true)) { return; } + + auto const start_value = pp->start_val; + + // if data size is known, can short circuit here + if ((chunks[pp->chunk_idx].data_type & 7) == FIXED_LEN_BYTE_ARRAY) { + if (t == 0) { + pp->str_bytes = pp->num_valids * s->dtype_len_in; + + // only need temp space if we're skipping values + if (start_value > 0) { + // just need to parse the header of the first delta binary block to get values_per_mb + delta_binary_decoder db; + db.init_binary_block(s->data_start, s->data_end); + // save enough for one mini-block plus some extra to save the last_string + pp->temp_string_size = s->dtype_len_in * (db.values_per_mb + 1); + } + } + } else { + // now process string info in the range [start_value, end_value) + // set up for decoding strings...can be either plain or dictionary + uint8_t const* data = s->data_start; + uint8_t const* const end = s->data_end; + auto const end_value = pp->end_val; + + auto const [len, temp_bytes] = totalDeltaByteArraySize(data, end, start_value, end_value); + + if (t == 0) { + // TODO check for overflow + pp->str_bytes = len; + + // only need temp space if we're skipping values + if (start_value > 0) { pp->temp_string_size = temp_bytes; } + } + } +} + +/** + * @brief Kernel for computing string page output size information. + * + * This call ignores non-string columns. On exit the `str_bytes` field of the `PageInfo` struct will + * be populated. + * + * @param pages All pages to be decoded + * @param chunks All chunks to be decoded + * @param min_rows crop all rows below min_row + * @param num_rows Maximum number of rows to read + */ +__global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSizes( + PageInfo* pages, device_span chunks, size_t min_row, size_t num_rows) +{ + __shared__ __align__(16) page_state_s state_g; + + page_state_s* const s = &state_g; + int const page_idx = blockIdx.x; + int const t = threadIdx.x; + PageInfo* const pp = &pages[page_idx]; + + // whether or not we have repetition levels (lists) + bool const has_repetition = chunks[pp->chunk_idx].max_level[level_type::REPETITION] > 0; + + // setup page info + if (!setupLocalPageInfo( + s, pp, chunks, min_row, num_rows, mask_filter{decode_kernel_mask::STRING}, true)) { + return; + } + + bool const is_bounds_pg = is_bounds_page(s, min_row, num_rows, has_repetition); auto const& col = s->col; size_t str_bytes = 0; @@ -530,6 +725,8 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz uint8_t const* const end = s->data_end; uint8_t const* dict_base = nullptr; int dict_size = 0; + auto const start_value = pp->start_val; + auto const end_value = pp->end_val; switch (pp->encoding) { case Encoding::PLAIN_DICTIONARY: @@ -561,6 +758,9 @@ __global__ void __launch_bounds__(preprocess_block_size) gpuComputePageStringSiz if (t == 0) { // TODO check for overflow pp->str_bytes = str_bytes; + + // only need temp space for delta + pp->temp_string_size = 0; } } @@ -586,6 +786,7 @@ __global__ void __launch_bounds__(decode_block_size) size_t num_rows, int32_t* error_code) { + using cudf::detail::warp_size; __shared__ __align__(16) page_state_s state_g; __shared__ __align__(4) size_type last_offset; __shared__ __align__(16) @@ -596,10 +797,12 @@ __global__ void __launch_bounds__(decode_block_size) auto* const sb = &state_buffers; int const page_idx = blockIdx.x; int const t = threadIdx.x; + int const lane_id = t % warp_size; [[maybe_unused]] null_count_back_copier _{s, t}; + auto const mask = decode_kernel_mask::STRING; if (!setupLocalPageInfo( - s, &pages[page_idx], chunks, min_row, num_rows, mask_filter{KERNEL_MASK_STRING}, true)) { + s, &pages[page_idx], chunks, min_row, num_rows, mask_filter{mask}, true)) { return; } @@ -630,6 +833,7 @@ __global__ void __launch_bounds__(decode_block_size) target_pos = min(s->nz_count, src_pos + decode_block_size - out_thread0); if (out_thread0 > 32) { target_pos = min(target_pos, s->dict_pos); } } + // TODO(ets): see if this sync can be removed __syncthreads(); if (t < 32) { // decode repetition and definition levels. @@ -643,9 +847,9 @@ __global__ void __launch_bounds__(decode_block_size) // WARP1: Decode dictionary indices, booleans or string positions if (s->dict_base) { - src_target_pos = gpuDecodeDictionaryIndices(s, sb, src_target_pos, t & 0x1f).first; + src_target_pos = gpuDecodeDictionaryIndices(s, sb, src_target_pos, lane_id).first; } else { - gpuInitStringDescriptors(s, sb, src_target_pos, t & 0x1f); + gpuInitStringDescriptors(s, sb, src_target_pos, lane_id); } if (t == 32) { *(volatile int32_t*)&s->dict_pos = src_target_pos; } } else { @@ -748,6 +952,19 @@ __global__ void __launch_bounds__(decode_block_size) if (t == 0 and s->error != 0) { set_error(s->error, error_code); } } +// Functor used to set the `temp_string_buf` pointer for each page. `data` points to a buffer +// to be used when skipping rows in the delta_byte_array decoder. Given a page and an offset, +// set the page's `temp_string_buf` to be `data + offset`. +struct page_tform_functor { + uint8_t* const data; + + __device__ PageInfo operator()(PageInfo& page, int64_t offset) + { + if (page.temp_string_size != 0) { page.temp_string_buf = data + offset; } + return page; + } +}; + } // anonymous namespace /** @@ -755,20 +972,81 @@ __global__ void __launch_bounds__(decode_block_size) */ void ComputePageStringSizes(cudf::detail::hostdevice_vector& pages, cudf::detail::hostdevice_vector const& chunks, + rmm::device_uvector& temp_string_buf, size_t min_row, size_t num_rows, int level_type_size, + uint32_t kernel_mask, rmm::cuda_stream_view stream) { - dim3 dim_block(preprocess_block_size, 1); - dim3 dim_grid(pages.size(), 1); // 1 threadblock per page + dim3 const dim_block(preprocess_block_size, 1); + dim3 const dim_grid(pages.size(), 1); // 1 threadblock per page if (level_type_size == 1) { - gpuComputePageStringSizes + gpuComputeStringPageBounds <<>>(pages.device_ptr(), chunks, min_row, num_rows); } else { - gpuComputePageStringSizes + gpuComputeStringPageBounds <<>>(pages.device_ptr(), chunks, min_row, num_rows); } + + // kernel mask may contain other kernels we don't need to count + int const count_mask = + kernel_mask & BitOr(decode_kernel_mask::DELTA_BYTE_ARRAY, decode_kernel_mask::STRING); + int const nkernels = std::bitset<32>(count_mask).count(); + auto const streams = cudf::detail::fork_streams(stream, nkernels); + + int s_idx = 0; + if (BitAnd(kernel_mask, decode_kernel_mask::DELTA_BYTE_ARRAY) != 0) { + dim3 dim_delta(delta_preproc_block_size, 1); + gpuComputeDeltaPageStringSizes<<>>( + pages.device_ptr(), chunks, min_row, num_rows); + } + if (BitAnd(kernel_mask, decode_kernel_mask::STRING) != 0) { + gpuComputePageStringSizes<<>>( + pages.device_ptr(), chunks, min_row, num_rows); + } + + // synchronize the streams + cudf::detail::join_streams(streams, stream); + + // check for needed temp space for DELTA_BYTE_ARRAY + auto const need_sizes = thrust::any_of( + rmm::exec_policy(stream), pages.d_begin(), pages.d_end(), [] __device__(auto& page) { + return page.temp_string_size != 0; + }); + + if (need_sizes) { + // sum up all of the temp_string_sizes + auto const page_sizes = [] __device__(PageInfo const& page) { return page.temp_string_size; }; + auto const total_size = thrust::transform_reduce(rmm::exec_policy(stream), + pages.d_begin(), + pages.d_end(), + page_sizes, + 0L, + thrust::plus{}); + + // now do an exclusive scan over the temp_string_sizes to get offsets for each + // page's chunk of the temp buffer + rmm::device_uvector page_string_offsets(pages.size(), stream); + thrust::transform_exclusive_scan(rmm::exec_policy_nosync(stream), + pages.d_begin(), + pages.d_end(), + page_string_offsets.begin(), + page_sizes, + 0L, + thrust::plus{}); + + // allocate the temp space + temp_string_buf.resize(total_size, stream); + + // now use the offsets array to set each page's temp_string_buf pointers + thrust::transform(rmm::exec_policy_nosync(stream), + pages.d_begin(), + pages.d_end(), + page_string_offsets.begin(), + pages.d_begin(), + page_tform_functor{temp_string_buf.data()}); + } } /** diff --git a/cpp/src/io/parquet/parquet_gpu.hpp b/cpp/src/io/parquet/parquet_gpu.hpp index 68851e72663..129d4e4d28c 100644 --- a/cpp/src/io/parquet/parquet_gpu.hpp +++ b/cpp/src/io/parquet/parquet_gpu.hpp @@ -35,6 +35,7 @@ #include +#include #include namespace cudf::io::parquet::detail { @@ -64,7 +65,8 @@ constexpr bool is_supported_encoding(Encoding enc) case Encoding::PLAIN_DICTIONARY: case Encoding::RLE: case Encoding::RLE_DICTIONARY: - case Encoding::DELTA_BINARY_PACKED: return true; + case Encoding::DELTA_BINARY_PACKED: + case Encoding::DELTA_BYTE_ARRAY: return true; default: return false; } } @@ -86,13 +88,15 @@ constexpr void set_error(int32_t error, int32_t* error_code) * These values are used as bitmasks, so they must be powers of 2. */ enum class decode_error : int32_t { - DATA_STREAM_OVERRUN = 0x1, - LEVEL_STREAM_OVERRUN = 0x2, - UNSUPPORTED_ENCODING = 0x4, - INVALID_LEVEL_RUN = 0x8, - INVALID_DATA_TYPE = 0x10, - EMPTY_PAGE = 0x20, - INVALID_DICT_WIDTH = 0x40, + DATA_STREAM_OVERRUN = 0x1, + LEVEL_STREAM_OVERRUN = 0x2, + UNSUPPORTED_ENCODING = 0x4, + INVALID_LEVEL_RUN = 0x8, + INVALID_DATA_TYPE = 0x10, + EMPTY_PAGE = 0x20, + INVALID_DICT_WIDTH = 0x40, + DELTA_PARAM_MISMATCH = 0x80, + DELTA_PARAMS_UNSUPPORTED = 0x100, }; /** @@ -145,6 +149,17 @@ constexpr uint32_t BitAnd(T1 a, T2 b) return static_cast(a) & static_cast(b); } +template ::value and std::is_same_v) or + (is_scoped_enum::value and std::is_same_v) or + (is_scoped_enum::value and std::is_same_v)>* = + nullptr> +constexpr uint32_t BitOr(T1 a, T2 b) +{ + return static_cast(a) | static_cast(b); +} + /** * @brief Enums for the flags in the page header */ @@ -168,10 +183,12 @@ enum level_type { * * Used to control which decode kernels to run. */ -enum kernel_mask_bits { - KERNEL_MASK_GENERAL = (1 << 0), // Run catch-all decode kernel - KERNEL_MASK_STRING = (1 << 1), // Run decode kernel for string data - KERNEL_MASK_DELTA_BINARY = (1 << 2) // Run decode kernel for DELTA_BINARY_PACKED data +enum class decode_kernel_mask { + NONE = 0, + GENERAL = (1 << 0), // Run catch-all decode kernel + STRING = (1 << 1), // Run decode kernel for string data + DELTA_BINARY = (1 << 2), // Run decode kernel for DELTA_BINARY_PACKED data + DELTA_BYTE_ARRAY = (1 << 3) // Run decode kernel for DELTA_BYTE_ARRAY encoded data }; /** @@ -252,9 +269,11 @@ struct PageInfo { int32_t num_input_values; int32_t chunk_row; // starting row of this page relative to the start of the chunk int32_t num_rows; // number of rows in this page - // the next two are calculated in gpuComputePageStringSizes + // the next four are calculated in gpuComputePageStringSizes int32_t num_nulls; // number of null values (V2 header), but recalculated for string cols int32_t num_valids; // number of non-null values, taking into account skip_rows/num_rows + int32_t start_val; // index of first value of the string data stream to use + int32_t end_val; // index of last value in string data stream int32_t chunk_idx; // column chunk this page belongs to int32_t src_col_schema; // schema index of this column uint8_t flags; // PAGEINFO_FLAGS_XXX @@ -291,7 +310,11 @@ struct PageInfo { // level decode buffers uint8_t* lvl_decode_buf[level_type::NUM_LEVEL_TYPES]; - uint32_t kernel_mask; + // temporary space for decoding DELTA_BYTE_ARRAY encoded strings + int64_t temp_string_size; + uint8_t* temp_string_buf; + + decode_kernel_mask kernel_mask; }; /** @@ -597,16 +620,20 @@ void ComputePageSizes(cudf::detail::hostdevice_vector& pages, * * @param[in,out] pages All pages to be decoded * @param[in] chunks All chunks to be decoded + * @param[out] temp_string_buf Temporary space needed for decoding DELTA_BYTE_ARRAY strings * @param[in] min_rows crop all rows below min_row * @param[in] num_rows Maximum number of rows to read * @param[in] level_type_size Size in bytes of the type for level decoding + * @param[in] kernel_mask Mask of kernels to run * @param[in] stream CUDA stream to use */ void ComputePageStringSizes(cudf::detail::hostdevice_vector& pages, cudf::detail::hostdevice_vector const& chunks, + rmm::device_uvector& temp_string_buf, size_t min_row, size_t num_rows, int level_type_size, + uint32_t kernel_mask, rmm::cuda_stream_view stream); /** @@ -665,7 +692,7 @@ void DecodeStringPageData(cudf::detail::hostdevice_vector& pages, * @param[in] min_row Minimum number of rows to read * @param[in] level_type_size Size in bytes of the type for level decoding * @param[out] error_code Error code for kernel failures - * @param[in] stream CUDA stream to use, default 0 + * @param[in] stream CUDA stream to use */ void DecodeDeltaBinary(cudf::detail::hostdevice_vector& pages, cudf::detail::hostdevice_vector const& chunks, @@ -675,6 +702,28 @@ void DecodeDeltaBinary(cudf::detail::hostdevice_vector& pages, int32_t* error_code, rmm::cuda_stream_view stream); +/** + * @brief Launches kernel for reading the DELTA_BYTE_ARRAY column data stored in the pages + * + * The page data will be written to the output pointed to in the page's + * associated column chunk. + * + * @param[in,out] pages All pages to be decoded + * @param[in] chunks All chunks to be decoded + * @param[in] num_rows Total number of rows to read + * @param[in] min_row Minimum number of rows to read + * @param[in] level_type_size Size in bytes of the type for level decoding + * @param[out] error_code Error code for kernel failures + * @param[in] stream CUDA stream to use + */ +void DecodeDeltaByteArray(cudf::detail::hostdevice_vector& pages, + cudf::detail::hostdevice_vector const& chunks, + size_t num_rows, + size_t min_row, + int level_type_size, + int32_t* error_code, + rmm::cuda_stream_view stream); + /** * @brief Launches kernel for initializing encoder row group fragments * diff --git a/cpp/src/io/parquet/reader_impl.cpp b/cpp/src/io/parquet/reader_impl.cpp index 11c20d0e540..6e799424d01 100644 --- a/cpp/src/io/parquet/reader_impl.cpp +++ b/cpp/src/io/parquet/reader_impl.cpp @@ -21,7 +21,6 @@ #include #include #include -#include #include #include @@ -30,10 +29,15 @@ namespace cudf::io::parquet::detail { void reader::impl::decode_page_data(size_t skip_rows, size_t num_rows) { - auto& chunks = _pass_itm_data->chunks; - auto& pages = _pass_itm_data->pages_info; - auto& page_nesting = _pass_itm_data->page_nesting_info; - auto& page_nesting_decode = _pass_itm_data->page_nesting_decode_info; + auto& chunks = _pass_itm_data->chunks; + auto& pages = _pass_itm_data->pages_info; + auto& page_nesting = _pass_itm_data->page_nesting_info; + auto& page_nesting_decode = _pass_itm_data->page_nesting_decode_info; + auto const level_type_size = _pass_itm_data->level_type_size; + + // temporary space for DELTA_BYTE_ARRAY decoding. this only needs to live until + // gpu::DecodeDeltaByteArray returns. + rmm::device_uvector delta_temp_buf(0, _stream); // Should not reach here if there is no page data. CUDF_EXPECTS(pages.size() > 0, "There is no page to decode"); @@ -52,11 +56,12 @@ void reader::impl::decode_page_data(size_t skip_rows, size_t num_rows) // doing a gather operation later on. // TODO: This step is somewhat redundant if size info has already been calculated (nested schema, // chunked reader). - auto const has_strings = (kernel_mask & KERNEL_MASK_STRING) != 0; + auto const has_strings = + (kernel_mask & BitOr(decode_kernel_mask::STRING, decode_kernel_mask::DELTA_BYTE_ARRAY)) != 0; std::vector col_sizes(_input_columns.size(), 0L); if (has_strings) { ComputePageStringSizes( - pages, chunks, skip_rows, num_rows, _pass_itm_data->level_type_size, _stream); + pages, chunks, delta_temp_buf, skip_rows, num_rows, level_type_size, kernel_mask, _stream); col_sizes = calculate_page_string_offsets(); @@ -163,6 +168,7 @@ void reader::impl::decode_page_data(size_t skip_rows, size_t num_rows) chunks.host_to_device_async(_stream); chunk_nested_valids.host_to_device_async(_stream); chunk_nested_data.host_to_device_async(_stream); + if (has_strings) { chunk_nested_str_data.host_to_device_async(_stream); } // create this before we fork streams kernel_error error_code(_stream); @@ -171,25 +177,27 @@ void reader::impl::decode_page_data(size_t skip_rows, size_t num_rows) int const nkernels = std::bitset<32>(kernel_mask).count(); auto streams = cudf::detail::fork_streams(_stream, nkernels); - auto const level_type_size = _pass_itm_data->level_type_size; - // launch string decoder int s_idx = 0; - if (has_strings) { - auto& stream = streams[s_idx++]; - chunk_nested_str_data.host_to_device_async(stream); + if (BitAnd(kernel_mask, decode_kernel_mask::STRING) != 0) { DecodeStringPageData( - pages, chunks, num_rows, skip_rows, level_type_size, error_code.data(), stream); + pages, chunks, num_rows, skip_rows, level_type_size, error_code.data(), streams[s_idx++]); + } + + // launch delta byte array decoder + if (BitAnd(kernel_mask, decode_kernel_mask::DELTA_BYTE_ARRAY) != 0) { + DecodeDeltaByteArray( + pages, chunks, num_rows, skip_rows, level_type_size, error_code.data(), streams[s_idx++]); } // launch delta binary decoder - if ((kernel_mask & KERNEL_MASK_DELTA_BINARY) != 0) { + if (BitAnd(kernel_mask, decode_kernel_mask::DELTA_BINARY) != 0) { DecodeDeltaBinary( pages, chunks, num_rows, skip_rows, level_type_size, error_code.data(), streams[s_idx++]); } // launch the catch-all page decoder - if ((kernel_mask & KERNEL_MASK_GENERAL) != 0) { + if (BitAnd(kernel_mask, decode_kernel_mask::GENERAL) != 0) { DecodePageData( pages, chunks, num_rows, skip_rows, level_type_size, error_code.data(), streams[s_idx++]); } diff --git a/cpp/src/io/parquet/reader_impl_preprocess.cu b/cpp/src/io/parquet/reader_impl_preprocess.cu index 80a4d00a5a2..0bc492546e9 100644 --- a/cpp/src/io/parquet/reader_impl_preprocess.cu +++ b/cpp/src/io/parquet/reader_impl_preprocess.cu @@ -1416,7 +1416,7 @@ std::vector reader::impl::calculate_page_string_offsets() page_index.begin(), page_to_string_size{pages.device_ptr(), chunks.device_ptr()}); // do scan by key to calculate string offsets for each page - thrust::exclusive_scan_by_key(rmm::exec_policy(_stream), + thrust::exclusive_scan_by_key(rmm::exec_policy_nosync(_stream), page_keys.begin(), page_keys.end(), val_iter, @@ -1424,7 +1424,7 @@ std::vector reader::impl::calculate_page_string_offsets() // now sum up page sizes rmm::device_uvector reduce_keys(col_sizes.size(), _stream); - thrust::reduce_by_key(rmm::exec_policy(_stream), + thrust::reduce_by_key(rmm::exec_policy_nosync(_stream), page_keys.begin(), page_keys.end(), val_iter, diff --git a/python/cudf/cudf/tests/data/parquet/delta_byte_arr.parquet b/python/cudf/cudf/tests/data/parquet/delta_byte_arr.parquet new file mode 100644 index 0000000000000000000000000000000000000000..7f6006a75bff0498e373d61f1265ff2e8458c917 GIT binary patch literal 5783 zcmcII2V4``(;E^(DAHut1UN+xLQyH9NIw=L9WDW-1Uw8R0VI)_glYp74;un177(#e z&J(bKD5wY)5JcrpKtWMdEC^Ttm9zYJ7lNGU-TlAc_kEjRcIVB!nKy4{cHYb~U_UdE z2-e7fXJkS_4hSHBWsm>@YT7}n@c6~>rDG6z0t5SuVE_Q&;0!8EJ_7DL2*80fb*P}w zw)j)`;!Xt*u6&`|i}u2Fc{z85&!0o*d-R73?#eZWXB;>bf2^`g3SGt^N3An!Yfq@j z)~58>9NjXK(j#lBbH6?KH8)2uj2EobJjd%?^YKSb*rQ=*aLx)j9!43TfIaZ$P1zo)Qrj5hs?!3b>-MXbz|ZaiA&^REejO?EkxMA~wp$ zgH@P~2$+qOP#akh4DkNZA~K#XGDd8HKhMD7Wyb6wAONj$+PaxG-(C6Rw}fN2Z4%2R zAC|At>mA)(xoyga(l?J7-ZDlcH3d3xb>*Iqx4V|K-9v!kgzoCkZ{ozhmSW$PgS~kJ zS@!bIu1=BL&i2Oqjs-XcfX9Ei zZiOfJ{t4x(X$y!3%YzyPnRh8d8L26hBErOMew>6YiH=>0S>ib!T8F3JAlUzTH^`Rj zWWhSZ*I^+S7|rKLvjt>%n9ogQFXeDWlfs~a1W-Yj*^CW)lug;IR6_T<@{^*PK#T00 z$LcdLEiNzMOD7D!8pb`;kZ+Sd8pVr<+)As*Xv- z_{8d+j6CHBaqC>%&rVrRC>iCIQd2wg6dE)u)>F)1rJRbUWnEenW4|c&J}&kxlU_zy zz0>&0ju$U>CSMRYdcDX{t6SkhKP$WWS^9U1(+;Ry5y-b~iJWC#&(D20>eg`P*vVs0 zu2j89pVb<2qI!hfF8*o2knB7suUXhW{fS=O!!paQbMK5aZ=T(&9^6$yl>u-En^V*$ z*`F6J5=+9CvIShs{`o{GUzS>lgyJ_%i}>M;EZBqj;m@Hb8yFQW;gX3kKb$L)pz-(i zOZE@MNy%U83zos1X~ zg;M8}pL~dP#2*&GtGPEW#s2s+qr9v6aq$+t>2#rj6N6zHHHTqYkXlqIcw-vx1ds}x z$PhVVvs0npGV20+g6?$1$Hy3CFKg{xGVCDImL!vXx{W5B`KmJ{Tbmy@osZsHp7rk2 zN8p63=*;)V3*;*5xV)vapB~tEwKDJ6b&u=T@9$4tkrWeGJ-}X-k@)`ntpLZjosPSs zPHwM=+S;#8JL}gfHh9bP{H3!%&*q1-+wA@rQ12b-&=S9o-gaK0`N#_6?9w+OT?xFC zR}{|a`;`h>uh$MD%QzDQ4!#XLczWOB*RGP1ma8KhH)!Di6Tm>~6vc=zK6|oHGKo-? zEcLQ7`c8apX2c;SIB*vZeKG#*5L_@B=lfM2=*IyZd zP#2E+R0CCZTm?kOjU@1nEUF`MQobS@Xasopxq0~Nc=&t6zI6XNItz{cjmdJbw=?YP z0VC_;PkjlZVFe)P4udB2%2xiMr6$~QUlEWey(Hsm6?f%U`V7pSnU(3#hI+g*Ms>U2yUlYy_n_!xO`o=^Q>y=Z0+H&J zt{T1A;D(M2ZY@a;G!0EGxUgo;lqXJ3swAgZwdr|c<&<9P{%SeNTN`EN%`OOvaL4-B?tLHh z+oFxjx5X*i;(u&-x<$})%3lASKfkazPqzBd6Yd%xi)^P$?;7mu-+A5C&yW7NqF!Sq zjeaIZOnJ2Hb%o#1FI5KT*sA9gHafoQfnF4M9Nrge{4B{PZ7{VgY1FVX+4-l8U2RPr zhE`Lkv}tm5gWHuB1*#inj#O~AX1M`v$y(pw;^68gGQ*k(cjn+jvUUw*kQ!P z0&V;zT~iJT`q;?mu|t3k1yHkR*jsS`01?#&f_18>Fd0l32|Md#40Og*zcG^e1?;6x z5pEosh1C+4$2_`(W*iYKAkUgO6O6IsDSl7%_| zYoV^Hom0#$9DXD~pIN^EoR7;@_FYfBI}&i*&co2|YX3f69HBe4&nmQC5cFHku_zA$ zfAOZY2c{~!--i@=5O4!PAE2y0bzfSAVcP7x?Y^6yXPI0ax+Cv0gAxG0t|{x6XOYi^ z7%?4QA2jyV)w5VzaKt7lzV4c672CV7Pt(`x)Wd`OOODMR?%Qyw*mOaKzBJ7JU(=U|FI#`ujp&d*>v%%30_$DN6{B{$ow|35sYN^)cyQaGf#N*1gwkZ| z3ynOL$l(T4rgitJ{?Pi$gQeM3_inS}tkN6k589iiw<^@z+1&a**A_gj?2}fUy(jeZ znWCwzJckz+v|{{r(lqZ3Sp`#0CSA#zwm^6kUb9wajf^HK1@aai-J$NTxaD#RPEpgI z0HoZAve2Zc$M*D(@+z=!V}?i9Yqv z-NEyZ)TWLH1LxiH_?LCdJSu7iJ!i^x{__0VaEt2ufg=Oj>s99G@y`jJ&HQDO0`_XX zI#Z~9e-a( zQX#TRb)OFzq(S!Gh;hOek7hgFT>D(G_eUZ0V7MhdF1jS4+L?b568j7{DCg>BpRP?$ zA(YSS(@ZlfEsr_umeYoRI2A5n5{qUr9qtK6Yv&Gp(we$xu*P!7QQhm#N3^Dztu#Q~ z|Bz{`Pag5|dcb+}_^#Zi>%m=bY#L7AOL~}+zPiz>KSz5(R(t2gIqyPPyKZjxa}d#g ztBLJvmz}%IDlk`k_Khd_eW8-u*%oWXdWwF3({M|kQs0na>*nUr>!TIBViHfhoa<{9 z*xXRaSa#Fp!obD)Z7L_MU6l0Gw1yUJJ(5bz*?BrJsO3DrDEZMX*{2p4tZi+c55#p6 zJA3j6$trfYzW=OGuUm0QvQo$JM(-=5X?r(T8SVuG4qs`vJ9W$0Xi?9NH|<*3UK1^nYrz@zBYn+_bwPT@FW9 zI$Uh*Rmkl3{nZ`%HLRv{aHvjrHN)TGb#M4pp(a;7xve~)INfe^lE~JXHI`p7D(=?+~!x7 zkTavtH?)XrIZp;4C~e2-NmCS(7Lh~~%@1R9I3liCjHSgr)1dY0)Dj(gtN*i%grvMZ zc<{zwS|B-^g(c0(!kUZ$<}@=Tc}W|degbgZo2r*-_^djnvdD9lVtyXku{1H{*NnZNv*TCf?6wVgRm1Avb$cke zTX@f!Wi;PlRzO$%iBpjunRYiKROM(K^%v(XMh3Y zV2KX6UVR+SumDRiksyI`#9fIk62&Hji6f$rw**v2T&^ecRj;cDcxxltV3a?#eqRy3 z^dyrBlk4Ki_-_}ZlV^Ng&`hRD3y;aTuix~^v~NW4I2qzlv?+Pib9ilOwEwR=eJI}n zgrK9=;7V$VvPS^+CHcml~ZK{Is(*B zQzzt1_6R2CpE?!j5h`QE2DAXE7oiCm-}W9R?HjEG^Z^xe*cO}xP*bK)2>IH<82kUE zVuiLlB66E;!36;I>5K`v(%``4O{z+vBaVn)3hYod%4ieQzm6M>?f>C~fa*|f>+HY- zIDlF?Z({v##3rV~KW10x4l3-CJ!lI+n}0yFlr~uj_VfEO(01%yCU?l1pjTp~3-aZS zIsP5a3x*S5mq^$XMIlcZb{?lNkueN=pd$n&eIlb6^v6{9LAj92*_kPo2g5>?o?wllmjuxQB*B4vaS+WlhUMlgucv$()E!|CeKFb*j9thY@*pgo~CS9nl$A1~NE0JVbvvBy%B&S9Yb`}_T~PDc}d`?Cc6q%XYa%5;qW_x_{9lldcPu3}D* zUlhlUA0LkT6b%ArCK@e2!m#yth()4fL{)5p)Sk}HekgM6LIMk!g6-EA#US=s!U^&< z4X2s$BltE#7B> + lists_per_row = 3 + list_size = 4 + num_rows = nrows + include_validity = add_nulls + + def list_gen_wrapped(x, y): + return list_row_gen( + int_gen, x * list_size * lists_per_row, list_size, lists_per_row + ) + + def string_list_gen_wrapped(x, y): + return list_row_gen( + string_gen, + x * list_size * lists_per_row, + list_size, + lists_per_row, + include_validity, + ) + + data = struct_gen( + [int_gen, string_gen, list_gen_wrapped, string_list_gen_wrapped], + 0, + num_rows, + include_validity, + ) + test_pdf = pa.Table.from_pydict({"sol": data}).to_pandas() + pdf_fname = tmpdir.join("pdfdeltaba.parquet") + test_pdf.to_parquet( + pdf_fname, + version="2.6", + column_encoding={ + "sol.col0": "DELTA_BINARY_PACKED", + "sol.col1": str_encoding, + "sol.col2.list.element.list.element": "DELTA_BINARY_PACKED", + "sol.col3.list.element.list.element": str_encoding, + }, + data_page_version="2.0", + data_page_size=64 * 1024, + engine="pyarrow", + use_dictionary=False, + ) + # sanity check to verify file is written properly + assert_eq(test_pdf, pd.read_parquet(pdf_fname)) + cdf = cudf.read_parquet(pdf_fname) + assert_eq(cdf, cudf.from_pandas(test_pdf)) + + @pytest.mark.parametrize( "data", [ From d2069f45cc518c78ca901ea65370eff5c474cfd1 Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Fri, 17 Nov 2023 16:11:22 +0000 Subject: [PATCH 23/44] Match pandas join ordering obligations in pandas-compatible mode (#14428) If we pass sort=True to merges we are on the hook to sort the result in order with respect to the key columns. If those key columns have repeated values there is still some space for ambiguity. Currently we get a result back whose order (for the repeated key values) is determined by the gather map that libcudf returns for the join. This does not come with any ordering guarantees. When sort=False, pandas has join-type dependent ordering guarantees which we also do not match. To fix this, in pandas-compatible mode only, reorder the gather maps according to the order of the input keys. When sort=False this means that our result matches pandas ordering. When sort=True, it ensures that (if we use a stable sort) the tie-break for equal sort keys is the input dataframe order. While we're here, switch from argsort + gather to sort_by_key when sorting results. - Closes #14001 Authors: - Lawrence Mitchell (https://github.com/wence-) Approvers: - Ashwin Srinath (https://github.com/shwina) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14428 --- python/cudf/cudf/core/_compat.py | 1 + python/cudf/cudf/core/join/join.py | 125 +++++++++-- python/cudf/cudf/tests/test_join_order.py | 261 ++++++++++++++++++++++ 3 files changed, 373 insertions(+), 14 deletions(-) create mode 100644 python/cudf/cudf/tests/test_join_order.py diff --git a/python/cudf/cudf/core/_compat.py b/python/cudf/cudf/core/_compat.py index 888b94e070c..e257b7a1fa1 100644 --- a/python/cudf/cudf/core/_compat.py +++ b/python/cudf/cudf/core/_compat.py @@ -11,3 +11,4 @@ PANDAS_LT_153 = PANDAS_VERSION < version.parse("1.5.3") PANDAS_GE_200 = PANDAS_VERSION >= version.parse("2.0.0") PANDAS_GE_210 = PANDAS_VERSION >= version.parse("2.1.0") +PANDAS_GE_220 = PANDAS_VERSION >= version.parse("2.2.0") diff --git a/python/cudf/cudf/core/join/join.py b/python/cudf/cudf/core/join/join.py index b94f8f583f4..20f5b7989eb 100644 --- a/python/cudf/cudf/core/join/join.py +++ b/python/cudf/cudf/core/join/join.py @@ -1,11 +1,13 @@ # Copyright (c) 2020-2023, NVIDIA CORPORATION. from __future__ import annotations +import itertools import warnings from typing import Any, ClassVar, List, Optional import cudf from cudf import _lib as libcudf +from cudf._lib.types import size_type_dtype from cudf.core.copy_types import GatherMap from cudf.core.join._join_helpers import ( _coerce_to_tuple, @@ -94,7 +96,44 @@ def __init__( self.lhs = lhs.copy(deep=False) self.rhs = rhs.copy(deep=False) self.how = how - self.sort = sort + # If the user requests that the result is sorted or we're in + # pandas-compatible mode we have various obligations on the + # output order: + # + # compat-> | False | True + # sort | | + # ---------+--------------------------+------------------------------- + # False| no obligation | ordering as per pandas docs(*) + # True | sorted lexicographically | sorted lexicographically(*) + # + # (*) If two keys are equal, tiebreak is to use input table order. + # + # In pandas-compat mode, we have obligations on the order to + # match pandas (even if sort=False), see + # pandas.pydata.org/docs/reference/api/pandas.DataFrame.merge.html. + # The ordering requirements differ depending on which join + # type is specified: + # + # - left: preserve key order (only keeping left keys) + # - right: preserve key order (only keeping right keys) + # - inner: preserve key order (of left keys) + # - outer: sort keys lexicographically + # - cross (not supported): preserve key order (of left keys) + # + # Moreover, in all cases, whenever there is a tiebreak + # situation (for sorting or otherwise), the deciding order is + # "input table order" + self.sort = sort or ( + cudf.get_option("mode.pandas_compatible") and how == "outer" + ) + self.preserve_key_order = cudf.get_option( + "mode.pandas_compatible" + ) and how in { + "inner", + "outer", + "left", + "right", + } self.lsuffix, self.rsuffix = suffixes # At this point validation guarantees that if on is not None we @@ -160,6 +199,55 @@ def __init__( } ) + def _gather_maps(self, left_cols, right_cols): + # Produce gather maps for the join, optionally reordering to + # match pandas-order in compat mode. + maps = self._joiner( + left_cols, + right_cols, + how=self.how, + ) + if not self.preserve_key_order: + return maps + # We should only get here if we're in a join on which + # pandas-compat places some ordering obligation (which + # precludes a semi-join) + # We must perform this reordering even if sort=True since the + # obligation to ensure tiebreaks appear in input table order + # means that the gather maps must be permuted into an original + # order. + assert self.how in {"inner", "outer", "left", "right"} + # And hence both maps returned from the libcudf join should be + # non-None. + assert all(m is not None for m in maps) + lengths = [len(left_cols[0]), len(right_cols[0])] + # Only nullify those maps that need it. + nullify = [ + self.how not in {"inner", "left"}, + self.how not in {"inner", "right"}, + ] + # To reorder maps so that they are in order of the input + # tables, we gather from iota on both right and left, and then + # sort the gather maps with those two columns as key. + key_order = list( + itertools.chain.from_iterable( + libcudf.copying.gather( + [cudf.core.column.arange(n, dtype=size_type_dtype)], + map_, + nullify=null, + ) + for map_, n, null in zip(maps, lengths, nullify) + ) + ) + return libcudf.sort.sort_by_key( + list(maps), + # If how is right, right map is primary sort key. + key_order[:: -1 if self.how == "right" else 1], + [True] * len(key_order), + ["last"] * len(key_order), + stable=True, + ) + def perform_merge(self) -> cudf.DataFrame: left_join_cols = [] right_join_cols = [] @@ -184,12 +272,9 @@ def perform_merge(self) -> cudf.DataFrame: left_key.set(self.lhs, lcol_casted, validate=False) right_key.set(self.rhs, rcol_casted, validate=False) - left_rows, right_rows = self._joiner( - left_join_cols, - right_join_cols, - how=self.how, + left_rows, right_rows = self._gather_maps( + left_join_cols, right_join_cols ) - gather_kwargs = { "keep_index": self._using_left_index or self._using_right_index, } @@ -305,6 +390,11 @@ def _sort_result(self, result: cudf.DataFrame) -> cudf.DataFrame: # same order as given in 'on'. If the indices are used as # keys, the index will be sorted. If one index is specified, # the key columns on the other side will be used to sort. + # In pandas-compatible mode, tie-breaking for multiple equal + # sort keys is to produce output in input dataframe order. + # This is taken care of by using a stable sort here, and (in + # pandas-compat mode) reordering the gather maps before + # producing the input result. by: List[Any] = [] if self._using_left_index and self._using_right_index: by.extend(result._index._data.columns) @@ -313,15 +403,22 @@ def _sort_result(self, result: cudf.DataFrame) -> cudf.DataFrame: if not self._using_right_index: by.extend([result._data[col.name] for col in self._right_keys]) if by: - to_sort = cudf.DataFrame._from_data(dict(enumerate(by))) - sort_order = GatherMap.from_column_unchecked( - cudf.core.column.as_column(to_sort.argsort()), - len(result), - nullify=False, + keep_index = self._using_left_index or self._using_right_index + if keep_index: + to_sort = [*result._index._columns, *result._columns] + index_names = result._index.names + else: + to_sort = [*result._columns] + index_names = None + result_columns = libcudf.sort.sort_by_key( + to_sort, + by, + [True] * len(by), + ["last"] * len(by), + stable=True, ) - result = result._gather( - sort_order, - keep_index=self._using_left_index or self._using_right_index, + result = result._from_columns_like_self( + result_columns, result._column_names, index_names ) return result diff --git a/python/cudf/cudf/tests/test_join_order.py b/python/cudf/cudf/tests/test_join_order.py new file mode 100644 index 00000000000..61a2ed239cb --- /dev/null +++ b/python/cudf/cudf/tests/test_join_order.py @@ -0,0 +1,261 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. + +import itertools +import operator +import string +from collections import defaultdict + +import numpy as np +import pytest + +import cudf +from cudf.core._compat import PANDAS_GE_220 +from cudf.testing._utils import assert_eq + + +@pytest.fixture(params=[False, True], ids=["unsorted", "sorted"]) +def sort(request): + return request.param + + +@pytest.fixture +def left(): + left_key = [1, 3, 2, 1, 1, 2, 5, 1, 4, 5, 8, 12, 12312, 1] * 100 + left_val = list(range(len(left_key))) + return cudf.DataFrame({"key": left_key, "val": left_val}) + + +@pytest.fixture +def right(): + right_key = [12312, 12312, 3, 2, 1, 1, 5, 7, 2] * 200 + right_val = list( + itertools.islice(itertools.cycle(string.ascii_letters), len(right_key)) + ) + return cudf.DataFrame({"key": right_key, "val": right_val}) + + +if PANDAS_GE_220: + # Behaviour in sort=False case didn't match documentation in many + # cases prior to https://github.com/pandas-dev/pandas/pull/54611 + # (released as part of pandas 2.2) + def expected(left, right, sort, *, how): + left = left.to_pandas() + right = right.to_pandas() + return left.merge(right, on="key", how=how, sort=sort) + +else: + + def expect_inner(left, right, sort): + left_key = left.key.values_host.tolist() + left_val = left.val.values_host.tolist() + right_key = right.key.values_host.tolist() + right_val = right.val.values_host.tolist() + + right_have = defaultdict(list) + for i, k in enumerate(right_key): + right_have[k].append(i) + keys = [] + val_x = [] + val_y = [] + for k, v in zip(left_key, left_val): + if k not in right_have: + continue + for i in right_have[k]: + keys.append(k) + val_x.append(v) + val_y.append(right_val[i]) + + if sort: + # Python sort is stable, so this will preserve input order for + # equal items. + keys, val_x, val_y = zip( + *sorted(zip(keys, val_x, val_y), key=operator.itemgetter(0)) + ) + return cudf.DataFrame({"key": keys, "val_x": val_x, "val_y": val_y}) + + def expect_left(left, right, sort): + left_key = left.key.values_host.tolist() + left_val = left.val.values_host.tolist() + right_key = right.key.values_host.tolist() + right_val = right.val.values_host.tolist() + + right_have = defaultdict(list) + for i, k in enumerate(right_key): + right_have[k].append(i) + keys = [] + val_x = [] + val_y = [] + for k, v in zip(left_key, left_val): + if k not in right_have: + right_vals = [None] + else: + right_vals = [right_val[i] for i in right_have[k]] + + for rv in right_vals: + keys.append(k) + val_x.append(v) + val_y.append(rv) + + if sort: + # Python sort is stable, so this will preserve input order for + # equal items. + keys, val_x, val_y = zip( + *sorted(zip(keys, val_x, val_y), key=operator.itemgetter(0)) + ) + return cudf.DataFrame({"key": keys, "val_x": val_x, "val_y": val_y}) + + def expect_outer(left, right, sort): + left_key = left.key.values_host.tolist() + left_val = left.val.values_host.tolist() + right_key = right.key.values_host.tolist() + right_val = right.val.values_host.tolist() + right_have = defaultdict(list) + for i, k in enumerate(right_key): + right_have[k].append(i) + keys = [] + val_x = [] + val_y = [] + for k, v in zip(left_key, left_val): + if k not in right_have: + right_vals = [None] + else: + right_vals = [right_val[i] for i in right_have[k]] + for rv in right_vals: + keys.append(k) + val_x.append(v) + val_y.append(rv) + left_have = set(left_key) + for k, v in zip(right_key, right_val): + if k not in left_have: + keys.append(k) + val_x.append(None) + val_y.append(v) + + # Python sort is stable, so this will preserve input order for + # equal items. + # outer joins are always sorted, but we test both sort values + keys, val_x, val_y = zip( + *sorted(zip(keys, val_x, val_y), key=operator.itemgetter(0)) + ) + return cudf.DataFrame({"key": keys, "val_x": val_x, "val_y": val_y}) + + def expected(left, right, sort, *, how): + if how == "inner": + return expect_inner(left, right, sort) + elif how == "outer": + return expect_outer(left, right, sort) + elif how == "left": + return expect_left(left, right, sort) + elif how == "right": + return expect_left(right, left, sort).rename( + {"val_x": "val_y", "val_y": "val_x"}, axis=1 + ) + else: + raise NotImplementedError() + + +@pytest.mark.parametrize("how", ["inner", "left", "right", "outer"]) +def test_join_ordering_pandas_compat(left, right, sort, how): + with cudf.option_context("mode.pandas_compatible", True): + actual = left.merge(right, on="key", how=how, sort=sort) + expect = expected(left, right, sort, how=how) + assert_eq(expect, actual) + + +@pytest.mark.parametrize("how", ["left", "right", "inner", "outer"]) +@pytest.mark.parametrize("sort", [True, False]) +@pytest.mark.parametrize("on_index", [True, False]) +@pytest.mark.parametrize("left_unique", [True, False]) +@pytest.mark.parametrize("left_monotonic", [True, False]) +@pytest.mark.parametrize("right_unique", [True, False]) +@pytest.mark.parametrize("right_monotonic", [True, False]) +def test_merge_combinations( + request, + how, + sort, + on_index, + left_unique, + left_monotonic, + right_unique, + right_monotonic, +): + request.applymarker( + pytest.mark.xfail( + condition=how == "outer" + and on_index + and left_unique + and not left_monotonic + and right_unique + and not right_monotonic, + reason="https://github.com/pandas-dev/pandas/issues/55992", + ) + ) + left = [2, 3] + if left_unique: + left.append(4 if left_monotonic else 1) + else: + left.append(3 if left_monotonic else 2) + + right = [2, 3] + if right_unique: + right.append(4 if right_monotonic else 1) + else: + right.append(3 if right_monotonic else 2) + + left = cudf.DataFrame({"key": left}) + right = cudf.DataFrame({"key": right}) + + if on_index: + left = left.set_index("key") + right = right.set_index("key") + on_kwargs = {"left_index": True, "right_index": True} + else: + on_kwargs = {"on": "key"} + + with cudf.option_context("mode.pandas_compatible", True): + result = cudf.merge(left, right, how=how, sort=sort, **on_kwargs) + if on_index: + left = left.reset_index() + right = right.reset_index() + + if how in ["left", "right", "inner"]: + if how in ["left", "inner"]: + expected, other, other_unique = left, right, right_unique + else: + expected, other, other_unique = right, left, left_unique + if how == "inner": + keep_values = set(left["key"].values_host).intersection( + right["key"].values_host + ) + keep_mask = expected["key"].isin(keep_values) + expected = expected[keep_mask] + if sort: + expected = expected.sort_values("key") + if not other_unique: + other_value_counts = other["key"].value_counts() + repeats = other_value_counts.reindex( + expected["key"].values, fill_value=1 + ) + repeats = repeats.astype(np.intp) + expected = expected["key"].repeat(repeats.values) + expected = expected.to_frame() + elif how == "outer": + if on_index and left_unique and left["key"].equals(right["key"]): + expected = cudf.DataFrame({"key": left["key"]}) + else: + left_counts = left["key"].value_counts() + right_counts = right["key"].value_counts() + expected_counts = left_counts.mul(right_counts, fill_value=1) + expected_counts = expected_counts.astype(np.intp) + expected = expected_counts.index.values_host.repeat( + expected_counts.values_host + ) + expected = cudf.DataFrame({"key": expected}) + expected = expected.sort_values("key") + + if on_index: + expected = expected.set_index("key") + else: + expected = expected.reset_index(drop=True) + + assert_eq(result, expected) From ba5ec4080be38b795053d11bf46cb3688c201893 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 17 Nov 2023 10:36:30 -0600 Subject: [PATCH 24/44] Enable build concurrency for nightly and merge triggers. (#14441) --- .github/workflows/build.yaml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index 2539057c105..e27361ab263 100644 --- a/.github/workflows/build.yaml +++ b/.github/workflows/build.yaml @@ -22,7 +22,7 @@ on: default: nightly concurrency: - group: ${{ github.workflow }}-${{ github.ref }} + group: ${{ github.workflow }}-${{ github.ref }}-${{ github.event_name }} cancel-in-progress: true jobs: From 6c2e972cefff05f6ffbba4fd9ba894e6849b041e Mon Sep 17 00:00:00 2001 From: Trent Nelson Date: Fri, 17 Nov 2023 13:29:23 -0800 Subject: [PATCH 25/44] Implement user_datasource_wrapper is_empty() and is_device_read_preferred(). (#14357) These two routines are missing from the current `user_datasource_wrapper` impl. Authors: - Trent Nelson (https://github.com/tpn) - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Mike Wilson (https://github.com/hyperbolic2346) - Vukasin Milovanovic (https://github.com/vuule) URL: https://github.com/rapidsai/cudf/pull/14357 --- cpp/src/io/utilities/datasource.cpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/cpp/src/io/utilities/datasource.cpp b/cpp/src/io/utilities/datasource.cpp index 5cdd92ce3b7..a466ef84133 100644 --- a/cpp/src/io/utilities/datasource.cpp +++ b/cpp/src/io/utilities/datasource.cpp @@ -360,6 +360,11 @@ class user_datasource_wrapper : public datasource { return source->supports_device_read(); } + [[nodiscard]] bool is_device_read_preferred(size_t size) const override + { + return source->is_device_read_preferred(size); + } + size_t device_read(size_t offset, size_t size, uint8_t* dst, @@ -385,6 +390,8 @@ class user_datasource_wrapper : public datasource { [[nodiscard]] size_t size() const override { return source->size(); } + [[nodiscard]] bool is_empty() const override { return source->is_empty(); } + private: datasource* const source; ///< A non-owning pointer to the user-implemented datasource }; From 10218a972d3a4950d4639c41d9ae9116aff18c12 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Fri, 17 Nov 2023 14:14:55 -0800 Subject: [PATCH 26/44] JSON writer: avoid default stream use in `string_scalar` constructors (#14444) Added the true/false string scalars to `column_to_strings_fn` so they are created once, instead of creating new scalars for each boolean column (using default stream). Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Nghia Truong (https://github.com/ttnghia) - https://github.com/shrshi URL: https://github.com/rapidsai/cudf/pull/14444 --- cpp/src/io/json/write_json.cu | 10 +++++++--- 1 file changed, 7 insertions(+), 3 deletions(-) diff --git a/cpp/src/io/json/write_json.cu b/cpp/src/io/json/write_json.cu index c211d17f13a..938f9728fe8 100644 --- a/cpp/src/io/json/write_json.cu +++ b/cpp/src/io/json/write_json.cu @@ -500,7 +500,9 @@ struct column_to_strings_fn { struct_row_end_wrap("}", true, stream), list_value_separator(",", true, stream), list_row_begin_wrap("[", true, stream), - list_row_end_wrap("]", true, stream) + list_row_end_wrap("]", true, stream), + true_value(options_.get_true_value(), true, stream), + false_value(options_.get_false_value(), true, stream) { } @@ -526,8 +528,7 @@ struct column_to_strings_fn { std::enable_if_t, std::unique_ptr> operator()( column_view const& column) const { - return cudf::strings::detail::from_booleans( - column, options_.get_true_value(), options_.get_false_value(), stream_, mr_); + return cudf::strings::detail::from_booleans(column, true_value, false_value, stream_, mr_); } // strings: @@ -742,6 +743,9 @@ struct column_to_strings_fn { string_scalar const list_value_separator; // "," string_scalar const list_row_begin_wrap; // "[" string_scalar const list_row_end_wrap; // "]" + // bool converter constants + string_scalar const true_value; + string_scalar const false_value; }; } // namespace From 723c565f7a03e3e9a842526cd4cc94bcf6f582e5 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Fri, 17 Nov 2023 17:37:47 -0800 Subject: [PATCH 27/44] Fix intermediate type checking in expression parsing (#14445) When parsing expressions, device data references are reused if there are multiple that are identical. Equality is determined by comparing the fields of the reference, but previously the data type was omitted. For column and literal references, this is OK because the `data_index` uniquely identifies the reference. For intermediates, however, the index is not sufficient to disambiguate because an expression could reuse a given location even if the operation produces a different data type. Therefore, the data type must be part of the equality operator. Resolves #14409 Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - David Wendt (https://github.com/davidwendt) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14445 --- .../cudf/ast/detail/expression_parser.hpp | 4 +-- cpp/tests/ast/transform_tests.cpp | 27 +++++++++++++++++++ 2 files changed, 29 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/ast/detail/expression_parser.hpp b/cpp/include/cudf/ast/detail/expression_parser.hpp index db0abe435b0..a36a831a7aa 100644 --- a/cpp/include/cudf/ast/detail/expression_parser.hpp +++ b/cpp/include/cudf/ast/detail/expression_parser.hpp @@ -67,8 +67,8 @@ struct alignas(8) device_data_reference { bool operator==(device_data_reference const& rhs) const { - return std::tie(data_index, reference_type, table_source) == - std::tie(rhs.data_index, rhs.reference_type, rhs.table_source); + return std::tie(data_index, data_type, reference_type, table_source) == + std::tie(rhs.data_index, rhs.data_type, rhs.reference_type, rhs.table_source); } }; diff --git a/cpp/tests/ast/transform_tests.cpp b/cpp/tests/ast/transform_tests.cpp index c0109a40cec..624a781c5b9 100644 --- a/cpp/tests/ast/transform_tests.cpp +++ b/cpp/tests/ast/transform_tests.cpp @@ -316,6 +316,33 @@ TEST_F(TransformTest, ImbalancedTreeArithmetic) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity); } +TEST_F(TransformTest, ImbalancedTreeArithmeticDeep) +{ + auto c_0 = column_wrapper{4, 5, 6}; + auto table = cudf::table_view{{c_0}}; + + auto col_ref_0 = cudf::ast::column_reference(0); + + // expression: (c0 < c0) == (c0 < (c0 + c0)) + // {false, false, false} == (c0 < {8, 10, 12}) + // {false, false, false} == {true, true, true} + // {false, false, false} + auto expression_left_subtree = + cudf::ast::operation(cudf::ast::ast_operator::LESS, col_ref_0, col_ref_0); + auto expression_right_inner_subtree = + cudf::ast::operation(cudf::ast::ast_operator::ADD, col_ref_0, col_ref_0); + auto expression_right_subtree = + cudf::ast::operation(cudf::ast::ast_operator::LESS, col_ref_0, expression_right_inner_subtree); + + auto expression_tree = cudf::ast::operation( + cudf::ast::ast_operator::EQUAL, expression_left_subtree, expression_right_subtree); + + auto result = cudf::compute_column(table, expression_tree); + auto expected = column_wrapper{false, false, false}; + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity); +} + TEST_F(TransformTest, MultiLevelTreeComparator) { auto c_0 = column_wrapper{3, 20, 1, 50}; From 2afb784a24b3ccb50e15f664117f87ec77aa5272 Mon Sep 17 00:00:00 2001 From: Matthew Roeschke <10647082+mroeschke@users.noreply.github.com> Date: Mon, 20 Nov 2023 06:32:30 -1000 Subject: [PATCH 28/44] REF: Remove instances of pd.core (#14421) `pandas.core` is technically private and methods could be moved at any time. Avoiding places in the codepace where they could be avoided Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Bradley Dice (https://github.com/bdice) - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14421 --- python/cudf/cudf/core/column/datetime.py | 6 +++++- python/cudf/cudf/core/dataframe.py | 15 +++++++++++++-- python/cudf/cudf/core/index.py | 10 ++++------ python/cudf/cudf/tests/test_rolling.py | 9 ++++----- 4 files changed, 26 insertions(+), 14 deletions(-) diff --git a/python/cudf/cudf/core/column/datetime.py b/python/cudf/cudf/core/column/datetime.py index b03b3c905a4..33354e1c3bc 100644 --- a/python/cudf/cudf/core/column/datetime.py +++ b/python/cudf/cudf/core/column/datetime.py @@ -27,13 +27,17 @@ is_scalar, is_timedelta64_dtype, ) +from cudf.core._compat import PANDAS_GE_220 from cudf.core.buffer import Buffer, cuda_array_interface_wrapper from cudf.core.column import ColumnBase, as_column, column, string from cudf.core.column.timedelta import _unit_to_nanoseconds_conversion from cudf.utils.dtypes import _get_base_dtype from cudf.utils.utils import _all_bools_with_nulls -_guess_datetime_format = pd.core.tools.datetimes.guess_datetime_format +if PANDAS_GE_220: + _guess_datetime_format = pd.tseries.api.guess_datetime_format +else: + _guess_datetime_format = pd.core.tools.datetimes.guess_datetime_format # nanoseconds per time_unit _dtype_to_format_conversion = { diff --git a/python/cudf/cudf/core/dataframe.py b/python/cudf/cudf/core/dataframe.py index 16eead6ea81..fd4a15a3391 100644 --- a/python/cudf/cudf/core/dataframe.py +++ b/python/cudf/cudf/core/dataframe.py @@ -2313,8 +2313,19 @@ def to_dict( if orient == "series": # Special case needed to avoid converting # cudf.Series objects into pd.Series - into_c = pd.core.common.standardize_mapping(into) - return into_c((k, v) for k, v in self.items()) + if not inspect.isclass(into): + cons = type(into) # type: ignore[assignment] + if isinstance(into, defaultdict): + cons = functools.partial(cons, into.default_factory) + elif issubclass(into, abc.Mapping): + cons = into # type: ignore[assignment] + if issubclass(into, defaultdict): + raise TypeError( + "to_dict() only accepts initialized defaultdicts" + ) + else: + raise TypeError(f"unsupported type: {into}") + return cons(self.items()) # type: ignore[misc] return self.to_pandas().to_dict(orient=orient, into=into) diff --git a/python/cudf/cudf/core/index.py b/python/cudf/cudf/core/index.py index 9f0c66a5c74..277b5d3bb17 100644 --- a/python/cudf/cudf/core/index.py +++ b/python/cudf/cudf/core/index.py @@ -330,17 +330,15 @@ def _data(self): @_cudf_nvtx_annotate def __contains__(self, item): - if not isinstance( + if isinstance(item, bool) or not isinstance( item, tuple(np.sctypes["int"] + np.sctypes["float"] + [int, float]) ): return False try: - item = pd.core.dtypes.common.ensure_python_int(item) - except TypeError: - return False - if not item % 1 == 0: + int_item = int(item) + return int_item == item and int_item in self._range + except (ValueError, OverflowError): return False - return item in range(self._start, self._stop, self._step) @_cudf_nvtx_annotate def copy(self, name=None, deep=False, dtype=None, names=None): diff --git a/python/cudf/cudf/tests/test_rolling.py b/python/cudf/cudf/tests/test_rolling.py index 43fa83e1735..19714b7b9d3 100644 --- a/python/cudf/cudf/tests/test_rolling.py +++ b/python/cudf/cudf/tests/test_rolling.py @@ -8,7 +8,7 @@ import pytest import cudf -from cudf.core._compat import PANDAS_GE_150, PANDAS_LT_140 +from cudf.core._compat import PANDAS_GE_150 from cudf.testing._utils import ( _create_pandas_series_float64_default, assert_eq, @@ -536,10 +536,9 @@ def get_window_bounds( "indexer", [ pd.api.indexers.FixedForwardWindowIndexer(window_size=2), - pd.core.window.expanding.ExpandingIndexer(), - pd.core.window.indexers.FixedWindowIndexer(window_size=3) - if PANDAS_LT_140 - else pd.core.indexers.objects.FixedWindowIndexer(window_size=3), + pd.api.indexers.VariableOffsetWindowIndexer( + index=pd.date_range("2020", periods=5), offset=pd.offsets.BDay(1) + ), ], ) def test_rolling_indexer_support(indexer): From 58387ffee6970115b77793f3706478d41eab328d Mon Sep 17 00:00:00 2001 From: Paul Taylor <178183+trxcllnt@users.noreply.github.com> Date: Mon, 20 Nov 2023 08:51:10 -0800 Subject: [PATCH 29/44] Add cuDF devcontainers (#14015) * add devcontainers * fix tag for CUDA 12.0 * use CUDA 11.8 for now * default to CUDA 12.0 * install cuda-cupti-dev in conda environment * remove MODIFY_PREFIX_PATH so the driver is found * install cuda-nvtx-dev in conda environment * update conda env * add MODIFY_PREFIX_PATH back * temporarily default to my branch with the fix for MODIFY_PREFIX_PATH in conda envs * remove temporary rapids-cmake pin * build all RAPIDS archs to take maximum advantage of sccache * add clangd and nsight vscode customizations * copy in default clangd config * remove options for pip vs. conda unless using the launch script * fix unified mounts * ensure dirs exist before mounting * add compile_commands to .gitignore * allow defining cudf and cudf_kafka include dirs via envvars * add kvikio * use volumes for isolated devcontainer source dirs * update README.md * update to rapidsai/devcontainers 23.10 * update rapids-build-utils version to 23.10 * add .clangd config file * update RAPIDS versions in devcontainer files * ensure the directory for the generated jitify kernels is exists after configuring * add clang and clang-tools 16 * remove isolated and unified devcontainers, make single the default * separate CUDA 11.8 and 12.0 devcontainers * fix version string for requirements.txt * update conda envs * clean up envvars, mounts, and build args, add codespaces post-attach command workaround * consolidate common vscode customizations * enumerate CUDA 11 packages, include up to CUDA 12.2 * include protoc-wheel when generating requirements.txt * default to cuda-python for cu11 * separate devcontainer mounts by CUDA version * add devcontainer build jobs to PR workflow * use pypi.nvidia.com instead of pypi.ngc.nvidia.com * fix venvs mount path * fix lint * ensure rmm-cuXX is included in pip requirements * disable libcudf_kakfa build for now * build dask-cudf * be more explicit in update-versions.sh, make devcontainer build required in pr jobs * revert rename devcontainer job * install librdkafka-dev in pip containers so we can build libcudf_kafka and cudf_kafka * separate cupy, cudf, and cudf_kafka matrices for CUDA 11 and 12 * add fallback include path for RMM * fallback to CUDA_PATH if CUDA_HOME is not set * define envvars in dockerfile * define envvars for cudf_kafka * build verbose * include wheel and setuptools in requirements.txt * switch workflow to branch-23.10 * update clang-tools version to 16.0.6 * fix update-version.sh * Use 24.02 branches. * fix version numbers * Fix dependencies.yaml. * Update .devcontainer/Dockerfile --------- Co-authored-by: Bradley Dice --- .devcontainer/Dockerfile | 35 +++ .devcontainer/README.md | 35 +++ .../cuda11.8-conda/devcontainer.json | 37 ++++ .devcontainer/cuda11.8-pip/devcontainer.json | 36 ++++ .../cuda12.0-conda/devcontainer.json | 37 ++++ .devcontainer/cuda12.0-pip/devcontainer.json | 36 ++++ .github/workflows/pr.yaml | 9 + .gitignore | 4 + .pre-commit-config.yaml | 2 +- ci/release/update-version.sh | 6 + .../all_cuda-118_arch-x86_64.yaml | 2 + .../all_cuda-120_arch-x86_64.yaml | 2 + cpp/.clangd | 65 ++++++ .../Modules/JitifyPreprocessKernels.cmake | 6 +- dependencies.yaml | 203 +++++++++++++----- 15 files changed, 458 insertions(+), 57 deletions(-) create mode 100644 .devcontainer/Dockerfile create mode 100644 .devcontainer/README.md create mode 100644 .devcontainer/cuda11.8-conda/devcontainer.json create mode 100644 .devcontainer/cuda11.8-pip/devcontainer.json create mode 100644 .devcontainer/cuda12.0-conda/devcontainer.json create mode 100644 .devcontainer/cuda12.0-pip/devcontainer.json create mode 100644 cpp/.clangd diff --git a/.devcontainer/Dockerfile b/.devcontainer/Dockerfile new file mode 100644 index 00000000000..c19bb68986f --- /dev/null +++ b/.devcontainer/Dockerfile @@ -0,0 +1,35 @@ +# syntax=docker/dockerfile:1.5 + +ARG BASE +ARG PYTHON_PACKAGE_MANAGER=conda + +FROM ${BASE} as pip-base + +RUN apt update -y \ + && DEBIAN_FRONTEND=noninteractive apt install -y \ + librdkafka-dev \ + && rm -rf /tmp/* /var/tmp/* /var/cache/apt/* /var/lib/apt/lists/*; + +ENV DEFAULT_VIRTUAL_ENV=rapids + +FROM ${BASE} as conda-base + +ENV DEFAULT_CONDA_ENV=rapids + +FROM ${PYTHON_PACKAGE_MANAGER}-base + +ARG CUDA +ENV CUDAARCHS="RAPIDS" +ENV CUDA_VERSION="${CUDA_VERSION:-${CUDA}}" + +ARG PYTHON_PACKAGE_MANAGER +ENV PYTHON_PACKAGE_MANAGER="${PYTHON_PACKAGE_MANAGER}" + +ENV PYTHONSAFEPATH="1" +ENV PYTHONUNBUFFERED="1" +ENV PYTHONDONTWRITEBYTECODE="1" + +ENV SCCACHE_REGION="us-east-2" +ENV SCCACHE_BUCKET="rapids-sccache-devs" +ENV VAULT_HOST="https://vault.ops.k8s.rapids.ai" +ENV HISTFILE="/home/coder/.cache/._bash_history" diff --git a/.devcontainer/README.md b/.devcontainer/README.md new file mode 100644 index 00000000000..91ee7ef85f7 --- /dev/null +++ b/.devcontainer/README.md @@ -0,0 +1,35 @@ +# cuDF Development Containers + +This directory contains [devcontainer configurations](https://containers.dev/implementors/json_reference/) for using VSCode to [develop in a container](https://code.visualstudio.com/docs/devcontainers/containers) via the `Remote Containers` [extension](https://marketplace.visualstudio.com/items?itemName=ms-vscode-remote.remote-containers) or [GitHub Codespaces](https://github.com/codespaces). + +This container is a turnkey development environment for building and testing the cuDF C++ and Python libraries. + +## Table of Contents + +* [Prerequisites](#prerequisites) +* [Host bind mounts](#host-bind-mounts) +* [Launch a Dev Container](#launch-a-dev-container) + +## Prerequisites + +* [VSCode](https://code.visualstudio.com/download) +* [VSCode Remote Containers extension](https://marketplace.visualstudio.com/items?itemName=ms-vscode-remote.remote-containers) + +## Host bind mounts + +By default, the following directories are bind-mounted into the devcontainer: + +* `${repo}:/home/coder/cudf` +* `${repo}/../.aws:/home/coder/.aws` +* `${repo}/../.local:/home/coder/.local` +* `${repo}/../.cache:/home/coder/.cache` +* `${repo}/../.conda:/home/coder/.conda` +* `${repo}/../.config:/home/coder/.config` + +This ensures caches, configurations, dependencies, and your commits are persisted on the host across container runs. + +## Launch a Dev Container + +To launch a devcontainer from VSCode, open the cuDF repo and select the "Reopen in Container" button in the bottom right:
+ +Alternatively, open the VSCode command palette (typically `cmd/ctrl + shift + P`) and run the "Rebuild and Reopen in Container" command. diff --git a/.devcontainer/cuda11.8-conda/devcontainer.json b/.devcontainer/cuda11.8-conda/devcontainer.json new file mode 100644 index 00000000000..6c3322dfc61 --- /dev/null +++ b/.devcontainer/cuda11.8-conda/devcontainer.json @@ -0,0 +1,37 @@ +{ + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "11.8", + "PYTHON_PACKAGE_MANAGER": "conda", + "BASE": "rapidsai/devcontainers:24.02-cpp-llvm16-cuda11.8-mambaforge-ubuntu22.04" + } + }, + "hostRequirements": {"gpu": "optional"}, + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.2": {} + }, + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + "initializeCommand": ["/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}-cuda11.8-envs}"], + "postAttachCommand": ["/bin/bash", "-c", "if [ ${CODESPACES:-false} = 'true' ]; then . devcontainer-utils-post-attach-command; . rapids-post-attach-command; fi"], + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/cudf,type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}-cuda11.8-envs,target=/home/coder/.conda/envs,type=bind,consistency=consistent" + ], + "customizations": { + "vscode": { + "extensions": [ + "ms-python.flake8", + "nvidia.nsight-vscode-edition" + ] + } + } +} diff --git a/.devcontainer/cuda11.8-pip/devcontainer.json b/.devcontainer/cuda11.8-pip/devcontainer.json new file mode 100644 index 00000000000..d2afd3a497e --- /dev/null +++ b/.devcontainer/cuda11.8-pip/devcontainer.json @@ -0,0 +1,36 @@ +{ + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "11.8", + "PYTHON_PACKAGE_MANAGER": "pip", + "BASE": "rapidsai/devcontainers:24.02-cpp-llvm16-cuda11.8-ubuntu22.04" + } + }, + "hostRequirements": {"gpu": "optional"}, + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.2": {} + }, + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + "initializeCommand": ["/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/share/${localWorkspaceFolderBasename}-cuda11.8-venvs}"], + "postAttachCommand": ["/bin/bash", "-c", "if [ ${CODESPACES:-false} = 'true' ]; then . devcontainer-utils-post-attach-command; . rapids-post-attach-command; fi"], + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/cudf,type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.local/share/${localWorkspaceFolderBasename}-cuda11.8-venvs,target=/home/coder/.local/share/venvs,type=bind,consistency=consistent" + ], + "customizations": { + "vscode": { + "extensions": [ + "ms-python.flake8", + "nvidia.nsight-vscode-edition" + ] + } + } +} diff --git a/.devcontainer/cuda12.0-conda/devcontainer.json b/.devcontainer/cuda12.0-conda/devcontainer.json new file mode 100644 index 00000000000..42ed334fe03 --- /dev/null +++ b/.devcontainer/cuda12.0-conda/devcontainer.json @@ -0,0 +1,37 @@ +{ + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "12.0", + "PYTHON_PACKAGE_MANAGER": "conda", + "BASE": "rapidsai/devcontainers:24.02-cpp-mambaforge-ubuntu22.04" + } + }, + "hostRequirements": {"gpu": "optional"}, + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.2": {} + }, + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + "initializeCommand": ["/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}-cuda12.0-envs}"], + "postAttachCommand": ["/bin/bash", "-c", "if [ ${CODESPACES:-false} = 'true' ]; then . devcontainer-utils-post-attach-command; . rapids-post-attach-command; fi"], + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/cudf,type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}-cuda12.0-envs,target=/home/coder/.conda/envs,type=bind,consistency=consistent" + ], + "customizations": { + "vscode": { + "extensions": [ + "ms-python.flake8", + "nvidia.nsight-vscode-edition" + ] + } + } +} diff --git a/.devcontainer/cuda12.0-pip/devcontainer.json b/.devcontainer/cuda12.0-pip/devcontainer.json new file mode 100644 index 00000000000..306a2065ef0 --- /dev/null +++ b/.devcontainer/cuda12.0-pip/devcontainer.json @@ -0,0 +1,36 @@ +{ + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "12.0", + "PYTHON_PACKAGE_MANAGER": "pip", + "BASE": "rapidsai/devcontainers:24.02-cpp-llvm16-cuda12.0-ubuntu22.04" + } + }, + "hostRequirements": {"gpu": "optional"}, + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.2": {} + }, + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + "initializeCommand": ["/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/share/${localWorkspaceFolderBasename}-cuda12.0-venvs}"], + "postAttachCommand": ["/bin/bash", "-c", "if [ ${CODESPACES:-false} = 'true' ]; then . devcontainer-utils-post-attach-command; . rapids-post-attach-command; fi"], + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/cudf,type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.local/share/${localWorkspaceFolderBasename}-cuda12.0-venvs,target=/home/coder/.local/share/venvs,type=bind,consistency=consistent" + ], + "customizations": { + "vscode": { + "extensions": [ + "ms-python.flake8", + "nvidia.nsight-vscode-edition" + ] + } + } +} diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 462fad6e938..5a77c6749fe 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -25,6 +25,7 @@ jobs: - wheel-tests-cudf - wheel-build-dask-cudf - wheel-tests-dask-cudf + - devcontainer - unit-tests-cudf-pandas - pandas-tests #- pandas-tests-diff @@ -130,6 +131,14 @@ jobs: matrix_filter: map(select(.ARCH == "amd64" and .PY_VER == "3.10" and (.CUDA_VER == "11.8.0" or .CUDA_VER == "12.0.1"))) build_type: pull-request script: ci/test_wheel_dask_cudf.sh + devcontainer: + secrets: inherit + uses: rapidsai/shared-action-workflows/.github/workflows/build-in-devcontainer.yaml@branch-24.02 + with: + build_command: | + sccache -z; + build-all -DBUILD_BENCHMARKS=ON -DNVBench_ENABLE_CUPTI=OFF --verbose; + sccache -s; unit-tests-cudf-pandas: needs: wheel-build-cudf secrets: inherit diff --git a/.gitignore b/.gitignore index 4a4a6a98efb..243ba73e5b0 100644 --- a/.gitignore +++ b/.gitignore @@ -173,3 +173,7 @@ jupyter_execute # cibuildwheel /wheelhouse + +# clang tooling +compile_commands.json +.clangd/ diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index a62104722e8..7db8d9ab52f 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -145,7 +145,7 @@ repos: ^CHANGELOG.md$ ) - repo: https://github.com/rapidsai/dependency-file-generator - rev: v1.5.1 + rev: v1.7.1 hooks: - id: rapids-dependency-file-generator args: ["--clean"] diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 16742465c32..f9b1436495a 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -115,3 +115,9 @@ NEXT_FULL_JAVA_TAG="${NEXT_SHORT_TAG}.${PATCH_PEP440}-SNAPSHOT" sed_runner "s|.*-SNAPSHOT|${NEXT_FULL_JAVA_TAG}|g" java/pom.xml sed_runner "s/branch-.*/branch-${NEXT_SHORT_TAG}/g" java/ci/README.md sed_runner "s/cudf-.*-SNAPSHOT/cudf-${NEXT_FULL_JAVA_TAG}/g" java/ci/README.md + +# .devcontainer files +find .devcontainer/ -type f -name devcontainer.json -print0 | while IFS= read -r -d '' filename; do + sed_runner "s@rapidsai/devcontainers:[0-9.]*@rapidsai/devcontainers:${NEXT_SHORT_TAG}@g" "${filename}" + sed_runner "s@rapidsai/devcontainers/features/rapids-build-utils:[0-9.]*@rapidsai/devcontainers/features/rapids-build-utils:${NEXT_SHORT_TAG_PEP440}@" "${filename}" +done diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 1a944fbdb21..0bafdd13af5 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -14,6 +14,8 @@ dependencies: - botocore>=1.24.21 - c-compiler - cachetools +- clang-tools=16.0.6 +- clang==16.0.6 - cmake>=3.26.4 - cramjam - cubinlinker diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index 21837b652f4..27dede9b519 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -14,6 +14,8 @@ dependencies: - botocore>=1.24.21 - c-compiler - cachetools +- clang-tools=16.0.6 +- clang==16.0.6 - cmake>=3.26.4 - cramjam - cuda-cudart-dev diff --git a/cpp/.clangd b/cpp/.clangd new file mode 100644 index 00000000000..7c4fe036ddf --- /dev/null +++ b/cpp/.clangd @@ -0,0 +1,65 @@ +# https://clangd.llvm.org/config + +# Apply a config conditionally to all C files +If: + PathMatch: .*\.(c|h)$ + +--- + +# Apply a config conditionally to all C++ files +If: + PathMatch: .*\.(c|h)pp + +--- + +# Apply a config conditionally to all CUDA files +If: + PathMatch: .*\.cuh? +CompileFlags: + Add: + - "-x" + - "cuda" + # No error on unknown CUDA versions + - "-Wno-unknown-cuda-version" + # Allow variadic CUDA functions + - "-Xclang=-fcuda-allow-variadic-functions" +Diagnostics: + Suppress: + - "variadic_device_fn" + - "attributes_not_allowed" + +--- + +# Tweak the clangd parse settings for all files +CompileFlags: + Add: + # report all errors + - "-ferror-limit=0" + - "-fmacro-backtrace-limit=0" + - "-ftemplate-backtrace-limit=0" + # Skip the CUDA version check + - "--no-cuda-version-check" + Remove: + # remove gcc's -fcoroutines + - -fcoroutines + # remove nvc++ flags unknown to clang + - "-gpu=*" + - "-stdpar*" + # remove nvcc flags unknown to clang + - "-arch*" + - "-gencode*" + - "--generate-code*" + - "-ccbin*" + - "-t=*" + - "--threads*" + - "-Xptxas*" + - "-Xcudafe*" + - "-Xfatbin*" + - "-Xcompiler*" + - "--diag-suppress*" + - "--diag_suppress*" + - "--compiler-options*" + - "--expt-extended-lambda" + - "--expt-relaxed-constexpr" + - "-forward-unknown-to-host-compiler" + - "-Werror=cross-execution-space-call" diff --git a/cpp/cmake/Modules/JitifyPreprocessKernels.cmake b/cpp/cmake/Modules/JitifyPreprocessKernels.cmake index df285bdea55..baabffceeac 100644 --- a/cpp/cmake/Modules/JitifyPreprocessKernels.cmake +++ b/cpp/cmake/Modules/JitifyPreprocessKernels.cmake @@ -1,5 +1,5 @@ # ============================================================================= -# Copyright (c) 2021-2022, NVIDIA CORPORATION. +# Copyright (c) 2021-2023, NVIDIA 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 @@ -54,6 +54,10 @@ function(jit_preprocess_files) ) endfunction() +if(NOT (EXISTS "${CUDF_GENERATED_INCLUDE_DIR}/include")) + make_directory("${CUDF_GENERATED_INCLUDE_DIR}/include") +endif() + jit_preprocess_files( SOURCE_DIRECTORY ${CUDF_SOURCE_DIR}/src FILES binaryop/jit/kernel.cu transform/jit/kernel.cu rolling/jit/kernel.cu diff --git a/dependencies.yaml b/dependencies.yaml index f00273a5db1..a84a8fa2bf2 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -19,12 +19,12 @@ files: - run_common - run_cudf - run_dask_cudf - - run_cudf_kafka - run_custreamz - test_cpp - test_python_common - test_python_cudf - test_python_dask_cudf + - depends_on_cupy test_cpp: output: none includes: @@ -83,6 +83,7 @@ files: - run_common - run_cudf - pyarrow_run + - depends_on_cupy py_test_cudf: output: pyproject pyproject_dir: python/cudf @@ -123,6 +124,8 @@ files: includes: - run_common - run_dask_cudf + - depends_on_cudf + - depends_on_cupy py_test_dask_cudf: output: pyproject pyproject_dir: python/dask_cudf @@ -146,7 +149,7 @@ files: extras: table: project includes: - - run_cudf_kafka + - depends_on_cudf py_test_cudf_kafka: output: pyproject pyproject_dir: python/cudf_kafka @@ -169,6 +172,8 @@ files: table: project includes: - run_custreamz + - depends_on_cudf + - depends_on_cudf_kafka py_test_custreamz: output: pyproject pyproject_dir: python/custreamz @@ -229,16 +234,14 @@ dependencies: - nvcc_linux-aarch64=11.8 build_cpp: common: - - output_types: [conda, requirements] - packages: - - librmm==24.2.* - - libkvikio==24.2.* - output_types: conda packages: - fmt>=9.1.0,<10 - &gbench benchmark==1.8.0 - >est gtest>=1.13.0 - &gmock gmock>=1.13.0 + - librmm==24.2.* + - libkvikio==24.2.* # Hard pin the patch version used during the build. This must be kept # in sync with the version pinned in get_arrow.cmake. - libarrow-all==14.0.1.* @@ -248,7 +251,7 @@ dependencies: - spdlog>=1.11.0,<1.12 build_wheels: common: - - output_types: pyproject + - output_types: [requirements, pyproject] packages: - wheel - setuptools @@ -267,15 +270,40 @@ dependencies: - pyarrow==14.0.1.* build_python_cudf: common: - - output_types: [conda, requirements, pyproject] - packages: - - rmm==24.2.* - output_types: conda packages: + - &rmm_conda rmm==24.2.* - &protobuf protobuf>=4.21,<5 - - output_types: pyproject + - pip + - pip: + - git+https://github.com/python-streamz/streamz.git@master + - output_types: [requirements, pyproject] packages: - protoc-wheel + - output_types: requirements + packages: + # pip recognizes the index as a global option for the requirements.txt file + # This index is needed for rmm-cu{11,12}. + - --extra-index-url=https://pypi.nvidia.com + - git+https://github.com/python-streamz/streamz.git@master + specific: + - output_types: [requirements, pyproject] + matrices: + - matrix: {cuda: "12.2"} + packages: &build_python_packages_cu12 + - &rmm_cu12 rmm-cu12==24.2.* + - {matrix: {cuda: "12.1"}, packages: *build_python_packages_cu12} + - {matrix: {cuda: "12.0"}, packages: *build_python_packages_cu12} + - matrix: {cuda: "11.8"} + packages: &build_python_packages_cu11 + - &rmm_cu11 rmm-cu11==24.2.* + - {matrix: {cuda: "11.5"}, packages: *build_python_packages_cu11} + - {matrix: {cuda: "11.4"}, packages: *build_python_packages_cu11} + - {matrix: {cuda: "11.2"}, packages: *build_python_packages_cu11} + - {matrix: null, packages: null } + - output_types: pyproject + matrices: + - {matrix: null, packages: [*rmm_conda] } libarrow_run: common: - output_types: conda @@ -385,6 +413,8 @@ dependencies: - identify>=2.5.20 - output_types: conda packages: + - clang==16.0.6 + - clang-tools=16.0.6 - &doxygen doxygen=1.9.1 # pre-commit hook needs a specific version. docs: common: @@ -444,72 +474,65 @@ dependencies: - nvtx>=0.2.1 - packaging - rich - - rmm==24.2.* - typing_extensions>=4.0.0 - *protobuf - output_types: conda packages: - - cupy>=12.0.0 - - pip - - pip: - - git+https://github.com/python-streamz/streamz.git@master + - *rmm_conda - output_types: requirements packages: # pip recognizes the index as a global option for the requirements.txt file # This index is needed for rmm, cubinlinker, ptxcompiler. - - --extra-index-url=https://pypi.ngc.nvidia.com - - git+https://github.com/python-streamz/streamz.git@master - - &cupy_pip cupy-cuda11x>=12.0.0 - - output_types: pyproject - packages: - - cubinlinker - - *cupy_pip - - ptxcompiler + - --extra-index-url=https://pypi.nvidia.com specific: - output_types: [conda, requirements, pyproject] matrices: - - matrix: - cuda: "12.0" - packages: + - matrix: {cuda: "12.2"} + packages: &run_cudf_packages_all_cu12 - cuda-python>=12.0,<13.0a0 - - matrix: # All CUDA 11 versions - packages: + - {matrix: {cuda: "12.1"}, packages: *run_cudf_packages_all_cu12} + - {matrix: {cuda: "12.0"}, packages: *run_cudf_packages_all_cu12} + - matrix: {cuda: "11.8"} + packages: &run_cudf_packages_all_cu11 - cuda-python>=11.7.1,<12.0a0 - - output_types: [conda, pyproject] + - {matrix: {cuda: "11.5"}, packages: *run_cudf_packages_all_cu11} + - {matrix: {cuda: "11.4"}, packages: *run_cudf_packages_all_cu11} + - {matrix: {cuda: "11.2"}, packages: *run_cudf_packages_all_cu11} + - {matrix: null, packages: *run_cudf_packages_all_cu11} + - output_types: conda matrices: - - matrix: - cuda: "12.0" - packages: - - matrix: # All CUDA 11 versions - packages: + - matrix: {cuda: "11.8"} + packages: &run_cudf_packages_conda_cu11 - cubinlinker - ptxcompiler - - output_types: requirements + - {matrix: {cuda: "11.5"}, packages: *run_cudf_packages_conda_cu11} + - {matrix: {cuda: "11.4"}, packages: *run_cudf_packages_conda_cu11} + - {matrix: {cuda: "11.2"}, packages: *run_cudf_packages_conda_cu11} + - {matrix: null, packages: null} + - output_types: [requirements, pyproject] matrices: - - matrix: - cuda: "12.0" - packages: - - matrix: # All CUDA 11 versions - packages: + - matrix: {cuda: "12.2"} + packages: &run_cudf_packages_pip_cu12 + - rmm-cu12==24.2.* + - {matrix: {cuda: "12.1"}, packages: *run_cudf_packages_pip_cu12} + - {matrix: {cuda: "12.0"}, packages: *run_cudf_packages_pip_cu12} + - matrix: {cuda: "11.8"} + packages: &run_cudf_packages_pip_cu11 + - rmm-cu11==24.2.* - cubinlinker-cu11 - ptxcompiler-cu11 + - {matrix: {cuda: "11.5"}, packages: *run_cudf_packages_pip_cu11} + - {matrix: {cuda: "11.4"}, packages: *run_cudf_packages_pip_cu11} + - {matrix: {cuda: "11.2"}, packages: *run_cudf_packages_pip_cu11} + - {matrix: null, packages: null} + - output_types: pyproject + matrices: + - {matrix: null, packages: [cubinlinker, ptxcompiler, *rmm_conda] } run_dask_cudf: common: - output_types: [conda, requirements, pyproject] packages: - rapids-dask-dependency==24.2.* - - output_types: conda - packages: - - cupy>=12.0.0 - - output_types: pyproject - packages: - - &cudf cudf==24.2.* - - *cupy_pip - run_cudf_kafka: - common: - - output_types: [requirements, pyproject] - packages: - - *cudf run_custreamz: common: - output_types: conda @@ -521,8 +544,6 @@ dependencies: - output_types: [requirements, pyproject] packages: - confluent-kafka>=1.9.0,<1.10.0a0 - - *cudf - - cudf_kafka==24.2.* test_cpp: common: - output_types: conda @@ -606,6 +627,78 @@ dependencies: packages: - dask-cuda==24.2.* - *numba + depends_on_cudf: + common: + - output_types: conda + packages: + - &cudf_conda cudf==24.2.* + - output_types: requirements + packages: + # pip recognizes the index as a global option for the requirements.txt file + # This index is needed for rmm, cubinlinker, ptxcompiler. + - --extra-index-url=https://pypi.nvidia.com + specific: + - output_types: [requirements, pyproject] + matrices: + - matrix: {cuda: "12.2"} + packages: &cudf_packages_pip_cu12 + - cudf-cu12==24.2.* + - {matrix: {cuda: "12.1"}, packages: *cudf_packages_pip_cu12} + - {matrix: {cuda: "12.0"}, packages: *cudf_packages_pip_cu12} + - matrix: {cuda: "11.8"} + packages: &cudf_packages_pip_cu11 + - cudf-cu11==24.2.* + - {matrix: {cuda: "11.5"}, packages: *cudf_packages_pip_cu11} + - {matrix: {cuda: "11.4"}, packages: *cudf_packages_pip_cu11} + - {matrix: {cuda: "11.2"}, packages: *cudf_packages_pip_cu11} + - {matrix: null, packages: [*cudf_conda]} + depends_on_cudf_kafka: + common: + - output_types: conda + packages: + - &cudf_kafka_conda cudf_kafka==24.2.* + - output_types: requirements + packages: + # pip recognizes the index as a global option for the requirements.txt file + # This index is needed for rmm, cubinlinker, ptxcompiler. + - --extra-index-url=https://pypi.nvidia.com + specific: + - output_types: [requirements, pyproject] + matrices: + - matrix: {cuda: "12.2"} + packages: &cudf_kafka_packages_pip_cu12 + - cudf_kafka-cu12==24.2.* + - {matrix: {cuda: "12.1"}, packages: *cudf_kafka_packages_pip_cu12} + - {matrix: {cuda: "12.0"}, packages: *cudf_kafka_packages_pip_cu12} + - matrix: {cuda: "11.8"} + packages: &cudf_kafka_packages_pip_cu11 + - cudf_kafka-cu11==24.2.* + - {matrix: {cuda: "11.5"}, packages: *cudf_kafka_packages_pip_cu11} + - {matrix: {cuda: "11.4"}, packages: *cudf_kafka_packages_pip_cu11} + - {matrix: {cuda: "11.2"}, packages: *cudf_kafka_packages_pip_cu11} + - {matrix: null, packages: [*cudf_kafka_conda]} + depends_on_cupy: + common: + - output_types: conda + packages: + - cupy>=12.0.0 + specific: + - output_types: [requirements, pyproject] + matrices: + # All CUDA 12 versions + - matrix: {cuda: "12.2"} + packages: &cupy_packages_cu12 + - cupy-cuda12x>=12.0.0 + - {matrix: {cuda: "12.1"}, packages: *cupy_packages_cu12} + - {matrix: {cuda: "12.0"}, packages: *cupy_packages_cu12} + # All CUDA 11 versions + - matrix: {cuda: "11.8"} + packages: &cupy_packages_cu11 + - cupy-cuda11x>=12.0.0 + - {matrix: {cuda: "11.5"}, packages: *cupy_packages_cu11} + - {matrix: {cuda: "11.4"}, packages: *cupy_packages_cu11} + - {matrix: {cuda: "11.2"}, packages: *cupy_packages_cu11} + - {matrix: null, packages: *cupy_packages_cu11} test_python_pandas_cudf: common: - output_types: pyproject From 3ef13d07057e87cff1cad4e0aa9460b3b5c45459 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 20 Nov 2023 13:02:29 -0600 Subject: [PATCH 30/44] Fix io reference in docs. (#14452) cuDF CI is failing to build docs due to an ambiguous reference `io`. This PR makes that reference unambiguous. ``` /__w/cudf/cudf/docs/cudf/source/user_guide/data-types.md:139: WARNING: Multiple matches found for 'io': pandas:std:label:io, pyarrow:std:label:io, python:py:module:io [myst.iref_ambiguous] ``` I used this output to help me find the object inventory that lists this. ```bash python -m sphinx.ext.intersphinx https://docs.rapids.ai/api/cudf/stable/objects.inv ``` I also looked at the MyST docs on external references. https://mystmd.org/guide/external-references Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14452 --- docs/cudf/source/user_guide/data-types.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/cudf/source/user_guide/data-types.md b/docs/cudf/source/user_guide/data-types.md index 1f4cfbc7366..e6fe3109c57 100644 --- a/docs/cudf/source/user_guide/data-types.md +++ b/docs/cudf/source/user_guide/data-types.md @@ -136,7 +136,7 @@ dtype: struct StructDtype({'a': dtype('int64'), 'b': dtype('int64')}) ``` -Or by reading them from disk, using a [file format that supports nested data](io). +Or by reading them from disk, using a [file format that supports nested data](/user_guide/io/index.md). ```python >>> pdf = pd.DataFrame({"a": [[1, 2], [3, 4, 5], [6, 7, 8]]}) From 823d3214a9489e3c496aa31041b5d29f650e94b3 Mon Sep 17 00:00:00 2001 From: brandon-b-miller <53796099+brandon-b-miller@users.noreply.github.com> Date: Mon, 20 Nov 2023 16:33:28 -0600 Subject: [PATCH 31/44] Use `pynvjitlink` for CUDA 12+ MVC (#13650) Fixes https://github.com/rapidsai/cudf/issues/12822 This PR provides minor version compatibility in the CUDA 12.x range through `nvjitlink` via the preliminary [nvjiitlink python binding](https://github.com/gmarkall/nvjitlink). Thus far this PR merely leverages a local installation of the library and should not be merged until `nvjitlink` is hosted on `conda-forge` and cuDF's dependencies are adjusted accordingly, likely as part of this PR. Authors: - https://github.com/brandon-b-miller - Ashwin Srinath (https://github.com/shwina) Approvers: - Bradley Dice (https://github.com/bdice) - Ashwin Srinath (https://github.com/shwina) URL: https://github.com/rapidsai/cudf/pull/13650 --- python/cudf/cudf/tests/test_mvc.py | 99 +++++++++++++++++++++ python/cudf/cudf/tests/test_numba_import.py | 48 ---------- python/cudf/cudf/utils/_numba.py | 53 ++++++----- 3 files changed, 128 insertions(+), 72 deletions(-) create mode 100644 python/cudf/cudf/tests/test_mvc.py delete mode 100644 python/cudf/cudf/tests/test_numba_import.py diff --git a/python/cudf/cudf/tests/test_mvc.py b/python/cudf/cudf/tests/test_mvc.py new file mode 100644 index 00000000000..7dd25ebc500 --- /dev/null +++ b/python/cudf/cudf/tests/test_mvc.py @@ -0,0 +1,99 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. +import subprocess +import sys + +import pytest + +IS_CUDA_11 = False +IS_CUDA_12 = False +try: + from ptxcompiler.patch import safe_get_versions +except ModuleNotFoundError: + from cudf.utils._ptxcompiler import safe_get_versions + +# do not test cuda 12 if pynvjitlink isn't present +HAVE_PYNVJITLINK = False +try: + import pynvjitlink # noqa: F401 + + HAVE_PYNVJITLINK = True +except ModuleNotFoundError: + pass + + +versions = safe_get_versions() +driver_version, runtime_version = versions + +if (11, 0) <= driver_version < (12, 0): + IS_CUDA_11 = True +if (12, 0) <= driver_version < (13, 0): + IS_CUDA_12 = True + + +TEST_BODY = """ +@numba.cuda.jit +def test_kernel(x): + id = numba.cuda.grid(1) + if id < len(x): + x[id] += 1 + +s = cudf.Series([1, 2, 3]) +with _CUDFNumbaConfig(): + test_kernel.forall(len(s))(s) +""" + +CUDA_11_TEST = ( + """ +import numba.cuda +import cudf +from cudf.utils._numba import _CUDFNumbaConfig, patch_numba_linker_cuda_11 + + +patch_numba_linker_cuda_11() +""" + + TEST_BODY +) + + +CUDA_12_TEST = ( + """ +import numba.cuda +import cudf +from cudf.utils._numba import _CUDFNumbaConfig +from pynvjitlink.patch import ( + patch_numba_linker as patch_numba_linker_pynvjitlink, +) + +patch_numba_linker_pynvjitlink() +""" + + TEST_BODY +) + + +@pytest.mark.parametrize( + "test", + [ + pytest.param( + CUDA_11_TEST, + marks=pytest.mark.skipif( + not IS_CUDA_11, + reason="Minor Version Compatibility test for CUDA 11", + ), + ), + pytest.param( + CUDA_12_TEST, + marks=pytest.mark.skipif( + not IS_CUDA_12 or not HAVE_PYNVJITLINK, + reason="Minor Version Compatibility test for CUDA 12", + ), + ), + ], +) +def test_numba_mvc(test): + cp = subprocess.run( + [sys.executable, "-c", test], + capture_output=True, + cwd="/", + ) + + assert cp.returncode == 0 diff --git a/python/cudf/cudf/tests/test_numba_import.py b/python/cudf/cudf/tests/test_numba_import.py deleted file mode 100644 index 238a32a94fa..00000000000 --- a/python/cudf/cudf/tests/test_numba_import.py +++ /dev/null @@ -1,48 +0,0 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. -import subprocess -import sys - -import pytest - -IS_CUDA_11 = False -try: - from ptxcompiler.patch import NO_DRIVER, safe_get_versions - - versions = safe_get_versions() - if versions != NO_DRIVER: - driver_version, runtime_version = versions - if driver_version < (12, 0): - IS_CUDA_11 = True -except ModuleNotFoundError: - pass - -TEST_NUMBA_MVC_ENABLED = """ -import numba.cuda -import cudf -from cudf.utils._numba import _CUDFNumbaConfig, _patch_numba_mvc - - -_patch_numba_mvc() - -@numba.cuda.jit -def test_kernel(x): - id = numba.cuda.grid(1) - if id < len(x): - x[id] += 1 - -s = cudf.Series([1, 2, 3]) -with _CUDFNumbaConfig(): - test_kernel.forall(len(s))(s) -""" - - -@pytest.mark.skipif( - not IS_CUDA_11, reason="Minor Version Compatibility test for CUDA 11" -) -def test_numba_mvc_enabled_cuda_11(): - cp = subprocess.run( - [sys.executable, "-c", TEST_NUMBA_MVC_ENABLED], - capture_output=True, - cwd="/", - ) - assert cp.returncode == 0 diff --git a/python/cudf/cudf/utils/_numba.py b/python/cudf/cudf/utils/_numba.py index 09afb5680bd..bc0d6f37d89 100644 --- a/python/cudf/cudf/utils/_numba.py +++ b/python/cudf/cudf/utils/_numba.py @@ -7,6 +7,19 @@ from numba import config as numba_config +try: + from pynvjitlink.patch import ( + patch_numba_linker as patch_numba_linker_pynvjitlink, + ) +except ImportError: + + def patch_numba_linker_pynvjitlink(): + warnings.warn( + "CUDA Toolkit is newer than CUDA driver. " + "Numba features will not work in this configuration. " + ) + + CC_60_PTX_FILE = os.path.join( os.path.dirname(__file__), "../core/udf/shim_60.ptx" ) @@ -65,7 +78,7 @@ def _get_ptx_file(path, prefix): return regular_result[1] -def _patch_numba_mvc(): +def patch_numba_linker_cuda_11(): # Enable the config option for minor version compatibility numba_config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY = 1 @@ -106,29 +119,19 @@ def _setup_numba(): versions = safe_get_versions() if versions != NO_DRIVER: driver_version, runtime_version = versions - if driver_version >= (12, 0) and runtime_version > driver_version: - warnings.warn( - f"Using CUDA toolkit version {runtime_version} with CUDA " - f"driver version {driver_version} requires minor version " - "compatibility, which is not yet supported for CUDA " - "driver versions 12.0 and above. It is likely that many " - "cuDF operations will not work in this state. Please " - f"install CUDA toolkit version {driver_version} to " - "continue using cuDF." - ) - else: - # Support MVC for all CUDA versions in the 11.x range - ptx_toolkit_version = _get_cuda_version_from_ptx_file( - CC_60_PTX_FILE - ) - # Numba thinks cubinlinker is only needed if the driver is older - # than the CUDA runtime, but when PTX files are present, it might - # also need to patch because those PTX files may be compiled by - # a CUDA version that is newer than the driver as well - if (driver_version < ptx_toolkit_version) or ( - driver_version < runtime_version - ): - _patch_numba_mvc() + ptx_toolkit_version = _get_cuda_version_from_ptx_file(CC_60_PTX_FILE) + + # MVC is required whenever any PTX is newer than the driver + # This could be the shipped PTX file or the PTX emitted by + # the version of NVVM on the user system, the latter aligning + # with the runtime version + if (driver_version < ptx_toolkit_version) or ( + driver_version < runtime_version + ): + if driver_version < (12, 0): + patch_numba_linker_cuda_11() + else: + patch_numba_linker_pynvjitlink() def _get_cuda_version_from_ptx_file(path): @@ -171,6 +174,8 @@ def _get_cuda_version_from_ptx_file(path): "7.8": (11, 8), "8.0": (12, 0), "8.1": (12, 1), + "8.2": (12, 2), + "8.3": (12, 3), } cuda_ver = ver_map.get(version) From 5831beb80dab9cc23668b5a701d9a92a4797fe70 Mon Sep 17 00:00:00 2001 From: Vukasin Milovanovic Date: Mon, 20 Nov 2023 15:39:35 -0800 Subject: [PATCH 32/44] Remove the use of `volatile` in Parquet (#14448) `volatile` should no be required in our code, unless there are compiler or synchronization issues. This PR removes the use in Parquet reader and writer. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - David Wendt (https://github.com/davidwendt) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/14448 --- cpp/src/io/parquet/decode_preprocess.cu | 2 +- cpp/src/io/parquet/page_data.cu | 31 +++++++++--------------- cpp/src/io/parquet/page_decode.cuh | 29 ++++++++++------------ cpp/src/io/parquet/page_enc.cu | 2 +- cpp/src/io/parquet/page_string_decode.cu | 4 +-- 5 files changed, 28 insertions(+), 40 deletions(-) diff --git a/cpp/src/io/parquet/decode_preprocess.cu b/cpp/src/io/parquet/decode_preprocess.cu index 544c93ee616..d9f91ed564c 100644 --- a/cpp/src/io/parquet/decode_preprocess.cu +++ b/cpp/src/io/parquet/decode_preprocess.cu @@ -61,7 +61,7 @@ __device__ size_type gpuDecodeTotalPageStringSize(page_state_s* s, int t) } else if ((s->col.data_type & 7) == BYTE_ARRAY) { str_len = gpuInitStringDescriptors(s, nullptr, target_pos, t); } - if (!t) { *(int32_t volatile*)&s->dict_pos = target_pos; } + if (!t) { s->dict_pos = target_pos; } return str_len; } diff --git a/cpp/src/io/parquet/page_data.cu b/cpp/src/io/parquet/page_data.cu index 0c53877f7c7..1a94f05498e 100644 --- a/cpp/src/io/parquet/page_data.cu +++ b/cpp/src/io/parquet/page_data.cu @@ -39,10 +39,7 @@ constexpr int rolling_buf_size = decode_block_size * 2; * @param[in] dstv Pointer to row output data (string descriptor or 32-bit hash) */ template -inline __device__ void gpuOutputString(volatile page_state_s* s, - volatile state_buf* sb, - int src_pos, - void* dstv) +inline __device__ void gpuOutputString(page_state_s* s, state_buf* sb, int src_pos, void* dstv) { auto [ptr, len] = gpuGetStringData(s, sb, src_pos); // make sure to only hash `BYTE_ARRAY` when specified with the output type size @@ -69,7 +66,7 @@ inline __device__ void gpuOutputString(volatile page_state_s* s, * @param[in] dst Pointer to row output data */ template -inline __device__ void gpuOutputBoolean(volatile state_buf* sb, int src_pos, uint8_t* dst) +inline __device__ void gpuOutputBoolean(state_buf* sb, int src_pos, uint8_t* dst) { *dst = sb->dict_idx[rolling_index(src_pos)]; } @@ -143,8 +140,8 @@ inline __device__ void gpuStoreOutput(uint2* dst, * @param[out] dst Pointer to row output data */ template -inline __device__ void gpuOutputInt96Timestamp(volatile page_state_s* s, - volatile state_buf* sb, +inline __device__ void gpuOutputInt96Timestamp(page_state_s* s, + state_buf* sb, int src_pos, int64_t* dst) { @@ -218,8 +215,8 @@ inline __device__ void gpuOutputInt96Timestamp(volatile page_state_s* s, * @param[in] dst Pointer to row output data */ template -inline __device__ void gpuOutputInt64Timestamp(volatile page_state_s* s, - volatile state_buf* sb, +inline __device__ void gpuOutputInt64Timestamp(page_state_s* s, + state_buf* sb, int src_pos, int64_t* dst) { @@ -301,10 +298,7 @@ __device__ void gpuOutputByteArrayAsInt(char const* ptr, int32_t len, T* dst) * @param[in] dst Pointer to row output data */ template -__device__ void gpuOutputFixedLenByteArrayAsInt(volatile page_state_s* s, - volatile state_buf* sb, - int src_pos, - T* dst) +__device__ void gpuOutputFixedLenByteArrayAsInt(page_state_s* s, state_buf* sb, int src_pos, T* dst) { uint32_t const dtype_len_in = s->dtype_len_in; uint8_t const* data = s->dict_base ? s->dict_base : s->data_start; @@ -338,10 +332,7 @@ __device__ void gpuOutputFixedLenByteArrayAsInt(volatile page_state_s* s, * @param[in] dst Pointer to row output data */ template -inline __device__ void gpuOutputFast(volatile page_state_s* s, - volatile state_buf* sb, - int src_pos, - T* dst) +inline __device__ void gpuOutputFast(page_state_s* s, state_buf* sb, int src_pos, T* dst) { uint8_t const* dict; uint32_t dict_pos, dict_size = s->dict_size; @@ -371,7 +362,7 @@ inline __device__ void gpuOutputFast(volatile page_state_s* s, */ template static __device__ void gpuOutputGeneric( - volatile page_state_s* s, volatile state_buf* sb, int src_pos, uint8_t* dst8, int len) + page_state_s* s, state_buf* sb, int src_pos, uint8_t* dst8, int len) { uint8_t const* dict; uint32_t dict_pos, dict_size = s->dict_size; @@ -512,7 +503,7 @@ __global__ void __launch_bounds__(decode_block_size) (s->col.data_type & 7) == FIXED_LEN_BYTE_ARRAY) { gpuInitStringDescriptors(s, sb, src_target_pos, t & 0x1f); } - if (t == 32) { *(volatile int32_t*)&s->dict_pos = src_target_pos; } + if (t == 32) { s->dict_pos = src_target_pos; } } else { // WARP1..WARP3: Decode values int const dtype = s->col.data_type & 7; @@ -601,7 +592,7 @@ __global__ void __launch_bounds__(decode_block_size) } } - if (t == out_thread0) { *(volatile int32_t*)&s->src_pos = target_pos; } + if (t == out_thread0) { s->src_pos = target_pos; } } __syncthreads(); } diff --git a/cpp/src/io/parquet/page_decode.cuh b/cpp/src/io/parquet/page_decode.cuh index 4db9bd3904b..a521f4af039 100644 --- a/cpp/src/io/parquet/page_decode.cuh +++ b/cpp/src/io/parquet/page_decode.cuh @@ -71,15 +71,15 @@ struct page_state_s { // points to either nesting_decode_cache above when possible, or to the global source otherwise PageNestingDecodeInfo* nesting_info{}; - inline __device__ void set_error_code(decode_error err) volatile + inline __device__ void set_error_code(decode_error err) { - cuda::atomic_ref ref{const_cast(error)}; + cuda::atomic_ref ref{error}; ref.fetch_or(static_cast(err), cuda::std::memory_order_relaxed); } - inline __device__ void reset_error_code() volatile + inline __device__ void reset_error_code() { - cuda::atomic_ref ref{const_cast(error)}; + cuda::atomic_ref ref{error}; ref.store(0, cuda::std::memory_order_release); } }; @@ -185,8 +185,8 @@ inline __device__ bool is_page_contained(page_state_s* const s, size_t start_row * @return A pair containing a pointer to the string and its length */ template -inline __device__ cuda::std::pair gpuGetStringData(page_state_s volatile* s, - state_buf volatile* sb, +inline __device__ cuda::std::pair gpuGetStringData(page_state_s* s, + state_buf* sb, int src_pos) { char const* ptr = nullptr; @@ -232,8 +232,10 @@ inline __device__ cuda::std::pair gpuGetStringData(page_sta * additional values. */ template -__device__ cuda::std::pair gpuDecodeDictionaryIndices( - page_state_s volatile* s, [[maybe_unused]] state_buf volatile* sb, int target_pos, int t) +__device__ cuda::std::pair gpuDecodeDictionaryIndices(page_state_s* s, + [[maybe_unused]] state_buf* sb, + int target_pos, + int t) { uint8_t const* end = s->data_end; int dict_bits = s->dict_bits; @@ -349,10 +351,7 @@ __device__ cuda::std::pair gpuDecodeDictionaryIndices( * @return The new output position */ template -inline __device__ int gpuDecodeRleBooleans(page_state_s volatile* s, - state_buf volatile* sb, - int target_pos, - int t) +inline __device__ int gpuDecodeRleBooleans(page_state_s* s, state_buf* sb, int target_pos, int t) { uint8_t const* end = s->data_end; int64_t pos = s->dict_pos; @@ -420,10 +419,8 @@ inline __device__ int gpuDecodeRleBooleans(page_state_s volatile* s, * @return Total length of strings processed */ template -__device__ size_type gpuInitStringDescriptors(page_state_s volatile* s, - [[maybe_unused]] state_buf volatile* sb, - int target_pos, - int t) +__device__ size_type +gpuInitStringDescriptors(page_state_s* s, [[maybe_unused]] state_buf* sb, int target_pos, int t) { int pos = s->dict_pos; int total_len = 0; diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu index 2b7980c93e9..d75608930d5 100644 --- a/cpp/src/io/parquet/page_enc.cu +++ b/cpp/src/io/parquet/page_enc.cu @@ -90,7 +90,7 @@ struct page_enc_state_s { uint32_t rle_rpt_count; uint32_t page_start_val; uint32_t chunk_start_val; - volatile uint32_t rpt_map[num_encode_warps]; + uint32_t rpt_map[num_encode_warps]; EncPage page; EncColumnChunk ck; parquet_column_device_view col; diff --git a/cpp/src/io/parquet/page_string_decode.cu b/cpp/src/io/parquet/page_string_decode.cu index e29db042401..916eaa3d681 100644 --- a/cpp/src/io/parquet/page_string_decode.cu +++ b/cpp/src/io/parquet/page_string_decode.cu @@ -851,7 +851,7 @@ __global__ void __launch_bounds__(decode_block_size) } else { gpuInitStringDescriptors(s, sb, src_target_pos, lane_id); } - if (t == 32) { *(volatile int32_t*)&s->dict_pos = src_target_pos; } + if (t == 32) { s->dict_pos = src_target_pos; } } else { int const me = t - out_thread0; @@ -934,7 +934,7 @@ __global__ void __launch_bounds__(decode_block_size) } } - if (t == out_thread0) { *(volatile int32_t*)&s->src_pos = target_pos; } + if (t == out_thread0) { s->src_pos = target_pos; } } __syncthreads(); } From 947081f5b10ca972826942b84c5c2530050325d8 Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Tue, 21 Nov 2023 00:45:37 +0000 Subject: [PATCH 33/44] Remove warning in dask-cudf docs (#14454) Move `from_delayed` and `concat` to appropriate subsections. - Closes #14299 Authors: - Lawrence Mitchell (https://github.com/wence-) - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Bradley Dice (https://github.com/bdice) - Vyas Ramasubramani (https://github.com/vyasr) URL: https://github.com/rapidsai/cudf/pull/14454 --- docs/dask_cudf/source/api.rst | 13 +++++-------- 1 file changed, 5 insertions(+), 8 deletions(-) diff --git a/docs/dask_cudf/source/api.rst b/docs/dask_cudf/source/api.rst index 893f5dd7434..db32f4bbcb3 100644 --- a/docs/dask_cudf/source/api.rst +++ b/docs/dask_cudf/source/api.rst @@ -19,6 +19,7 @@ data reading facilities, followed by calling :members: from_cudf, from_dask_dataframe, + from_delayed, read_csv, read_json, read_orc, @@ -26,14 +27,6 @@ data reading facilities, followed by calling read_text, read_parquet -.. warning:: - - FIXME: where should the following live? - - .. autofunction:: dask_cudf.concat - - .. autofunction:: dask_cudf.from_delayed - Grouping ======== @@ -77,3 +70,7 @@ identical. The full API is provided below. :members: :inherited-members: :show-inheritance: + +.. automodule:: dask_cudf + :members: + concat From fcc89503c1f1e15ec287519959013adcf2bf8a52 Mon Sep 17 00:00:00 2001 From: Matthew Roeschke <10647082+mroeschke@users.noreply.github.com> Date: Tue, 21 Nov 2023 05:19:11 -1000 Subject: [PATCH 34/44] Preserve DataFrame(columns=).columns dtype during empty-like construction (#14381) `.column` used to always return `pd.Index([], dtype=object)` even if an empty-dtyped columns was passed into the DataFrame constructor e.g. `DatetimeIndex([])`. Needed to preserved some information about what column dtype was passed in so we can return a correctly type Index Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14381 --- python/cudf/cudf/core/column_accessor.py | 14 +++++++++++++- python/cudf/cudf/core/dataframe.py | 6 ++++++ python/cudf/cudf/tests/test_dataframe.py | 11 +++++++++++ 3 files changed, 30 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/core/column_accessor.py b/python/cudf/cudf/core/column_accessor.py index 93105b4a252..b106b8bbb02 100644 --- a/python/cudf/cudf/core/column_accessor.py +++ b/python/cudf/cudf/core/column_accessor.py @@ -27,6 +27,7 @@ from cudf.core import column if TYPE_CHECKING: + from cudf._typing import Dtype from cudf.core.column import ColumnBase @@ -99,6 +100,9 @@ class ColumnAccessor(abc.MutableMapping): rangeindex : bool, optional Whether the keys should be returned as a RangeIndex in `to_pandas_index` (default=False). + label_dtype : Dtype, optional + What dtype should be returned in `to_pandas_index` + (default=None). """ _data: "Dict[Any, ColumnBase]" @@ -111,8 +115,10 @@ def __init__( multiindex: bool = False, level_names=None, rangeindex: bool = False, + label_dtype: Dtype | None = None, ): self.rangeindex = rangeindex + self.label_dtype = label_dtype if data is None: data = {} # TODO: we should validate the keys of `data` @@ -123,6 +129,7 @@ def __init__( self.multiindex = multiindex self._level_names = level_names self.rangeindex = data.rangeindex + self.label_dtype = data.label_dtype else: # This code path is performance-critical for copies and should be # modified with care. @@ -292,7 +299,12 @@ def to_pandas_index(self) -> pd.Index: self.names[0], self.names[-1] + diff, diff ) return pd.RangeIndex(new_range, name=self.name) - result = pd.Index(self.names, name=self.name, tupleize_cols=False) + result = pd.Index( + self.names, + name=self.name, + tupleize_cols=False, + dtype=self.label_dtype, + ) return result def insert( diff --git a/python/cudf/cudf/core/dataframe.py b/python/cudf/cudf/core/dataframe.py index fd4a15a3391..43ae9b9e81e 100644 --- a/python/cudf/cudf/core/dataframe.py +++ b/python/cudf/cudf/core/dataframe.py @@ -734,6 +734,7 @@ def __init__( rangeindex = isinstance( columns, (range, pd.RangeIndex, cudf.RangeIndex) ) + label_dtype = getattr(columns, "dtype", None) self._data = ColumnAccessor( { k: column.column_empty( @@ -745,6 +746,7 @@ def __init__( if isinstance(columns, pd.Index) else None, rangeindex=rangeindex, + label_dtype=label_dtype, ) elif isinstance(data, ColumnAccessor): raise TypeError( @@ -995,12 +997,15 @@ def _init_from_list_like(self, data, index=None, columns=None): self._data.rangeindex = isinstance( columns, (range, pd.RangeIndex, cudf.RangeIndex) ) + self._data.label_dtype = getattr(columns, "dtype", None) @_cudf_nvtx_annotate def _init_from_dict_like( self, data, index=None, columns=None, nan_as_null=None ): + label_dtype = None if columns is not None: + label_dtype = getattr(columns, "dtype", None) # remove all entries in data that are not in columns, # inserting new empty columns for entries in columns that # are not in data @@ -1069,6 +1074,7 @@ def _init_from_dict_like( if isinstance(columns, pd.Index) else self._data._level_names ) + self._data.label_dtype = label_dtype @classmethod def _from_data( diff --git a/python/cudf/cudf/tests/test_dataframe.py b/python/cudf/cudf/tests/test_dataframe.py index 5677f97408a..74165731683 100644 --- a/python/cudf/cudf/tests/test_dataframe.py +++ b/python/cudf/cudf/tests/test_dataframe.py @@ -4566,6 +4566,17 @@ def test_dataframe_columns_returns_rangeindex_single_col(): assert_eq(result, expected) +@pytest.mark.parametrize("dtype", ["int64", "datetime64[ns]", "int8"]) +@pytest.mark.parametrize("idx_data", [[], [1, 2]]) +@pytest.mark.parametrize("data", [None, [], {}]) +def test_dataframe_columns_empty_data_preserves_dtype(dtype, idx_data, data): + result = cudf.DataFrame( + data, columns=cudf.Index(idx_data, dtype=dtype) + ).columns + expected = pd.Index(idx_data, dtype=dtype) + assert_eq(result, expected) + + @pytest.mark.parametrize( "data", [ From cfa2d513667edabda6c4487f15f251f757f0c94d Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Wed, 22 Nov 2023 10:21:52 -0500 Subject: [PATCH 35/44] Expose stream parameter in public nvtext APIs (#14456) Add stream parameter to public APIs: - `nvtext::is_letter()` - `nvtext::porter_stemmer_measure` - `nvtext::edit_distance()` - `nvtext::edit_distance_matrix()` Also cleaned up some of the doxygen comments and added stream gtests. Reference #13744 Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Mark Harris (https://github.com/harrism) URL: https://github.com/rapidsai/cudf/pull/14456 --- cpp/include/nvtext/edit_distance.hpp | 42 +++++++++-------- cpp/include/nvtext/stemmer.hpp | 46 +++++++++++-------- cpp/src/text/edit_distance.cu | 10 ++-- cpp/src/text/stemmer.cu | 20 ++++---- cpp/tests/CMakeLists.txt | 10 +++- cpp/tests/streams/text/edit_distance_test.cpp | 33 +++++++++++++ cpp/tests/streams/text/stemmer_test.cpp | 42 +++++++++++++++++ 7 files changed, 148 insertions(+), 55 deletions(-) create mode 100644 cpp/tests/streams/text/edit_distance_test.cpp create mode 100644 cpp/tests/streams/text/stemmer_test.cpp diff --git a/cpp/include/nvtext/edit_distance.hpp b/cpp/include/nvtext/edit_distance.hpp index 953ecf7734d..9a24662455b 100644 --- a/cpp/include/nvtext/edit_distance.hpp +++ b/cpp/include/nvtext/edit_distance.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -30,7 +30,7 @@ namespace nvtext { /** * @brief Compute the edit distance between individual strings in two strings columns. * - * The `output[i]` is the edit distance between `strings[i]` and `targets[i]`. + * The `output[i]` is the edit distance between `input[i]` and `targets[i]`. * This edit distance calculation uses the Levenshtein algorithm as documented here: * https://www.cuelogic.com/blog/the-levenshtein-algorithm * @@ -42,23 +42,25 @@ namespace nvtext { * d is now [1, 7, 0] * @endcode * - * Any null entries for either `strings` or `targets` is ignored and the edit distance + * Any null entries for either `input` or `targets` is ignored and the edit distance * is computed as though the null entry is an empty string. * - * The `targets.size()` must equal `strings.size()` unless `targets.size()==1`. - * In this case, all `strings` will be computed against the single `targets[0]` string. + * The `targets.size()` must equal `input.size()` unless `targets.size()==1`. + * In this case, all `input` will be computed against the single `targets[0]` string. * - * @throw cudf::logic_error if `targets.size() != strings.size()` and + * @throw cudf::logic_error if `targets.size() != input.size()` and * if `targets.size() != 1` * - * @param strings Strings column of input strings - * @param targets Strings to compute edit distance against `strings` - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return New strings columns of with replaced strings. + * @param input Strings column of input strings + * @param targets Strings to compute edit distance against `input` + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * @return New strings columns of with replaced strings */ std::unique_ptr edit_distance( - cudf::strings_column_view const& strings, + cudf::strings_column_view const& input, cudf::strings_column_view const& targets, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -67,7 +69,7 @@ std::unique_ptr edit_distance( * This uses the Levenshtein algorithm to calculate the edit distance between * two strings as documented here: https://www.cuelogic.com/blog/the-levenshtein-algorithm * - * The output is essentially a `strings.size() x strings.size()` square matrix of integers. + * The output is essentially a `input.size() x input.size()` square matrix of integers. * All values at diagonal `row == col` are 0 since the edit distance between two identical * strings is zero. All values above the diagonal are reflected below since the edit distance * calculation is also commutative. @@ -81,20 +83,22 @@ std::unique_ptr edit_distance( * [1, 2, 0]] * @endcode * - * Null entries for `strings` are ignored and the edit distance + * Null entries for `input` are ignored and the edit distance * is computed as though the null entry is an empty string. * - * The output is a lists column of size `strings.size()` and where each list item - * is `strings.size()` elements. + * The output is a lists column of size `input.size()` and where each list item + * is `input.size()` elements. * * @throw cudf::logic_error if `strings.size() == 1` * - * @param strings Strings column of input strings - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return New lists column of edit distance values. + * @param input Strings column of input strings + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * @return New lists column of edit distance values */ std::unique_ptr edit_distance_matrix( - cudf::strings_column_view const& strings, + cudf::strings_column_view const& input, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group diff --git a/cpp/include/nvtext/stemmer.hpp b/cpp/include/nvtext/stemmer.hpp index 0a57f8944d4..0e1759fdc5a 100644 --- a/cpp/include/nvtext/stemmer.hpp +++ b/cpp/include/nvtext/stemmer.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -56,7 +56,7 @@ enum class letter_type { * * A negative index value will check the character starting from the end * of each string. That is, for `character_index < 0` the letter checked for string - * `strings[i]` is at position `strings[i].length + index`. + * `input[i]` is at position `input[i].length + index`. * * @code{.pseudo} * Example: @@ -68,20 +68,22 @@ enum class letter_type { * A null input element at row `i` produces a corresponding null entry * for row `i` in the output column. * - * @param strings Strings column of words to measure. - * @param ltype Specify letter type to check. - * @param character_index The character position to check in each string. - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return New BOOL column. + * @param input Strings column of words to measure + * @param ltype Specify letter type to check + * @param character_index The character position to check in each string + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * @return New BOOL column */ std::unique_ptr is_letter( - cudf::strings_column_view const& strings, + cudf::strings_column_view const& input, letter_type ltype, cudf::size_type character_index, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** - * @brief Returns boolean column indicating if character at `indices[i]` of `strings[i]` + * @brief Returns boolean column indicating if character at `indices[i]` of `input[i]` * is a consonant or vowel. * * Determining consonants and vowels is described in the following @@ -116,19 +118,21 @@ std::unique_ptr is_letter( * A null input element at row `i` produces a corresponding null entry * for row `i` in the output column. * - * @throw cudf::logic_error if `indices.size() != strings.size()` + * @throw cudf::logic_error if `indices.size() != input.size()` * @throw cudf::logic_error if `indices` contain nulls. * - * @param strings Strings column of words to measure. - * @param ltype Specify letter type to check. - * @param indices The character positions to check in each string. - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return New BOOL column. + * @param input Strings column of words to measure + * @param ltype Specify letter type to check + * @param indices The character positions to check in each string + * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate the returned column's device memory + * @return New BOOL column */ std::unique_ptr is_letter( - cudf::strings_column_view const& strings, + cudf::strings_column_view const& input, letter_type ltype, cudf::column_view const& indices, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @@ -155,12 +159,14 @@ std::unique_ptr is_letter( * A null input element at row `i` produces a corresponding null entry * for row `i` in the output column. * - * @param strings Strings column of words to measure. - * @param mr Device memory resource used to allocate the returned column's device memory. - * @return New INT32 column of measure values. + * @param input Strings column of words to measure + * @param mr Device memory resource used to allocate the returned column's device memory + * @param stream CUDA stream used for device memory operations and kernel launches + * @return New INT32 column of measure values */ std::unique_ptr porter_stemmer_measure( - cudf::strings_column_view const& strings, + cudf::strings_column_view const& input, + rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); /** @} */ // end of group diff --git a/cpp/src/text/edit_distance.cu b/cpp/src/text/edit_distance.cu index 3d5f2d72e6f..a1d97409987 100644 --- a/cpp/src/text/edit_distance.cu +++ b/cpp/src/text/edit_distance.cu @@ -298,22 +298,24 @@ std::unique_ptr edit_distance_matrix(cudf::strings_column_view con /** * @copydoc nvtext::edit_distance */ -std::unique_ptr edit_distance(cudf::strings_column_view const& strings, +std::unique_ptr edit_distance(cudf::strings_column_view const& input, cudf::strings_column_view const& targets, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::edit_distance(strings, targets, cudf::get_default_stream(), mr); + return detail::edit_distance(input, targets, stream, mr); } /** * @copydoc nvtext::edit_distance_matrix */ -std::unique_ptr edit_distance_matrix(cudf::strings_column_view const& strings, +std::unique_ptr edit_distance_matrix(cudf::strings_column_view const& input, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::edit_distance_matrix(strings, cudf::get_default_stream(), mr); + return detail::edit_distance_matrix(input, stream, mr); } } // namespace nvtext diff --git a/cpp/src/text/stemmer.cu b/cpp/src/text/stemmer.cu index 2b2b8429d9c..bdcb0b2af32 100644 --- a/cpp/src/text/stemmer.cu +++ b/cpp/src/text/stemmer.cu @@ -250,36 +250,36 @@ std::unique_ptr is_letter(cudf::strings_column_view const& strings // external APIs -std::unique_ptr is_letter(cudf::strings_column_view const& strings, +std::unique_ptr is_letter(cudf::strings_column_view const& input, letter_type ltype, cudf::size_type character_index, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::is_letter(strings, - ltype, - thrust::make_constant_iterator(character_index), - cudf::get_default_stream(), - mr); + return detail::is_letter( + input, ltype, thrust::make_constant_iterator(character_index), stream, mr); } -std::unique_ptr is_letter(cudf::strings_column_view const& strings, +std::unique_ptr is_letter(cudf::strings_column_view const& input, letter_type ltype, cudf::column_view const& indices, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::is_letter(strings, ltype, indices, cudf::get_default_stream(), mr); + return detail::is_letter(input, ltype, indices, stream, mr); } /** * @copydoc nvtext::porter_stemmer_measure */ -std::unique_ptr porter_stemmer_measure(cudf::strings_column_view const& strings, +std::unique_ptr porter_stemmer_measure(cudf::strings_column_view const& input, + rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { CUDF_FUNC_RANGE(); - return detail::porter_stemmer_measure(strings, cudf::get_default_stream(), mr); + return detail::porter_stemmer_measure(input, stream, mr); } } // namespace nvtext diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 1be8566fb0f..b35c72b9e9d 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -664,8 +664,14 @@ ConfigureTest( testing ) ConfigureTest( - STREAM_TEXT_TEST streams/text/ngrams_test.cpp streams/text/replace_test.cpp - streams/text/tokenize_test.cpp STREAM_MODE testing + STREAM_TEXT_TEST + streams/text/edit_distance_test.cpp + streams/text/ngrams_test.cpp + streams/text/replace_test.cpp + streams/text/stemmer_test.cpp + streams/text/tokenize_test.cpp + STREAM_MODE + testing ) ConfigureTest(STREAM_UNARY_TEST streams/unary_test.cpp STREAM_MODE testing) diff --git a/cpp/tests/streams/text/edit_distance_test.cpp b/cpp/tests/streams/text/edit_distance_test.cpp new file mode 100644 index 00000000000..59206c39e69 --- /dev/null +++ b/cpp/tests/streams/text/edit_distance_test.cpp @@ -0,0 +1,33 @@ +/* + * Copyright (c) 2023, NVIDIA 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 + +#include +#include +#include + +class TextEditDistanceTest : public cudf::test::BaseFixture {}; + +TEST_F(TextEditDistanceTest, EditDistance) +{ + auto const input = cudf::test::strings_column_wrapper({"dog", "cat", "mouse", "pupper"}); + auto const input_view = cudf::strings_column_view(input); + auto const target = cudf::test::strings_column_wrapper({"hog", "cake", "house", "puppy"}); + auto const target_view = cudf::strings_column_view(target); + nvtext::edit_distance(input_view, target_view, cudf::test::get_default_stream()); + nvtext::edit_distance_matrix(input_view, cudf::test::get_default_stream()); +} diff --git a/cpp/tests/streams/text/stemmer_test.cpp b/cpp/tests/streams/text/stemmer_test.cpp new file mode 100644 index 00000000000..7aa51befa73 --- /dev/null +++ b/cpp/tests/streams/text/stemmer_test.cpp @@ -0,0 +1,42 @@ +/* + * Copyright (c) 2023, NVIDIA 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 + +#include +#include +#include + +class TextStemmerTest : public cudf::test::BaseFixture {}; + +TEST_F(TextStemmerTest, IsLetter) +{ + auto const input = + cudf::test::strings_column_wrapper({"abbey", "normal", "creates", "yearly", "trouble"}); + auto const view = cudf::strings_column_view(input); + auto const delimiter = cudf::string_scalar{" ", true, cudf::test::get_default_stream()}; + nvtext::is_letter(view, nvtext::letter_type::VOWEL, 0, cudf::test::get_default_stream()); + auto const indices = cudf::test::fixed_width_column_wrapper({0, 1, 3, 5, 4}); + nvtext::is_letter(view, nvtext::letter_type::VOWEL, indices, cudf::test::get_default_stream()); +} + +TEST_F(TextStemmerTest, Porter) +{ + auto const input = + cudf::test::strings_column_wrapper({"abbey", "normal", "creates", "yearly", "trouble"}); + auto const view = cudf::strings_column_view(input); + nvtext::porter_stemmer_measure(view, cudf::test::get_default_stream()); +} From f02fde9de9354a829d6f4425e086c84d36c076ae Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Wed, 22 Nov 2023 18:26:18 +0000 Subject: [PATCH 36/44] Correct dtype of count aggregations on empty dataframes (#14473) A count aggregation should always return an int64 column, even if the grouped dataframe is empty. Previously we did not do this because the short-circuiting for empty inputs was hit before handling the count case. Fix this by reordering the conditions. - Closes https://github.com/rapidsai/cudf/issues/14200 Authors: - Lawrence Mitchell (https://github.com/wence-) Approvers: - Bradley Dice (https://github.com/bdice) - GALI PREM SAGAR (https://github.com/galipremsagar) URL: https://github.com/rapidsai/cudf/pull/14473 --- python/cudf/cudf/core/groupby/groupby.py | 6 +++--- python/cudf/cudf/tests/groupby/test_agg.py | 18 ++++++++++++++++++ 2 files changed, 21 insertions(+), 3 deletions(-) create mode 100644 python/cudf/cudf/tests/groupby/test_agg.py diff --git a/python/cudf/cudf/core/groupby/groupby.py b/python/cudf/cudf/core/groupby/groupby.py index c48e5109ff2..73e6774f5ce 100644 --- a/python/cudf/cudf/core/groupby/groupby.py +++ b/python/cudf/cudf/core/groupby/groupby.py @@ -592,7 +592,9 @@ def agg(self, func): # Structs lose their labels which we reconstruct here col = col._with_type_metadata(cudf.ListDtype(orig_dtype)) - if ( + if agg_kind in {"COUNT", "SIZE"}: + data[key] = col.astype("int64") + elif ( self.obj.empty and ( isinstance(agg_name, str) @@ -609,8 +611,6 @@ def agg(self, func): ) ): data[key] = col.astype(orig_dtype) - elif agg_kind in {"COUNT", "SIZE"}: - data[key] = col.astype("int64") else: data[key] = col data = ColumnAccessor(data, multiindex=multilevel) diff --git a/python/cudf/cudf/tests/groupby/test_agg.py b/python/cudf/cudf/tests/groupby/test_agg.py new file mode 100644 index 00000000000..7919ee4a9f1 --- /dev/null +++ b/python/cudf/cudf/tests/groupby/test_agg.py @@ -0,0 +1,18 @@ +# Copyright (c) 2023, NVIDIA CORPORATION. +import numpy as np +import pytest + +import cudf + + +@pytest.mark.parametrize( + "empty", + [True, False], + ids=["empty", "nonempty"], +) +def test_agg_count_dtype(empty): + df = cudf.DataFrame({"a": [1, 2, 1], "c": ["a", "b", "c"]}) + if empty: + df = df.iloc[:0] + result = df.groupby("a").agg({"c": "count"}) + assert result["c"].dtype == np.dtype("int64") From ee3fb396d2933180c8f9df10d1d97ffe0042f4fa Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Wed, 22 Nov 2023 17:06:48 -0500 Subject: [PATCH 37/44] Remove unsanitized null from input strings column in rank_tests.cpp (#14475) Removes a non-empty null entry from a test strings column utility in `rank_tests.cpp`. Behavior with unsanitized nulls in at best UB and should not be included in unit tests. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Bradley Dice (https://github.com/bdice) - Vyas Ramasubramani (https://github.com/vyasr) URL: https://github.com/rapidsai/cudf/pull/14475 --- cpp/tests/reductions/rank_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/reductions/rank_tests.cpp b/cpp/tests/reductions/rank_tests.cpp index 4cc8d7e6600..3ab1fc01eaa 100644 --- a/cpp/tests/reductions/rank_tests.cpp +++ b/cpp/tests/reductions/rank_tests.cpp @@ -140,7 +140,7 @@ auto make_input_column() auto make_strings_column() { return cudf::test::strings_column_wrapper{ - {"0a", "0a", "2a", "2a", "3b", "5", "6c", "6c", "6c", "9", "9", "10d"}, + {"0a", "0a", "2a", "2a", "3b", "5", "6c", "6c", "", "9", "9", "10d"}, cudf::test::iterators::null_at(8)}; } From 168533a8ad4086bd020be4f7bf9264a08b6d2243 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Wed, 22 Nov 2023 17:26:32 -0600 Subject: [PATCH 38/44] Update README links with redirects. (#14378) A couple links in the README redirect to other pages. This PR replaces those links with the destination to which they redirect. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Ashwin Srinath (https://github.com/shwina) URL: https://github.com/rapidsai/cudf/pull/14378 --- README.md | 22 ++++------------------ img/GDF_community.png | Bin 81567 -> 0 bytes img/goai_logo.png | Bin 30695 -> 0 bytes img/rapids_arrow.png | Bin 192477 -> 0 bytes python/custreamz/README.md | 2 +- 5 files changed, 5 insertions(+), 19 deletions(-) delete mode 100644 img/GDF_community.png delete mode 100644 img/goai_logo.png delete mode 100644 img/rapids_arrow.png diff --git a/README.md b/README.md index 6dba012b9c5..2a2fa84f33a 100644 --- a/README.md +++ b/README.md @@ -51,10 +51,10 @@ print(tips_df.groupby("size").tip_percentage.mean()) ## Resources - [Try cudf.pandas now](https://nvda.ws/rapids-cudf): Explore `cudf.pandas` on a free GPU enabled instance on Google Colab! -- [Install](https://rapids.ai/start.html): Instructions for installing cuDF and other [RAPIDS](https://rapids.ai) libraries. +- [Install](https://docs.rapids.ai/install): Instructions for installing cuDF and other [RAPIDS](https://rapids.ai) libraries. - [cudf (Python) documentation](https://docs.rapids.ai/api/cudf/stable/) - [libcudf (C++/CUDA) documentation](https://docs.rapids.ai/api/libcudf/stable/) -- [RAPIDS Community](https://rapids.ai/community.html): Get help, contribute, and collaborate. +- [RAPIDS Community](https://rapids.ai/learn-more/#get-involved): Get help, contribute, and collaborate. ## Installation @@ -66,7 +66,7 @@ print(tips_df.groupby("size").tip_percentage.mean()) ### Conda -cuDF can be installed with conda (via [miniconda](https://conda.io/miniconda.html) or the full [Anaconda distribution](https://www.anaconda.com/download)) from the `rapidsai` channel: +cuDF can be installed with conda (via [miniconda](https://docs.conda.io/projects/miniconda/en/latest/) or the full [Anaconda distribution](https://www.anaconda.com/download) from the `rapidsai` channel: ```bash conda install -c rapidsai -c conda-forge -c nvidia \ @@ -78,7 +78,7 @@ of our latest development branch. Note: cuDF is supported only on Linux, and with Python versions 3.9 and later. -See the [Get RAPIDS version picker](https://rapids.ai/start.html) for more OS and version info. +See the [RAPIDS installation guide](https://docs.rapids.ai/install) for more OS and version info. ## Build/Install from Source See build [instructions](CONTRIBUTING.md#setting-up-your-build-environment). @@ -86,17 +86,3 @@ See build [instructions](CONTRIBUTING.md#setting-up-your-build-environment). ## Contributing Please see our [guide for contributing to cuDF](CONTRIBUTING.md). - -## Contact - -Find out more details on the [RAPIDS site](https://rapids.ai/community.html) - -##
Open GPU Data Science - -The RAPIDS suite of open source software libraries aim to enable execution of end-to-end data science and analytics pipelines entirely on GPUs. It relies on NVIDIA® CUDA® primitives for low-level compute optimization, but exposing that GPU parallelism and high-bandwidth memory speed through user-friendly Python interfaces. - -

- -### Apache Arrow on GPU - -The GPU version of [Apache Arrow](https://arrow.apache.org/) is a common API that enables efficient interchange of tabular data between processes running on the GPU. End-to-end computation on the GPU avoids unnecessary copying and converting of data off the GPU, reducing compute time and cost for high-performance analytics common in artificial intelligence workloads. As the name implies, cuDF uses the Apache Arrow columnar data format on the GPU. Currently, a subset of the features in Apache Arrow are supported. diff --git a/img/GDF_community.png b/img/GDF_community.png deleted file mode 100644 index 69c5edee6c2e4fdc9c8b9ee8a99b810ff6faadf5..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 81567 zcmc$`Wmr{R*9Hn(Kw3(W-gGz8u}SIflI{?YmJ+19yHi3sq@^1q1PMt&B%~Y8+TioP z-}8Mx&iQw)Yl$pi&o$?mW8C8&F;}RP!Yedn0%RB%7&K`qaTOStC%G^%pkPD<;3se9 z73zWiV4YQ7iNaKl5&Z&gklsjXJHx=>VL|`FT4a9(J_3P}78g--huzJ5YN)0@+kYRA zj3lTQyWAq9Twda53@ak`<#U){<8%ZA7wA362)hClz_`D~Lf4lmIYwDw&@-0G{> z0UVZ8>WJGI3daZ~M|>)`_fOxK8*RE{ZRh7l=KIa(i_WkJSYX)yPge*xD9Qad#sM{Lkz)Lwf^az-v{*aE=<}t)m`J;k%?ihYQO&RV5P}mRyKUWMy&|iW7}s!d#vLfT zut+g5|GpsQ7&%x342ofUl>cke{ybm^6(>3gc?$$Y`R}V*)KA8;T!5}29!|>dU)P!+ zOyUHOd&n;6-;aSo>OqF3&#gWxm{a)QvI%(tQWITYA*(_8Uyq^0V1O$L>>G+V=9m8W zW2C_7h)PJ20l@>z|35CsKS9fic0n2{oiR8@si0nCk>-jK-tT+$@X}H>1q4=`4Jva= z`&_T<#LpZYX>?}z2@7Y*TWgP!EUj=CL_i{xh9b9V6_a%u&(;s#IjmX}F)vde*@=`> z2uaK+*aQvp3!1&P9~|SZ^8C?yu}kYE0iwkC1{KNay4jW#*asru|5k9 zj>Gx)Np$UiH>wBG5dxoX1wagEBS&e%UqdT?XNFIuZLYa%!zNBh38T)Eii2=7X4YIY zmUneHu{Svc3md|BO$vkmT=vl%j{fZEvSx6UULTbP{i1UTTy=;Dv4a;@K@toN(gfDZ zChqsMiS49nP7T;eO;Df?yC(`!LQl!O5_aA1Koci)3O$|u>=1bMz zt4C4sj!RaPg3>r`3G?hw9S9)?=^;gCTJ$Ex!5}q}O;mQl=*^vpylY=|4>+wUBJJ$; zptn7KbXmZ3PT&w4ZaHQT?upW$`4V0YX*~fbFf1%2yLob4(tu@N8=Ow+R~ zXgmK+>R7&;li z6>C&;mff}-E>KkBtaE(B`8NdjSPuGs@u;*8yJt^l+?TXACU(#0zm~LiLbK&NsVJYm zBd*xVirhc)m=OAHOsV^|%KC@iFHMPsx^mC!J*tD|A^U_$>E)KlIght=K1=-lhuTWb z=+|wYfA+nHTr3T)QI!{4at$Jr=5*@_7EpZHN0Lu(PpLRd=FtNdgw0JWE^NKTD}TAn z;kk~80k$Rt08@8=+RIoNq`;8Qr)3A-e=6@G<3;mYE%&&Li+3#9);rUGRO(6=G!Z;x z1?Y46O4J;(2!0NM31||OTAu3Kc$Ct*m9Z9d!rt`a+rn_K)cdyD zBA%%GhE1?ADL=x1vvdYtpFltAu6Q}>T=U2Ji-H!@H8do+SV#*3UZ{&GvDlg-zX3P3 z$VkX1)%@ml;v2JXa?vs+?Zuw9s-Z-7rkbJ*0)qxNay^WApXS8)1x;vU&Rn;zr*)`< zJh=w@7{XaFKOEwD#T~Y2PQ}t>&}y50wZN_v-V-d@{2+MapxUcU=jof&s;l|-s%+S+ z(o1i_Lf>@P?bdVn2dl#<&*6=(EzACecNlFJMN3DRdA%3k+KK4!@2zP+uJu11@~W&3 zdpyY}L)^axT_SE9uiEjFiMvox#*-M;APmlHf=Drt4Wx`4%E))Rb08>Yuu2%-cEMjb zO_}o6m=RU(-b9q^ytDJ{{O$1Miu1lrSzJ@*ZBB)OFh4iZ`-an}1M>yaU$aUjuX5KH z{v0PYR-OJ4HOcYtDZ7qMiVO8AUSmo4or9SN+EdT9qgR9(UwZre7q>>+QB_7nS$f*s-)c;-gBr$a5vTaZGK@nW}f={B!Pq z7#6>7sjjy>91||l=N8Vcr(D^&0lv+AW6D>}^_h14;IF&1>a{$Es$sW|fV{7_k|e~G zzExKSyeKZN cvrq0b`u8I=MLSZZVjn6*v6FYf`fV!wB2BeXeUe^iDt=%SR>iD4E zYqR6%qWf;qwpd>&Hn)Il9(PexK_&U)Tc}67q zv*;af4N0%YBNsEr*-vgfI7(Y`c=g@`j{Cn_>(_~qO9Z3a`$;RjCz@uS&xdbF(Rl%W z!n@D9blJDs=zTcIdOW{lPo~T{6C% ze_Lvc=_f}NF8Mf-EiDuETC=5EfsOhcdrxN*!wMHb|800=tP$ymLozo}ud5Rv?Eo^yv@{7N22g)mD4zpwwVWstXmtC1*ZmT(@Z$+6Qsbn&h{_$=!aRGK;LA|}v5J2QG(Mpxc|s@xY=9i-1N+?YYsY-D z8vOv+AnYu|p=exf{8DX6RkdnH>|zFjAh*kw)jicufZ-5P-+e;6yFONGWuf+TU>;36 zM&7cH<2gTRH8wF_5vIfd`+F;(`+T1k7KT_&YTvle8*Q`66YTa+x?5XGT6k?#;vX1#F@e5RcBM2jr9gQZ%@Cz%Q}PfBXOCk{Qb_XrP0ZKRm$eaT5^{n z*A=Y~-wYMeV?RDAit4+gT6RC>!8@JM*DxT@esbz)Rvea}`;lS%?Sls30HIkd`M4(# zWr-@RJYhRt@IKLuKAnqxdHC~1$>$lg=r3X!6Ri2zy8+*qUNI0A)qQk}N+HoqnP*-; z)Ctjr98#`i`jOzBOoeuIbMm&FeLYk3WLY7bxy!PyDRz6w)Z7sKH8*-$x2~^5{M_8@ zyN&)QNd}Yq2J^bDZ#LI03pwVx(!%ojH^@I6u4$r@cDaU1)(cp;6}T!DZMb9@$rrOd zb~NEf2rB(i319q8$r|Ln^$8hIj>w@WSIvN3IYMMeBkE2vm>H;eKu26%2PZe*hjw;gO2^@ z`#41pY00&-+}`OQZX`I-mEWVP)s1hSJALk;CMhEvYOc3)R9-kFk=WU}%d`8`RM63} zqj(|qgC%lijDb%ua+qJpY`3hAR6cJxd)(g9qDMpQhk;`)ZTA5|5&{4$h~AA+1J_#o z6{Y#|{LzeJcSnT!SB2cp8cPi3t!R_YlV+AWWrX%v2eSk+8|6mcZhNcmXY(cey~lli zAD&-5IqpTha>;NUI(TD0EWEmv=^#S0{^HSdBSJlQ7Q%ukAVDmUC&`Zpvu~?fzjEo1 z?psV>v@hRq_ao=g=fL;!S1TcOsdntq1#E@g8mJ=&@rVlAw>SL}c-L<9p2*a9I@fOp z-|zFd`qMUt8z=ea{bJeSSbT1yFT_@-sb{gbB&=$z>np`k#mhqy*` z&`|r=pJe+;z40(PWgk>GM+`EjCyzN9yewvsu++Z_{rE$fbydV6wD2bdr9nQp{fjWJ zKKQ_-S3BRj%}}If71On^<~Q2GI5Yg0clob|TmN)x_gRGc_*h<_Y~$Q|7q}<)2s0eK z7T$F`dSUC3S9#L@dBe%0*nv&RVRtNJ7r8_s2mlKF5iOt0HmUj!VJ+_tqW-!cJkF;( zl`Qz>ADMa{>9Qn9fR;I;O8YrdU_j@#$3W7a1V^{BX8Ft{Pi?q+>>^7*t)O~Z<(AEk zc93V$_J00wSwS#Io>c`2ySp4Vc2BQl{pT$vhWsZdFP^%X8Vg7D>ZZpr{pz!}l-ADO z#nk<6UO)A|JZ;akP)*PU*9-gkv6a&c`MU&JR`w)5tb>EK=(2g6q z@UHp-y??LUuwM!+1!oyWNEw3zjeKw6cD-np>eZrU|0L~K%9X6S66^mY*Lk{i9Y(L= zr-#+LSYm9YX}n}n+pM4Hn4MfriJ=Gbd>d0~pN8Z~B)z)YW*!z)AqTc&tj9IB>H1Pt zpZFYzqGKSGh!aHISku}x|2scONeeh3q^p&AJE%+gpH~PR)IH^l@%^79pa{UgHo@G0=iEw0{n)ur1t;hC#94CM8-bps$GCW z%>TSFkf2yClIa2@u>O^YguKVf!Ai?{mbCE--@(T{(n-ZSP1|_fb;&d z%k}TEV35S$1HX|{vdH|`!-0v!vn`DQU{Tz06H(f1^^!4FVt0OIUNV{w=T=x>Z6oGmV6 zgVvvg0@E8tN&5##NRpPI|t! zc_Tnnkue5PU;hcr7ko(d9E?r)xQh${xwwfoHL$(2qSLW%wFOQD2aoCoq5$Xys}G|* zbn9oZFgO+^@@&rz;|^m`-Cr}(bvYjh1!fu?v4sM31OlSSg4CXo*(ge{rPCnd+WIHL67HYhZgN&tMDkf2OQrNYGk++zJRCXF0ZpRgE z<~*<~@4zv)l%J`KS9}Jqbd@1f9~o5ayPwT$8u2|Si61ZMt(<8Tp)Az4fV^f8N2TKR zkZC=qypMfCn-zmNekO+#`Pg|gmO1dmH?oi)6yDi!s@03`9q*JN>MC{j(M;YmIiJ1gQFMN3v%6$I#H)x->^ads0K(aHvYvT5T? zp5hYT=`^5x`l;B<7LA_@0U^nOuMMWJv^3K&G30zj9^3$;0FV9`w!UWYWV-T|$zUk1!L#=fHB)MBLz zhN|@kjj+^*hcuqV#;-{d2T*g0gsQb^wBc{ZdYho=#fkn zMTJvC@R_K0C<^@Zf)Tj2%HJiwmreKz9icN3QYGA|Uw9n#fLNViR?25CJ7z06Jkm&c z^g8hS-_}CihEA$eDClZZx*;G(iB58c?Bwu1ULDWp?BkP1Yg!c_N|LjK^BpCkGl4_| z*q;_kp~FZJ(xFcN2VBoPjp&eV+#rL>xWvg~l!moakMUd{)YEEYEJS?&U4a5Bi-Gy9 z&L>BAZ;`YmXyAtL_6pY$N;FG;=e~m)o0KU41@ark3dLtmt5z8rRcfshU`V>uCVtQX369i}M}zt!{|Ris4SDvJ12r2zI3l=muA7yn(9RDjgbGZ$ zfwG&ZJx?%8vee`JbV1HXRSueD2%)8&&h{LEbTzGC0$ft8O@Ft$}6U?Fw42!8^sBKRA<|`)&S#a^;CXya#Q(ty}nUdG?0GDd}1r`BJd4 z$f#2}t<1<~z;_1ZhnRkw0@5!&c`i3c4$P5<2Ham*o-E?J@Ww@Tgs4$xm{q`605`Zy zVO9*}k1>!H=rbtou8BZ5`XqK;QQJdtR-QBx{Vu`C;L+zg1BAxrA8@n+vCp9`v@5Z7 z+&?dE2VRdINXVZ=9P_`CDkxp<9FhJRk_|B?C};0t29ML--&mR_vIW%qrU4 zJiqoPS7PR@#9%?AeeyC2e?O#+=q_f*9GBc1;trfD)*odn0C(1*y6cHG!Jq6i(UTF#FlxBBK}&J~bK3Z;@k#On>pFY9EsH*AOOSL%qHs+>rH ziI7tGV>72hTGPa#>n7K>*5tpA(TEJNcE%KKB&gP9Q0A2JbA`LZpea{rDREOF)dp|m zDujH@(ZWI`^v{h~;vaoKf8|E}QKr|)@Ng|ziGcnxLG}@Q0qRu3@ND@3+T9TYmrp)9 zx5=&Z;ChNY8^<?Gp>Gs(80AQu3eOyQ<5|CW zGrW)mP0E5Tkb^?!uP9b6a|0X@CUV5c^7Kl9G~B{$=4a8)k(Vr<#EHRS$~B>bM~WmU7R+i*E*XZu zpW?UEu!1}PgSUPp%qK6RvP2=N-S>(_BRTpZ5{H1gr~tKV5}khMTe9c%IIu1i*a}gJirt#SuI7eX63{jHaka zXqOQgesid0!SqEtFJ$ba;St8O3QR@h){=Yw{BQ<)QE800oHgS0REv^Fl1A z{79w*t!#K`QW0T{)IFZVt(AwKzE_}`MPjNVK zpcht6lNNbby_TfM$AR5Esyic2C4=r@m|*Lj`m*qrxWQu}6&L|EyWB!UsC-MZAa`2$ zIN5(T=$>pe=lNyPIhK{NAdR>q4)>bupQroUqh0w$k#KURu6qd58n{vI2!h{X!UK!F(-N4zZ|R;4@Z$+X(OnqtO)z> z2RQl{5^)sct$kMpJ?nB)^6P zE5o8KbL3J&B@G19`~QuK$-q^{LT7FrFNRHHvKfk^elqGa^5ic-c9YwRP{bz0!{iJ> zyz!3Caq#yRxc+=AHQF&0*N_e_hJek-CFTVHTC(kqZVL}pNw^Q4N8?ou*ASd=a=wXj z2pxaPHJA~jm`QhJ$dDK74Rwql{B7?E{p-615Wd@uN1tWu24nIA%>x{qbmyX+hwWKn z1vFmv4mq57bd=@-D&HQB&!Dq`c3l9?V@FxpZpk1z0owkN!4r}>SLIQ6(R%j z=Rph^La$bajsyYGAeCR>5PnAcVx)|~EQtSmD-j+qyf)vV=hvj9uq?aRd#o3{m|HYd zt4xuTCDH_&vYxPk>q9_6j~I~_!ENfWTyO4CqW$TV=W0V( zRAQR3TbwTWnnwqW1}MCpX~X!@IhXU5<;H`Dd6>kcJU9k_XnP-D_itzhR%*a*92hlz z6_}03ZZZ$%C9Zv{C{mOykt*nqX!;{OEzfTZT?zKKunL2l5`)V%5$cApq5vfd^xD{E z1;kE{HIN-<;H+9V&tg@$aa7X)IFqOVO^CgUWtxwx9g`Fdw3YQosbcR zkhgs|9J)_78b*Vzi=w+vc*MH*A2~}C8Tbd@Ut|fD1iK*{B(eCVkDTe^1a&vCpQImx zXHycpoDM}+c91O7)80!Qolke#M)zP&{$2~&=BwS*ehl=$4TvdO)e#l;H4hX5+GAG! zp-Bv9@E~*jra4Ld(`V^YcLxAhE>xpD*2DeyKFY}Az7qGR&EdLfD#K{(!^1xddf+de zgaDge-hN~=YY7k^5*Zm8AZni^hJ9P@yoB1yylJJ54b!8=OfQ>#A6xu`!cSA6Y2|DyXPK z;Jq@h*`xTypTat-op4`@HZ9oNn7%wM=?$wbitI*)4#tJJArKmsgu^)P-BqM(+dVwh zw}K03M$Chlz|iCjcM&Nm;REUJA_GSm?FXa1PD`Xlwq;7^DQO0%UOKRKlZp^x#9aH8g^3=kptTZEGbk?OxNsPsLc= zl}pvS)aq#=YE5G%(&p%o%oaseg~>YB^YyX+2VRY{0n3>44SHMwU3=-i3SA?mrM6 zrdX|-fohkaK3X=4j3cqRYmaAA&m*W~uYoWufd@O!A)u#+(FXR#JL(uT?6#ZgUwsl^ zsQr*2lmbUVldAc_gOv~4$ch+}ykNG3Kqi(j=qxu}yHdj(@Mm>Mv3eLV%LrspK6PID z3$5|zdmKx)0Yi4~_I)n|tQrJ{q2U7{8$6)u-1qI+)B$RMU|P5#yj79;j`LD&qr6V@ z0xj0-ApuTKKp5nTD<1`>RfARY{ASsg2prw(TM*AvnaAHhLZRa;K;VmlDy|{WJp|Fj zWIrpbM{HhWa}$I4s$*tO?h7GN0mXFX#gc>3i>yrz(`a+{M!7_Y^-H_U;MXpO%2AV` zktYDx!T^S>AEeg2dcY7I&|23C#DgvxM$LP+8W!+%YV1~9q9OTBr92a_X!1!%ZUw8&ao-+L4}lzKy{ zWd94(O^Z2u9R!r150EkOleo};^G%u=+YiP;*9o*ais{gR-~q0Ic=Tt};qAM5{)e%U z{jDV2<#14@`fUcVtUaBumPfN9ibSk1J_H8(@x|3phEk9_GMY7jTrAD>10BtWYyYa$ z4U}x(kEaMeWC=TwHga#?U{*eVB>^ZDwnUvr4g7)W7TG8uRup1b7cf;+fq4k%`}em@ zYXdUGI;AEIQ|wxQ?QI`O-$!1*8bN>X)o&0-hQ&luN8hNa8EWLw=2_5u>p2FVtGA8F zIodOI^o|)hC+JI!lU^Zu1$09nmxTlY!ffmnvfA*6U1J2*6g&GiifA!ki@gf{LVkZu zYFb#~7VSURa>K!R)WInm_#Q~!>Cr@mP$Q7*5@G7|I}_o;kq4MP+AkKA)AYEH?NU8p zQzU>*Iuy(HvJ$DSb9G862yV?UUWcf(^a3i5itL|(w4F|Q?P4a=b>mUF>27?bRr%)W zroR#*^uv@efQD0dikag*c-=X;qUbKTdD1Hd7$=}@ZgPQwdUYBAMN2<8NkM;V>A@h7 zPIoyiFUs5fIurM4fEGNerIY}mme>cWwfzqZL;9+F{P@PPJ^%QA&att zAFy3we0tKZ9%U%qy-)+P_9Ta7-4m3l6KRC5)3ZMk5-{v*K;ox~ILD7R*o1Lm%84vH z%qoY2ITd9oj=4vH?r7kIIg(z5u{}pJE-kTzprkEgMHjRg-M#u>00%*f8-C9+^&iNZ zA*Q>90GHNKEJIkyFACnCmdOJE{@V@fL!Ca-#PGv64j^ya&1}8R7Rzds*MRt5j|FIL z3OzM>6u49;2P_};(MS;8<9-TZV6Zd-)TU2anVRL0tfiRoiI-mBr{wW|Nk#k@q(q=* zY{v%L91gSR-)~-Y2EDR~=pF8|!YqVB5+HK|Sk?VgY&5LvN9Uj<=)XAVu8NcXLoIts z44v?7p?sUeTY$@Bplcfj$y7M%=b+Vva+a!UZJsUH2v)zErmTlHh8chc57#|3TkmVm z_z?6_t_1nH(AeX`Z_;o}w@YZpP}w*i=%)sw{y=6;7KRJD!Td#M<<1R(fBT^x@tbs; z8syf0y!CNa%s}}6d1VC+I>#8GXZBLnaYQF-H7{iUC4(4FK~`$2`wXz-$8tFavaPh{ zO}7GDCDgq^0X5G~K0@mi)HLz`h>HP$1cO7?*hkM!3R_+b`6DhO=2Ob3hk2sJT&gdhgt z-k&5-s$Xj?F}ua@swU9p;kmZF2-kQgS=diwXu!Uy&iUzX^=Bag z^`kC9s}I{L0~Tr{Qc#9Ux$+EFd-R>eh>gQZv0Py*z>c@i$6l0I;Ed~%0{L$7;SI`e zOl1nv@TOrd@9&loPRX4-k3z>{73(}HFG(HJdSowPM?|9KF#aCK^J^GSJ(lP04%?K| z=n%;-xTEcOFXQ9cbg4lj3;^kQSBJq(#VGEY@#xJd0htHEF?@Jj--s7_@0yr;^wFGE zQV^|@K0Bkk{?L?DzhcPX=${%0|I)be*!=nP%z@K>XyTFT zr11Y`Qk~yQL}4mX*JDy-jU|prOdaKN^l`Aylm=k&tU67pLbd9|vY`}tu?IpCqYB0z z4+jJQnTU`9xlNHSDY8eM$U+Qp&b3>5YFtb!Attz43iL;IUFttw!;7u|__;MSD~kg^ zfJxe2T|;D4k#phAK}2;*>^9@P?}L>5Copn6R4toS9tt5b5GoKray^vA56W@&t}D=9 z4Gq1~xC_ww<(1nrM2Bq2KPiA0dg$a`E7%BopD9Gp!L6J1H;R+W;RF4&kS>cKdFTZK z^!^Qbl`HXn;=(&ZQZW8Tf?GrT!-X47qDOSJ>ugB$K z&;c+-A(3_ZfaUKowQjIvqVdReHNH2b?R`!;sQVP_Z&f3f@Z)tQ1`{3rMf9Iv`viN} znY&Mx?3nugzz51jD?igui}yz| zAMQ8dWpnY7(tylhSI;ZypzCGcvo@qZ@BHiXC|4Mu*zzB|147iH%ei@zs#O=4KyClL zkgO?h3=O7x9EOphkYp#o*pO3|?}}-2i*5x%K%QG8Y(Q>CZtmHm#J8ypuTU61xAe|F zYP>iS;8dL6OV4~P{5VMR<}`Jy%VCskX!h9HgXr$BAxh?rO{yK(bkjKH*X^@`xH z04sPkoJ)tHus3du&V~z*!x%?H`JtD&EgMu-d*Jh5Ly^7_2=_?Ow8X^MV=GW!jBE0yGSH)J00eAeclYj_H3)juuWFmQP6| zs!Uuf*UBGPZ;lgNF`bz-0KOM<@;(&g^I^Uy9He8*)Q`w!wrmD)$A3UT8(NP{BKms| z75I2n@in9zv(UOg+kb6kWdr~it4oWE-tq`7x?~O{3LxceaAmvvDFCf;gebhZ`m4p? zQ$^r*DBb+DHTaMVhR9)IY&sUw#+K%3WI8lF;X(;|+k!~?DJ7cN+)Ui7IfG_mDenC< z@oPEm1r2R@Z`Z3Aa-9EFaD^w_W__5w1wE~hU7!Yn@}Gbz2GR*Y6hcG{=rIO; zF9G%?LQhUIP8nceWl+VbcaG!ulQDl3<>Y}l=4OhqL3!RQ{_C6JtA(3ZjNZGH3e@hc z4Wn!zzRr9=JU_Vi3SgesRU!_!57`gE9n$-|k@V1OW8PNKDjan*gOL|1cvat|H-GIM z804`~D6;TSp?#VVi7Y6#tLCqUU@!s(9v~6t z^mnq7hCg{mFrS}HLICQ%==@Rcia{(&xcTM$dlvC0?V_WGw@e4r_l1Mf7fP+)qg>;E zzNB{pDc}XmesUjG^Il5;Qoj{fwj+hwMInnSSX8iy&&PMxYQ&$OEoMQAWHu!R97PO!8ixlULk>ipDz zXWn+R$)`{3r)G?IMfl%Rf>b9e?B()5tjT|y20S$XByU=tZ9zplk+D|Gaj@?2LT^+* z-E4(U>mz;9*pYfJzd;B>?eUud-sNN)lV~ey7C*$L?p8C)hNS z_hpq4njTyNLf)t(he^R!|7fE$)ZIvDyFA+Sg#-j=vJvpbo0??bMy^d0--HaLFd}b% zsFdD67_w^;*Dw)w%--SEcs`Usd9%XXKZv6&j+aa1HZWLFyBR(|^>KF)L{cd}c?Z;5 zbK@AvAKKWVn#c#u3Gij3=su)f|8hXAJ^GkqWov##{-%7PH208^1;d`h9RaQidGuL(Bd`h%R<@Y}h388upZjvwh1(25K#s?mD0+TuRaQ40`| z3-hPVJ@*{LY{gNgCWIzQ?xaS1up)al8Y_g&Mt(9it|-sfUits04+Ah*=qbC!9&OTx zK_E(qBQL_YBqqv4Ff}q`aDJS-eWacL@$1-UV&~&onteAvhvd5SCy$gzb_i5NB>`TQ zAc(>B&@my^=>`gJSaCi4vL63b=SQK&DBYIpm6=XD4g!_hRat(PY+cW7k<;G99Tyq^} z^HVzL8$&`9r(H?%9{(sJI-rQ%BB?Hq0Z3qipL_5UO_ZN&@JB3Vdeyk}+-Q~~XO&6y zqnSghPpJZHw7+tC3EC;FN9tnD9-&m~U->yH@D9?3Tu%%q3T~LwEO|oY={zO(jKh?^ zZ~77j-G%wIX~a<@`1n!pV<4Q+Of{OKzujZZ!YTo_VDAgX1`B2tWYNL;;2h1YSG=nZ zauqJA)5h6PI~v!`FEVs;<%+xAMyUNu{UJR_1>&vnwvU7as4hvb3<;Q;ny;xV_eox9 zbH?ln(Is=CjK#TrV2DuI4>w58_FrsXcVyK|ta&Wq0QtZ`$cqk|Csx#dQx`E1c>mpw z9C-jhhlNJ7R?UqWB!_=JGbtzx390jULg7uMy|45Yy z?#mSwK2xSS|D%kI<01Q$dKqz160u4UG~?1`c{>h2pH<*+ss5MzOe(Y=n{r=5=HAuV zKJs=5IVKRIN-3s$LnB}R7u4Wu(tFC^;iTbl$F+B0MCoP}NUB`)>sC7Fu$7YfG2<7*HU>LCXhIa;=nkK6qB6GCSuV1Bb#2zMX(q z{?Lnlw}`zb!Ie0v>y~J8H6TH5#Wf`V8$*cu0X--bFq!I)^1e_2K!&W!sOo@9G+8yF zb-B$!;$t$wJ12vU!goq>YWS7${$=mUBSinI2Y3KV6Yyz!9}xg9x$|T)rM?WYqW5yl zj>}N~P2n}&=z9X<@43k>Dw>KA-B5>JW+pwypa62hefkCeB#yq7`ZqJfp8u#DDZqtJ zgK_^>!o&kRfjn<=*bL?QYYG~plngXJ9Sy4SC~*^UOpW*jQigAtNkq-JIW;P+RYdOA zl9bmI;bJlgiSZk!a`Bx9*3}R;8bb;DL&V5Rza9| z_PD0KPRjVx=WVL@{lOZ9ML_>Va^gbw!!Z;* zl58gEsj1lspR%NlI}=V|4>m)tI;FO4jVOI2q6i+XC05vLT;cv=jjahkOAH$5bpudTTNj#s|W6r{%H%6Ge z%^n3uYRT%X$?zU`^SJ8dHz^g>#X`$7Ba5od{|p!y4oKilthwOhq>aeln31rvTdzs- zavTd?#@S$p$=GS=v3V2y{BE`-_eHx#QC-tKfW&murR+w)xt!?U)0Z&{t__6zH~8^q z#m?ciPw}tidqlbkteEFR1% zcdy=A?|I6-1|qM3?>-OxG*IP71hO>*A{XV4hmV0$|Ix9Lgpj!m8izBDp&u1~ey_&_ zhBh(U9I_N#6K!T5YZg_Qz0&)EwAThvbR7Mw>p6%87!Okcc8DD$ay1lDV^IB2(}KqN zL7gR`+tdvC&!oOOfEGKnSt~G39oKjglvS&wO&gq43l){~+>@jw)zgdRNTaEQo*M&= zb`K~Vq6}=e$w^oBLyrgq^b=`A;R{_2y=U}0q-E+MOMTV2@@*kIB8>=7kbDpEBr^xQ z>8CB_!;jZ8jz~rA*Q{LWv0d!{AchUJd)Y=YV(&56g_zQLQawi|+J_FL$2shp)HPnI zli96C8w%p79etuz#4Z-U!%1kZIVzR;?fk0T-ErTY=%kjX_8(l8gQhvBqN5`p85Lq+ zCz3<0;(7b?aaYS=y?Lf9zw(H)6N#b?RjQLj%KNdW`s()GQ}DWnSAZYdl@)}ZAAZcc zLLFVJpz_&YN9O?A@pfhV^BL!zpana^@X$|7WvQnc6x(w0-<~Uzo;AF)!?qBs0r9^I zSEEQC{AYFF;ND#-}gS-wo!^3s6CnZexr z`uX&TW8S#BJaQ+g=J6m(0Nh8P(mq3)xra835J^y8pbz;T0ZycyN%}jQs!xU}LN|`| z?OJE-VY1PmaauhlWZf^PboJL%;%!!L?7ywow4di|-=Z^HwTAM7f(v>rLdZX}{nQyl zfC?Fc1C?;^vtW5>E1fI@E#9jg1$rG2{3H~ zilu?*fXXVAM8t%o=j*={9SrxQ6$u&grKRbb%Fn%DOsq>56q~CJjV|zjSydAzGI+M1 zKEC!s8QOONIHV6;+*={KpXJL3H`TJ@8VZ25R+T zkz_z8QA|EAu7C>?Nx0H!ITtLVEzch+A^mK;nk@8RtguYKe;p=Pu6Dm%uR3LIEcvXm zs$YlIbQs~G^i(mi8#uEqq3vNArD?Zd2!l|6?Yv3v@leo$p6H=-Lf0`#9U;a5&ErUe zs-+Aa7`f8yqKHAj**nATE_gDVH_E?Y8*AK%^>M#?E@=!&VX;bgA_$TdT0_tECM4mh z%Oi@)nGZ2ATcjPgaJtUKaJ*p zy%0dIq5OIo8ss<94Bb8Qr~iKamwxTZu|DE|?fZ_%KYJD%XM)-d{h7zL} zaLJN~-9}UYTFw94f}}#%8mj+>Fl>+M=#rE{8fxLzol zCR?$u|Lo1;y=mQzSy(Qh`u0U#-vDB*g77ttBQ{O1d%&&0ojWRF_e<Ve4FB!F2QWMGfCC{l7`NzS!Mk)erA{6;+@0%y!C@ z&1fOC@u1V#O^rNgyv6?gY~)j%#O4YAsrjoLJVKhQf2)pIXYO`_ik(afj?4PCpJ*1wGAd}PqvNO+iRIyd zI!YpPmU$CEPr6yuH&ujzbnI*M3d$!6X;Q7w93}dTV`9H|QY~kvZxpDB<+jz2TfF#c zlqW2a&+ahnX+eO>VkwAwb&InDoRy?t-eZO3KELCfR#ygEN*=hiJC%~l&Uc%$IR$~n z8NsH!sVPYs18o5~n={{yda-Ezw@L|BX>)J?5VkgL>%CeA&i=3!tc%xG_`<$5+$3rs zG6+$e6^f{T+Uh%KQDBZwt-D8{7xMj=JQ2Aq7a=;)LP(<5kN#IS{rTlLO%zNSc%#;* zUn;yiV*Hh{3~qn1+X~&-@swFW*_E z3LjT9v$|L5o?B~q4bO&Mbh|0P;T_Le$@(%-w+d58Cyep^8Z~`?pmq{7GmVah@Wr_L zg(*TEet36(WBABcUJ4GY_q>bo`xaEQ*>%g9NrXst{(Nb3?u=R61UWwQ7e$|w2hXK} zS7EFP%T}UlzgDQ2!t)id$tCkWpG+hx3w{q}Z(vS1#*){yFHf$Er*%wiPkHoOrFJ8!{r)l1vd$IZO^r<&MfCB9G07k9eZlOsHd{n?dcR?0pW1AN z{z{`SakyhwIa9AFjboX{BqlW2+0NfT&BG$FUQ18}Rf|E+9rze>e%N_qI&0h_Y3PQ^ z(wq%vDcx#Aj~^}_F*7b7rpbOTPG<`4NFHoDm8_x9C&0D8$4i^uiPH5Ne^qhPC%O3l zF!jyxeKcRV&9`BrHfhYpww=bd+1O4R8;xz-Nn@j7W81dvyXkv>pL_q`ojo&W=HNMJ zo{dk>{3FAu6Wdl06#iuLYz$D=;CL)()QU@b>|u;(;umx+tmrVu$%OAu?ti*AOQBrD zni80VU&$T%XC|wF`vJ2o4O`bnb2-UmnWJl*i;@2y^^elSYg5fy16%&6sa1;xMXc{( z>H3hAf z&6Bs@s_KtQg3k|KZNh3R)fZh|H6?kRB57-no|?0^F<)*f@t}UJ_>85+<~}yNrcM`G zdER3c()wzFz3|_%v;p@Bxt`S^GZK0Kjh%yZk&J^DzR~t zc2&5i$ntJpG%ej@+qwpwbcz|#V)Kq2>vx|astV!=YzO==I!$Z_aOGA@C|iN<|8XAG zJa6WWo;h&#oJqT3HQYTd&KK;3dNazj_5SWV@CbSGc$*m*KclMioQ5c+2NQ01|w0CK6^6%%*DK)L)@DC`QvqA`7yku)*H5qni)n%;)&-{vl5j6j(Fika41A@yjq(a^0+lrJ@I1HQaU{LJR!-c z;C0`zPdO1j?hfr$TY!OrL*nVB9u4PiqfxdL?k2^2t#7>gG?||H>S$PWq+quK!*3M6 zlJ0h9R@&OkeRWLjg>=s%!aO}-E_k@0t|b)18vdf9MNT<|6}@1ne-MG=l3gc{bdLXu zqb*@98D=`tym@)Zcio95W=f0e)U0KTE-my4&Vy}t;7ZK~B* z@D`M+7G=6d1^W#eX>H{05Ro$c_#>h&y( z#DA2!zQ70A0K_5OgX7>gh*NuUSYcdkDLZL$}#l|~X5 z#IcrC!o_>n5{qv_$_|K|oEP|+CO+atYck6uUXq9|RCNY(8iAL^7}E~JYA}{Gw33O! ziCBhS3(Va@V2-FoZpB749ontpw+zYxZcI(W6*7#Fn4=HQ|89+K5RF$pj1(zseHDqn zGFWpZ^2L49Q~7TGSJIgZO6eKZ(yQ0zJnvR>VeSZrDf;sf)V;YwJ(tZ+zAl!l$NBsiW@G=PycvDo!OKA1){;m%&Oz~UF$NV_;mTdO{FDzvpSRvhnZu9onDo=*Pud`P&C;T1|06WxiLy<~u zzH`mXd@iU2!_i46{Z#AcVoQXbkG5dRD&Nw`C9w6Yq}J0Cehv{xtK=g8f6QHX(kcXK zAvakV##~2;?x*tz*;^JX_J`DIj)#^s14dGIl>YN@g|573W#>4^L%_7N?p?a$WqEy+ z659giU6#6)FXE1{75_~A?l9qI!F|rOk6AKbNn~#wS_7zuOv^%=N4AOOBO@jjh?ox? z5v0r<>6HzhGF>@W6O*me0~&$?F>u1**KgfkUS@WS0~ZqQu`22JlRmCH90`Vs)InlcFBdG9D{Jl6GUAMH$wZ54R}g-Cmw{H$55)X{_U} z;UV`x%*-c=JnoU4_%lV%eHYk}X!=-uZ>T^2X7s&YCfhYzgwGqg_q{W>T|750O{i!+ z|Beh=({gFaJF&}S*V)f_4PE3h5iIuObVOzPy^nIZiCtEKSGS{Ta;jpDf_@Q1@&vKPus(g$!i zgRty$y|I*LKf9)bgHCmn-0qd0c$jzuLwcwA-p1WnFigoBx;c1*fS5V}fsIMg8P-kQ z$lozRL9kxMdK23*<)I~Z2W>M z$)^+XqSd2y8B!Wb$w3QDL;sU%UbBN!_h;dxYcur?*7Y0~#3r(AL%om!ZI_dF=-Mo$ z2=ej9!|(Vv%A~39`QwBPzE%XATu|EKjAN5?G|u7yM(tSh{De*2^0+0ax&I=*H&J(Q zLMMz#%c7}PAYcqRo$G0Bw-DPfjrOsj5!q~4-Hbc1L1mc=*~~9oUx{f9paT;rBE`7* zi{|z_((mMMw8ODl<|go?5Ei@_opIVFr#@as^@BR9*)U`%eMKdimTL7_SyB>p;&1UH z;7<9mO;uVzVQB*9Zsi+q;sREBReg}}=E?1lnHSn$BMU7LSsCrRoz_;uUd|F~{jn^+ zBt+#k^CT=q&n9ISkF_L}Ir~`1zx)Tc@C&X?zOUEjp0M$Qev~V|U^MiOWl9uwvV5T}%^jguYAhjR^f?W<+!f?9|`b+=q?X`H5Q3qw#cQd_@eXK3@ZM7V!M0(0~@L>#d zx=L2=z*E1~{4A6v!=GWlORm36B8!Fu^ZB|>``a#rA{6RMwk z)$6l=14f=*ID4|SznRe$A7S1fK-P`|D;t5%iuOlK)?R`CZw`_&3_OdRRx#I=f7|_P z7qe0Xv+b%%F>OI*cBj~Z1-|2buD==1nnm&awPsU$X0Oj=CqSd|=ODf3wq0Y#VD%mR z@P_q~T+lhiPLrxk+d7#$Pc#m$i=JXoStT!!368_|e4$ z?9J?NSqPr5H3mggitc+wx)ZlvLruw7;r1t-%&mg;(6vFhyh>J-y@(ob!TcvzJ&ZqD zx5&yZAwsctuD^2+7X;U3Kt8F?xFD4`2TSK>A-)xDmSw6fk|^yv#8kGixwZU0aSrlp zp=FUizfV3mj@sWt)^<)jbLp4;enZx4ROiaB?CoI^K`2q`y7$Y@E!Va(Hp-PqZzoHn zNfJ|2o1Mez6j>Q2z2nukO23jVa%yTy5t%go;DUI6z^AzXa;+j)(Aq zx^3r$oPq(WG#lA^bT%kQO$8U(jWpc9brk70F&d31*-NnP2?KSstWq)8u?59ejQyf{ zuBDC$P!X0k?T(s97&&|O&&cY*4;lg($Hk>@3>dA)pOM*8lZF=j`i`8!bOS_Zu*z-~ z5MTsT_<(wpsbxr6*GH511wO^>Jkf@~Z;u`GY!g7CI7&AbrdI`ySn_p93k@&^%ZJgSA8wF%8pR5S*7Tvni9}(PSeMQGP|nCc?9izU?df z!}y8IYC26P6;0i;1qm(b^uDo3*+$qk)zi|Qa8*xNN$EXU(Mz2^vl3=j1bY77=zp0pahiD+jA zjM6PUKDEQYvE&27cx(x2pulf?vPYuQ| z6!mKk&!@o}6tLK}OWhyE8C*+OuZQOu0PVAb;D33NxfJwYLhw?b>wi?*X8L=uv?Ver zsYl}Vx3jyjv2WqS05^{j!5s+7*Y6R66!0B0wy3`XT!U^Mt0#YRYDOImK*NbyAq5`h zehAptR|>puPPl3$jJsa;w%!||ttMj%9erZNig(&d+AuR?7qIG)^uy8VQaN=xgJOtw zKM_V(9n-g7+9CieQy=>vh93q_l_IyW}?>Zy5=lBXF*Y!2g@wg`LX zR0Ai0cbq}(q%_Ml)?IHygcs_j@w?2QiXQetJdXC6awM}c{ktK(4wXW|=6L?Q8Ml=2 zG8<8cY`LBXs>$xnO2be#yqS3M6|~2VMf&Q=lJ4G;>}2P7uXEo|9h6E6U^7_l31`ed zf9PzeM|00=#Sm7wZEcKmmsFv@8>iWB`5*xSPG0?AQ$AD2I%f`)iwhV2DTxo5$Dc}Z ztB=ke_{}_Rx-snG-8{-Gi>|%rtGo#PnnsG!XBCYqAkk|WP^44bE#;o~Qt45xBT_&f zxwK#G-91E7GVY7_TT|AK#)5T|j)&?f%Ig?}Y}0-95r}6h8Z%_|PCk|PhO}K`v{Z7T zJh7B@zxLAdp9rFh;+Mb_K?JNYmXA8S!Iso|w%o9bpgJC}mm+c8v}gy`)W+puYL8jk zLq}eO0uUHB_c%3rJR)C*W`qtqK$?i{!ZPTYm)~N9Nr_UYv z?+{*ii5x(|38T`3m<+P`#o4P32EIwOcb|k92t0(0#hf`ky$qPcVv44?FR6fUZT2{K z@}Q1??f?A8_FFUycIDhx+JX2d4|#PuNWU6^bj)g$PImW}9T)wMX6cM#wWG-u#;|Pw zJ!9}Z@?Q$!K4TBrKP;>WJj7L8>8r2Y&%09>A!!?=5uIE(!uUy}i+c7lDM3jX{J|C1qgX_9T2ne7m&_ zMNGAP@!ch{=2DRWlUR%L!F>deT+>zCRUNNo&0hT3S!;@8)sQD1ZAW@JB&C<96~}e_ zk3W-TPKe#lit(~q`0OunYnUtV`OgmHUEWM~#q);0Z3aj7Ri6K-*42k=k?xOsWeLHJ z5q)(nUej3jB_4QeF=oLxYgxB0dZjeCb~$X-c@&r)RJgZx9OkK|+4E?( zkVd^hvxJbJW)eGw0Ie$F@w5%i1ae-B7%@r$Gd7zI7ZPvLnO`4b|4PNnoD8txXRiK~ z#0zZ*Z+#8||1PP!;`3lfvoj)0{RrXd6xy=v$_;1B$Q3fTV9QKRfeEiIps?g(MZAi^!G}~(^c9!|9vbit>beP7PG#6dMn&gpH*7J`J!$8j5 zF{Jmmx}Af1{S5jaH5Z3=Cj+?t0C?Urj=(~V58d5zT1+z+>1dX{1~947@CqPAOv+>8yz$mSU}(!?tDWkJ^ZnJM6th_{^x|YxoAgfb#|u|kEg`PTx4u_b33$-2UKKUFYLD4 zbEXts*d0ox+0XR&*1X*N^ z)YeXPb7U@O7*f;5c2B4TR)d;~V^EF;vpwEYFWxUA){CRcCI^)|vEHlNE+K2#8URGL$a)nH z(*o@^R*nc*C9u!x%x$v26!ueC@m#0(y8oY+*uZl$mrlojih=%~8!EDAoOotgI?r0_O(7&gCoR1!` z9v4hRYB8laMM`l}tzK9C`c?VZUZGtX9VGqz?l=J_da;{nosS#Yypj3FD|Ct^EP9Ko=CS3BB5&tyc4y@xBq7as_k`gM+>t1iSBL*47_ zs_14*zpL<%)+_W{EpLpibq;Z@M6RG;<~BSMVmc-_iY|@Lsf1G#F zr<}%37A31j$MKVlvC}E$!Cl8QY*|K-B@A zIjXZZzMK9y&zf7AORUk!qq4ZquN#3!)q?jrHM*kRw7lZFk!?csFDkWU>g^Z{QdRoG zA_^4kMM6Wy^~b4$H@&j-(}KT1GK+oO)#!w&NCKBbI=89_83Lo^#V|E(a$N^PU~g zAC{)Im9OA^6?HO_XSu2Hk{p{hT&lL8(_2}!Hi2mSi_e;f>;D?UHRRd$3Z}VwfFbZD zZkdw$DDAKL)K^!W*S@3sPlTq0L=RkrmXMV?e**lszn*R|^ZLa`TVss%#SCAi_a@RQ zRuFxQruwo=yAkE~Jme&$@3h>3j`t*Qn!wCkmqeH)Uc@e)p&pa}yM2yBG~*+f#qrRR zE?Lgy=kW-}Yj}AHqw<|U-(Ne4Zmtkh(FG1k!498!br_dz3;}?0^{gryQQL6nPVswh zz){X^C!>022z@Qg;=tsOw4r#Pnx4*w3e_6ZL#@fZ<*mHSiu!GqB`Zrh# ztZC`rdTg>!=5*v*h)JuHZectbJRrHD({p4<6yguCSJ_dNN-i{> ziMY33q8=O0{_rw%URn+0jxydFX;(YRsM^?URRKXljh8xiYVo z(5V(2IGx;QAcRyVp&c{ec4-PF8_e++hcWK3dU&h=Cy?sg=j&SxAEyo*d#@gB%$rp#U43v9ztX7cGb|# z?N9yBePK7kHp5Y-uLIg6fz|0#u~DK4GSNGI6cm(BUPeUPG3I7~*{~-LfqLd?Q6VY> zh(zVY_tSVYc0P{R6m;xCbx?M2^#8=dD-5`>onm-aii1;EqqJf`ka(7oP=__}Xk1}0 zHQ~2uns_v)eEPXp)}z2%FWen>>siAa0(#oo4Hh0QQ57)6QBc$S!--Rm{bqYL6`0oj zKuj;tVdaK=Ik`zdXZY2KPUDn_XJ`vOxz(nN*l;or~K&^(aXuR?r5d4K=I>^3K9N{7DKnOU3=CPWIXZYs#)5~ z8EX@C#=EsOHpjN_&mzIQT5x7Phk$hkXzfmy$+nN(*Uki<%J?j za&=m$ff`z`_O^S~E|2D{OHMRVu7KawAbIv>dK+5@ZQ%8NC-T`gD%q^M^^gtQljB|X znK=F=IC_Dd&AzkiC$)=A(L1GRcWKnK0ov~eeBB|K)86`}>}n+DZm+F7JMv@MuhJe; zw{U+`ChXu7!e zm);MI11XPo!Nk7AgxH(W&7nbZ;vI^B-`(+Yb%j3+Ycv$_Z3Eo)eRn9+@cpfP%_xqN z2mRCissDm@a#{zvL&?;YibdTm?MdQ>eq!hJ4~u6eED_P5(prr-1w6e}RFUpcRVfBi zsq+e9@D6Wdy>eN1EmW$lMrLMishRb+jQFh;D48XnSK$4v?_mzrXo`&K4O;vAGV!=C z`!hhW^LxW`6Do-+yqGw2xJyU7)g^S?l&;w$W5YJ-g4Z=-Uo##|QibNyb_%iWt(`Ud z-=l;MbLGyiI(|v$X)*~DNoe!JnLzKx&PhyHUTX`VCwG58U*t0*!y*tSLL3F*Gk63l z&bIHW-a=@HyrTQ`Uuhcp2c@Q8P~irSWktVyn|IcqLrT#5Bj_|L78V|DZZ!@%YkiTR z)4gLuvzG_BOkvqgv2VcBYJEMq*>hV#n2)AP9x&OOemje(m3K)cZzDasa5=C27^zd! z+VOnO(C2-Z2*zy20V3)j?qaSbB>7itfBm|S z$BfjRQHul}FbTD7!X##MunSe)*I8FD!^An;4!kQ166b?a?&A*si=Q{$K~e^$MQtnvHUotvBzKGdC-ndF@>K)PNWN4dji3E(IMpWFwbCMT!jM_{ib!(jvsr4;Hxa0lUTQ4*(YZQ6G(rZ+)* zfIXNXs6WYrZ29JIY28VcK*=}+kiLGm>%<{5mu%wnAkDdT_>F&KYTF{0+t_` zbr!gur0zL*ZQ;!oY$&{rjc>hHZBX}9$J9p8xmxmslrs9s)%0r|KxJD5Fu9gc)bxE$ ziVha&?bVE=5)PmdK~q0gWpUujXe2-`rpF9)D*d0JX5f=&LW5Bkf48}1%Q6#%%D1X< zP`0Qm*0@`atzOE&)!}3f7Ex@~KrU!{qv+ zBJ`UQrTl-R7JwJIYnxj8D4|l~W`k0`cKF^iu&{eD4%9 zCqss^WpmVHCZfnKcNTaJqiNRR#BQWCXPWIdqgt%|ZwHe3*dU>^f#D!~qq$T*NwDoe zQcARD@eK9&=q{Oa63b4(6cN`Jyn(KtyT9vRidyUu%;a_vPzips^Q&K|N>g*dLrKj4 z+-qw8--(a~(G%illOWlBhj4IdCCeHPds~Q$z9%x0t<@)%g@4~a@a=-!0r&WH#2+i^ z#s0K9qla(g82st3-&Gyxfluu3PI7?zpK1W=EVcEi$=~YVrHi~tASYRT9TDll_adMN zRVKu>Q^{DTJVz7i#sa8xhs>sg_$>zwhaZjou~ zp?K%(WPYE(-SMHTudB8=ehN5Ac-JQqDfE7S7(|HaO5|SY(2YgLC~?BKfCkxh!(f&) z8AFHVQS7iehs%wf(#ii$1HJ?3Ai2m;s#ic6_B6ltm((ui0C-a#>{)ml2!nZPBpdYT zif`F^bjaAi3cpfu0M0}&C-nxg{}7bk`KKIRGn$_AvU z0j2nh14m01#>|aaHn1}+F)EV>tYOg`PjYqH32&s=kAMVlThcQ9V)e1;)_Q;R9SX}e zgx&Y?3~v)*-56~rjHRJY)44EH1<*>6;sE1Hq^e+ZY1;r!e7rtUUg8faB#lMxGyg@@ z3rXw5={fbT)n3Zf?+-jPR4;mq)5i?$YsgPR;G}wc^r_Eo$9KZ91pU%j-S>+BBn$8+ zA^8L=jFJN`j?7l?e8ZW4_xLIZ^boI_uGUnquT!0bJ~Zw39LyrRH+<^WhOG!p4-O^S z^pQ9XIZ!wM!G?j7{off=3hs>Gl1mj|7c{be)A~C8D|{umEXUm;g9%Bd2}Sm&PU~32 zHurtjQ%%Ru)zHLBjUe0Ch$By6H-SjA62?1|O<{?3LpnDI!ip@rcWRATZsxpMJ2FX( zZAfNC9+;3@nx3o=w$1}bLK4Xf3nR_Gootrbf6|eewhX{B#b`5^hipd>%VJRm{UABN z2>HR5*3lg|&D~@!H8Uj+sGuj|!x^hUu`^G%98=`FsB+G)ySj<{RUJILOheqDAoc{b z=b7v)OZu~!?m>#*o<~RTCQcG?1=C+i%JTpj&v#6Y_o&Wz7woDlXVcUi>44iBE|qlc zMCM<2Cxz|=HF%e$gqUJ)j_za726-rQaGa4d{y_Pe08sdCn z!#B-M#R$NV&=Ylv&EPs=X<7T!5VF9LBG|8-_Rm>&UKxEA>Hpq3{B$((O5z&olh{PE z|NJXl1pSy(Uqon&aBr;xbs8z!UAw~6c1_e_LkHMJ0)WC@T^swBAf=OX@ewdyw zV5P)hmWSys7J||i%e+VP^$^r%cfvt|X!r}tP*0JsKJhJYi0lDwT)Ea!lej>uCST3s z%8v#TeXmuLe{(S(Se~X;Rp?H9@8kxEFA`T)iC%6+bmEMBr(}w6UU|#X?)B*dbwY)xU$=BJ!+K_;0!ls*e}({&-?}!| zt=$RBL~MQ!dtiwUS!V|lvVKtsi){>k?G66t_Ng7}we~*`jYTiHCJJwGW434T#ZN1J z{pV8&iRk!7M03|Xk`KmPfvb4A-J#X>kp-| zv%1QKU<6}7kV?4OS6{Cw9=pHE1sKx>u@$E_F{06q3>tu&KnLuOe{m;kk0Dt8MjPu$ zUh~=8Oo3Py2F$U)96SnF8MJ{fdrhgm8E%*Tce91c3CC_$qiMtL_SSUpsB13HCTuDZ zjw07Drc6~zrR0l8kcicVjgH3=5AkON5U@*n*4*B7$RQNBX_jAUK4FW<(l()39X#=N zd@7oj5-SQF;^g(~xOo-id8CG^|Ef8&L%eJ|=$`qhh{;0hjA_|rkhZbl=%}%F>fR1` zqe>0g=0}jE=JE#2!+wqC{44#k=cSA%ima)__pPCe$|;pBm6~iG!+mMu&sHE6;~ZQ8 z4((dC-*oby$fz@=SMPWLqJRsY`?O_50dB+#rkr8PXn}{Ph+TAcobM>5C-L4$=K`Il z#vP~1AaS57O0-MVh8MVyexGdB7!yWj5%B56=@S-%372+*Ivm;xau-Rfplr%a2n-JA zdV#degiwR8g&(n$g9MDfG*E!kP^Eg^lju!xuHE-ui=fi{dbiuILpYHTLOU1MS-mo) z+d!19a`jEMb%hztVSOLf%IjCG`YTTA-JAbT3vV{(5hgaxd~l${823Ztk9|=zZStIH zM=J-IS%b0jCF3z8FXjm9`CA&)%euticn;Q9bncW_6fI00Gi>vyNytK3D>+-9-Pa3H z!s%en(wS>F_1)#@SdL*jk8m5K#+si_eY+YByTUuV=X;BY;ka$;{yIAscD3MW0D-dm z7?>8y!6p539aprr@Q3aktBFr~A*!5T(ktzEaH_Pj*65Hu=XOkyB1E~*UdEf*oCQYf zg*JEJP?cQKS0tNcybcFt@lI(JGk~Ml4}?u6yb#2ni9u_TvyyNo_1A&d@mtVG>U{3qAz#SBJ=g+F{^0i#MaT4@%E^ zm#N2|tcKf^U+Jf%Edm?R-XM-{N3W-!oCFGtHCWo0?%7yQ-X?xWU^+$Vi!|V<8<72) zyG@fc-^ABHbFsuLl7)#E*D1#dUPvXyj@_(o4%xZsf} z!kL$nU>E9w6c^zVm%Xu4EbtdT5L4pov&pPmT=bKq zLTH8~V)Jb2Ie7ZF67!s|@ff5k6NEOdI)26dp;MUvZ$p4$;NO3vIb? z92c**XQJd;R=6eJ&0_m4NWZASJeNQtAgiQ0GT{#9-OqpgR z&nkkH0)V7gX(x5kmcj=gS`x>{nq6TfcEjeb=HX+lZUN-Cf-aJ{g>bYOH6V{9tQce2 zjD`B0+Xd+=3j@x#?sbGmfrf~|XP;X_rKDH?EB}00`7VpOW}8L)z}uC)Ty~o+0EW+!>rKKEJ12oyFuTjirCuA67eJmLiZJ?7#W zrFdDdEQbzqkYDC^-d7@>T_OWs1$f^g!J$1KSS@SPy$xmq*K*^p{#2GZK-T7tMAlT( zyLa!P_#P08dYh*OH|>bgpuT=X`{b1&zUBQ{=+VmRkE2*{qi+>I;@M#DP)%1<`S?$f zW)+wuQhbT@ZIpO2hgbDM_0+xEflhWy#ztl17x@{YDw$N*!o)@68lr?p>~h2rlK}@b zv3dum)WLY!pB^r>6g-D7at-m`rXUrA969w*8QYnI)QWr$w9i7C*5o(oY1m;Y8NtWr zcUee)CH{p5WY&E{tl3+#pf4o&NAcC6=8)mty^6v*n}mV~d^P9SIzbq&nbxDnzG2AC z0;nGlBUR&bzzHPQ8E87}P#GthhYDWos~q+ejq0_=3^nZ<;l$B27j0YnKOIz{o!mhd zpKxpKqMr6YwynAFq%O>R?z1w^)J7nYedC*3`bLf}y<<{T|s|9`(mc_$r%#_;^Ka+|SM{FZA?@ zOt3Ex>Nn-FM-dQU#}g9=|Y`qeW`;s?LgX6$A~B`u_7GH1T9qL>%oCy zSqdJt)Sde*{JSgDwA)d19cvw)3*JaLiQem3)q}-Nit%I3u4ureYR(2P;RNf8GnWIP zl!7lw7f6?+k)2x*55qY}J$qj`yCw}n|>6(us4fq}VzYA+_;YV*QjslUO8N*ON3(=qT>a@tQAg;|d9*LLa0k5nlAcW@V$#IioqV{Ixufy-d>A>rTqjjWXgHkE zg0MBOF{3}4+PQ8*dfTmqY-D%wxejeR9=n!S{Uk*kNTh`fz&V2s@#Zd4{+CdjN#GxS zv)CNa8l`GhQMGjTK_rfLUaaV+YOaJ6>!p^w_j=l`dk1k>yLQjz|L#tyBvY;D>>Q+=k|zhmlI(u+Ozu7TX!EI zHuAx8^#G;LL@=tqAs^j(gfcJaxTu6K)|18IRW?E-I=Tj8xE*-h6yLODdrcZ;^ zXW8ZAWiuVjc1N#a*n$$qysny8x#hv{YnBXc^@yl5&uS3s+=V~Y4l03^;=G%?lqP7U zgUZ2k`0`9SkOE`YLM zofA19oiFy<6(@PfmCBR;n7i4{ZiX_yb$NNVAeEwrKJ=zM!NC^yL_3+-Xcs?Dy#Cf* zT5@*MW>vT24OHcA3(akWZn0P1x=U9$pQL>`ziFOW4RQoJuo{i1QtT(s1X>_!2=cuz zJ>`0ydP5`h-tmL8FLe}I%-44Xh2`KB1fAwlsQ)l1c8~yx-UnLDu_`BXdumW+FH};1 zG?Z24A#B3bBM~fH^D}W99o0&5ZygSl%L&Zc?*H(A@4{yFh9%wTsAu=J|rgtFzEHgoJ{TL;J5tzO|V-R z-W5~i?iE(UjGoW+4{7CIu}A&k6W1>rBxRQr1W_urh*)=TvT`T_{3~`go&k<}wz3J< z(TS+U>r%S=3T%i@tvw!EZFfN~Zcsahl9)b--Bc>Qa9n^aU}S-IEtI6v2C^T8ND`sSmJZgrwWg3wK+Fe`r5?CmmJEj za!c!ik7mTNVaPb_`wWI(SKRCVVe;Sc0n-nCURw-ElMf?_^&Wo1rdW!PW{PezaY_J z0hW$LGbWfLV+_nn(+({L{f@oQsFy;!6B2TM*}+kJsv}xVxX;6gPHp+i6g+^P0+o?Y zL@=fAFWwU?z=XhM013V(Q}6i=O~oORqpeIm2sT6|1fZEh6k!dURANUiy5q;uyME%m zD|Fm^b8*@ad>;ZI4!n$XgGV|iKE!fw&&W}siNo%*;W?Y8{FyUVZdjFjyv`U2bh2Tp z4a0>DO=9Y(`J)iSfa-0dCvcj8uSx)ztL}g)C!kUaBK(>icN+{j#b2xLQOQa_cJL8H zg#MWQ(W(WNWivcB@b%`YpOY;y(Jn(w?TC~_fD6H%A8*gpfs`4eLWHy*NSR3 zN^=73gdkq8t@WxuekAMiUj2?8?_URa%la%<{~irIpv>g^E@sY-WWOHDKD4&FP$=6} zpU8O)mUYw{-CZ@eCdUTMjr;Aym}00J*Qd`a{8NL#yBdm2u>dvr65yK;T!8KI7=L?| z?TqAyDb9-~OOVubK-?-L3awDp8H(4(ARnXonKg{im)34KAt-2&?8l&WAOgHx^34Y< z+2nia4x^ED_^0(Lpd_%PrLyc?+#_9h$)0Yv=z<43$MQ5%2zvC_hibN~KA*A~P+|vm z0+~<9d;R1K4)=TfEC?;>@?j|1@B*8p^&Sh)(@8i7nId!)U|1Z< zlf_DBaKF^-jM={t2u;RPa>2Zlx|Q!=R-@`LAdys{5Bslf>GmbG_~$Q<3ManwfQz z2XFU=l0E9S;E1DYVxu#@)?J&2_#>8N!&JyWnd#mOnUeTQVSpC@eEEwR;OGANk+d9w z)K2)da6M0PIRw$ksk!EbGF;6;?J6R~JdtP6c47WiiUi^Ui&wolz*GJ`vPY&{HnVhQ zhDLx<c*OGP@KlNx!F>Z$i7|C+x+VP>@;vF%@X3?Hz>XtI8d_# z9C2mqe4O6KkD}?y&CA(p+aU#IADOYnOEQd2Ioll0TA#u4H zAyx{*y23DDXZ8;$0lgZ7>v=|pG2(Jy1A=V+`e1zGmCNADQM13=^~12=)x^EA)ns2# z6Yt2@a6uvHZ19nNKyaZ>S{%dEdvk z22UH#dAiLeaM&DvHMjANeX63w?ymWmA=o`1{@z?7j*t7VB*#_(rN_ZSI{TH=mVS!| zis4`>At^$L6F_xCHNT=LOTfiRr+_8o}Ag%(9!VU0kg54ECS-4;N96KYaDcYSOV4YsRX zOb0$`EA0Kr1U3|9E0>!}x~E+4vi|D8mGeF}dLHxSwT}1SfMo9j2E3PG#PJKjd2syN z!q9L=!LbYN@I!Oh(VKVb8(90hCyo_>I=DFgEA+YgIfIB2sTraDCLPN7m!sL-xx9&c z9^-WMh&Ymlbga9z%PH+ z0f&_8+ObE{(b|W(yv|{jMKAZP!}}*;AoIojMQ zp)~HIoao{F;~>4l$tMy9{SFdcrcssm-h(Xv&Y9#Q{kzX8DB2l!ADq6iGY0&egk9f~X?)b^9I55Wc0FVd!SA`1wb^jS31+kXGe zMxvI}*yyx=M4paf@x{#R-x}E<18cM_sPvA+(i27*mZ!+3VOxb9D`q^uo5xU8}h*dN(T2cz{vq?w2Qie^dZ;uw`pUMTM+g z=-#^yBYO8(mioS_1wbS67oY3Kr0F@CVPj4PoWB{pnD_susn38EaG#a@-=!ew?BR2w z)DDg}?DDuheiIT?-h3@Pdmn)o6a$}{@b`Sb+UORuRS1|t@Z#3+Od9a)Fps1nR{tfYKrl*Cb`V%h$BEL>GXCa*PP(gd%- znIn-a0SDO-4c()jvyX`eueM!f22lzB62jcb3cOQfNV`&e*Rr&dxBvkbz)Obftr@eK7$ zyTK~WnNY_e)oQyG4b2uxUB#{BfTz7SgRYt=QIeLv22ecbI(tGZ&>W}ttQnf;VA==$ z@G$DLm`U`1oA~+dy7`~gngYAPPe|v)DU}OykvEYem~0xdO-J~Qh%`Wt_|qE>Kz`1y zTCT?ls8;UjUs&nVPa1mCo<~e?J0fxMoVgE$&mh&NPqu!vY>8Tl7r5b_R9yyAWLlti zwvdu+GS^AK=9v>MSxr*<>ATYW-yGzB87%U^fe%Cxl@18=Z%8C1yl~8-d`)50Sz1+m zfmbek4x3^p?2I=HZ7gVbNBe4n$*uM~T-k#QE@JFwFfj*->*$_uZ$hSrxcZp!3QP*G z4>cTcIzGcipyPDCw2VN&M$KymQO3#1#Qt)3#(E#038w+s_H=!qbY*cT(Q>gq5A_K% zAY{wljh*K0D0|ANV2+i zW?NQV+>7juboIFQ50=HtCn<_G9_TJor1f-I)!(EAPj>}4AAxoG-t!}7uiEvG#dqP$ z#a5kqlCK=bIPe%pR(K4nLe}Ez65>&ecz)s7`b&?VSs83U+^H!>FIUaYhUhsk%;Gk! z&@(!$yXS9S(r5HUvZtjA5|%#(TsrDVpDfx`>Nf>COG~fXTziUZND%-2OIcPA`jtsq5eb8t{Ox7n>VxlFkv*p?s=IkuqWLE~`NKgAX4zLTc5sGxR$Dc2 zORL{8^!VoJ65jaDKlO@wahhZ0uEA`M^zQ7|`$S-W`i!>W!8d_TY`ufD_{AIUvF#SS zxCjgH$%^*kBYd+Fiw`IGr6@!T~7Cl4CI@mB2q)^F;|jEkor#&ZZYxw+R(h19v1Qq(kUvBeoAEp)SALBj zqh#k9D`{$rvkFpexz52mTlA@}yzkI39ngSO=ztlT?aG()R;9@_xYF|RtA}M@?6^@19ho|?^i$bhgOzwJaZ=*X3E}#MGWCqVgSxb5;>L56 zjrC3j5yW*f?5n&uZmiM2E>*|`>xmIZ&%yms8`yGXOT|YYGK0Rw8Tqq%AjcC^TLavO zHv<#^gyp!X00S&{JXh>mW(M_!O_vE7Zn3SyAJ?~;5LOPAaA$@Kw z-*LR=RP^C3Yiu*L<;y9$!6zzN4M}@skvMVkOz!AyjfEYd2*1*C>uD?UxDAQuEP%dD zKT4F3k7fJf6h08To2j@WZ~YN%3X0r3)Zu+ooB!Uf{hk~x?hnU!@dCn0l@_RLKaBOhEI+&3ncL)69&abqus>2# z&EM2*+$=TU*G7I;B6t6-DPd}oS9$E-R;7=yqHJuA8A15L6@T!g=qUP>LFPV-b8Yby znqrtnq6jEwhSa@6tHY1xKU(8Tnc9nx3ymf0EPOT~el^{Vi{0;-C`}wW$76jUjeAUE zExvA#jFy?Rnrm&jWK^j&W?!+itqemqcSH)YUYKkzwX%<6UU=Qg=_CNQ^m%>4+(dOx zQ2$ZpzQ{nJnu2WjBzyAJ_cDrPa{p^!>^)bJ7D6w z9x5=vLRQmgNs{G)(jAK@YuMmjx+J3^f+ZJVPq# zSEugW%!|y~U4$1lu(z1>krvK+QOgyl90au+zst&9^{{g%z+p7EWSOLR-$Dl>)v&5& zPoHs3Z(qjdPZsPVMIhpE(#x9Q+RfZoSV>-&Je-0=z_lu7Y0s-CeUBF`TRr!r&YKd- z#-1C67@<#7>MIAbytW^2Hh-;^XAYf4;3hHLh>Y5kuVyiw!Y~w<))QatXiN2Ey3JWy z@1(tEPu>%su-kigR$JIe*!^%Lk3W9+%&`weOM7^_d8;i*dBA#mLh(q0t(C6ZKzh3;ZKX+}X%ZksM^d~Ei|n{)R_Qdz zQVzW!m|$*w*6>SdytRMS1%Bs!LNOH_PQyP}((;ym&`STq@6f(lO$g@942ooTLlS02 z=?*Z(ipvs&z`^E=3~ZA#>m4b*Z(2ClPO>qRUU(nPDIz}Qf^}T>Nlft#+z>n~>zLHIffNDyyIOCt_N~SR-D#5%i$3Ns{+H{Iu*Z?? z>EIcu*l|X+H8{&28H1GZI~%vk8}2&kPUmA>y6#}m^-NP#UsJ_r`tUyPb)YuS`LUqp zhg9z0Nhy?-fI^s2`a1a|xWt-$r?D5yS)NTzh0AnslHt~M84-6Uz2|4BhR)nCA4EKC$gOZbiB`eo1X zX~gw;qZ&R`t>lL#pR`wy0)Hy8IAw%U^d7^{*OaBd2Y>OcwQI={Dtt1BH;7fy1It&K*xLe);qpx30ofIP^+# zGerCG-8gmDv@6}k(lpO6i;C+{Y@3+PR!bbY27}O+%Im4z8tFsKXO(}JK2pqn*A{xP z{E&eB7iWy}X=NbQ0K`vKeI2uuVLa7|OW#(Kk@joGY-<-*`_J>K1385aD&$j3x2f>Q z{89s@Oamityin7iN1ra!Ftb78e3VxMR$7~k*k}qRCCiXvU%2tjpJZhpYU+&icp|EU z{wsgZ5Fnlj)&5LP3s(Z7=fC3k?jV;X>cs9oT3NEgOR-UPdn&eLeFWDNW9}{9LIh&P zKdtr-^;i?7*UlU0Sc=Z8;uDKUN%P{w(p@w&H!e=q z0cWFtpC!5huGL0WzQ$Yc^JHL1)-Q8BWor$<< zW`ueE`Pj`jY}xv0XMY^Xxz42VDc?)@OLn-Z(KQ2X>t;d+^~ub`2OdpQ2SOO>u~k#w z;O5fXD$*C1NEKPf1y{*NJ5PspM-!_y*O(yBS(jK9*_PnO?f?}rAqWfA{0Ld;cvQe0 z)*vP9EAfw+-e!ofJ=h;6;|j@PGD(`BMz!y(m#egy{;uQ<>23&iEh@Pt2RpxTUKQu3 znK@CeWbr;ycJ@1x5+WyU{yZ$UGjsAbuIKqMDc?CJd|*pv39hxKo+0@`>ZEq6!=md* z&#;aAh!4&&FrVin{luR(Z8qw$3)` zl4e;0!H_!GZpcEbIDzy>)Di3`@=v}#e~c!fcs1MUR{41e2MzMg;GH26Pty6RN41AwKbZ zmtdj|q95Kt~DEm>i6y`0Jo9}3VEi5G}BjdF1yA#A#g)vmNSl`L7l zYy^hPH%<@>hwc-u&0W`+^MBL(QjI!5LHV(~9Xa&L-nHG3WR2cmbi-!tU?{2b7Oo$& z=sCQ(@#-N1yS?2>!Bao_xymlAdA?Mvp=?iqP*;K4XxLoy^b^Tv@moo%K-by|Ds@0) z`-z>oWj;n%IUF{Em_veUcGa8+>JIPO2&;M2%pIY3lzX(gt=O@(sI7I(Cjb*PcgCsK zH_(XPKgK7V5-Jq%4-|q41Xl8WS8g0o4-fnXV$|a27J!j(@KLRxe45)*)+sT*CdJ5ZmX7uI)3kpzZRv@mmYy_RHGe z=CEslmA-BF+=R`m75~0x8hwI8pIx*j)B#BKs!F4M?dDd$b~|hdiS48_heEW*RdGk5 z^9(s<@r;XaC$HWTJ(Bj!yRXD$SCdj(a5Mll+=5DVi`!)OdO>sT*#i{kb~@R5@D32X zMptTlFE{%2>4SovMf=hjB)67zh@;|I-PQ9&dA{{z=6to~O&-{T^lSIz5lz;ei}7cN z({lJz0F~|QgycQ9p}A^sI-8O_3QsuSX9Wzb49NSfREVgiru=4Ph~8A3YxNpT=6@h71aS^>GJ+ z$c8FQmY>;5o#Y1$s(#grIjqH36&Ubkf1xys6|Q6HZ)J* zIC`Ds@f}lsowCYy7(xv>{C)Jg0U)Ct!*B?)g;(uFhNcufN+E|xe4D};7M^p|dCQLs zB<|6ME;%rB0B2r#j)D2T_W3eN-_uiy%BG6foE|qk&Odz|z9^_nCmkQll<}Xwd;1B~zW$N+%!Xg$Te6<>pvD1)j<8QwOO`wL;Ww#c{%@*(=qb-+YLNuz@2d2;W$fx9kXQPBe6g8s3M>4@=e#9S3<&?}Rq+s$gXvMNKElcY z(~-my`gs)a0tHZ11y=L>8`uz2cLOL>fH6C)RP#f%)HlK{gpSZslF6LN>CHTUn~@8O zAPqDpQUekMqA|IKe5HWum6D<~F~wFI1C*(AsWSmZY9$Iqixg~Y+W)FN1Bn1ZK^T!{ z93TV)UYm|yRx=qcqEevJcQI>%z7!ey-oVL`>0y!^5GsN>xr=`^=2{UfXxm4%T5$C` z9(jG044kE&+!LJBhAlUf2coYtKncfY{z*xsNjk3l-=d(tElO|+z^!@%6QLsd?1w{} z{SkgRUZSCWP>>3b51Od)v4~g<0-z&&en=wc*!T`GmCJj-tG3T1N`oEh&~rW?Pw#`k zDgaz|B!m)$o`U+HL?@8?4No1?c6|eJAVI(SvL45S63uHj1-KKL;rJ9fzTbDVS&I9Y zO<~ZnngdnfFAxI?Lj&c?*=}m!mff%)?@BP4aFsJ|&{+&Z~< zyA2+0yh2Iw_EmtD`2JNKVBT<^E-%Fvz=UoP)7fD|<*HTuR`Pnh4;Itl(C6!u#`Htm z;9~rj{VEdE5uyaJ-ZbkmduYMZioc6T_!7}q)-U+8q&M^JW@SJQ1I7MnD)Q1dDx>6o zfacO`Dbeowd*zTsxZWUp4qJMb|G{^~B#+s29=dde)h=H~~DjK`LIgH}>KDMt5m7gEiCaRzc7+N3fSq z%UNn;b)v_NuEtK~JOnD5ucaGB;a?;x)X%drmBGDLHAZ=7W6mp2KLPn`#13zb<(ql* z;sMU^17Y^#n@44X5Nr4v^zHL4)9+uI^{T@|&GI6f0f`Y={_r|$O6zkG_rDT91a@Kk zWJ`g2V!@fOAyW`5yLn6^Rqd@e6>!=vj4cwuKjiE6L&uscPDDtX8OKsa+6Moo(k~6rpU!4cx?uB_hf|&VQeeyDmDLTm}^^Of!^z0$2 zTiJ-QX&<1_4sGg6vP1oQPrykkctXGZOE~W<*Ql;;%p?|$cw4AfT5Co+cwh1TR@J0= z3~BwMANm`X{B6)jOHfpaz05WRp}634ETn{Gu4K*NV}9XFzrQEy3jL^sl-_u&@08py za6`<$Kc*FT-n0LJ!|Awl)@nIvS zH!DHMeD07rb%jq5Gw#NFc{w6~;b%i!YShQ~ao>B4cL5PC*|>um4?pCc5bo-Gs^o(z zJ)>77HJ+@%U_DCl0u@z+Q#h?OvcaBE&6VmP9+nq9K_d@~vx^$T86y0JYUh>DcxU@U z?)G|^o|{h?FI{VUJ2yln+t)W6fCO^bp5-Z9>5T}iAOh^aQlF*bm>5{$D+u5RB?GkI z#u=XjHxEm!%uHhp4*Jjq%CYTV8;QIRI?oy++%HI?SfyAvAZN+fsvFXT8S9}*#K&#I zSESxR7gRvjI)Kunq$h~sVIx_YYevl+($UFt4(X|6`ku{n*3(`N>g6MJxnjqucMY@V zoR5K)C0?m3r_!h2UJ~E`mjUws!PL9peyXSSASW+M_ep1Z{vq&q5ijutT(6BoS~P_4 zOD7u}f~D0Oe>I`u%5X_Q#ekW6DTXGwevs>IIe(Qoze}UQ+{Ld2M~z2dnYO`$WmA zV#zWk@m)_3d}&}F;)rLA2`+qJe_0t>S0o;cBH^&J(J21&VB<3i;;r(4D92vw84*>q zq54aXq5W)f#~@x;_Pj-6IXFb`i7`m*R`pBbNcbB4j59sM8~MX=rG$!=B&_2yp5Jcw zI(T#?io69%J0{UrQJ_!X&V>Y-a4JL`Fz;{FJfOJx1NT7;ID|0BpD!T8_exyR=PDjC z(d|M-l|%1y4v&vxY^?`lCsfvpBlI5&^%P_Za^!+=Mw$5vH%)>ZLBblO306_QLOhq& zT+_ZnS@S2W{SBiCn%1eABvcHj>PDM0zC!0O5+7gh;4yR|NqZa=iIVy}^^{g_-Di@n z-Z@MN-MY-OB@ga1XxsCz7)wH_o#$|&4H-t$DL5}`z1ex0FMz?LB+TgpOc6{H@7jjw zXHHx~suc6RKcm_L==%a*4K3$ik9kCFCW6Z9%YD$L@n$%igKfg*$Aby3wJ!3zU3~Y2 z-H_HJc#+YmpV?oK(ZL0B;E>OhMc|Ol!#$M32{$&(M>a8=2%h_=1$L(=c-+0bSN`&F zHVzwWf^xIlYIMRZLK#*zvKU_hQk^16OyLY@G=Tfzg-izk+C@B&fF&Uif1{?B>q9?T z)arpa#6v`l?UL=^8J~SU&I7nzqLo*6v@}f$Y9cNUPUdG~XVMgf!e;Sv7Av9j5FXly zA;yM1sTKRKz4l7+6`}@8Do)FGT?AjWi7yPEgu4igMwl7Se?QbJvGQ}Ne1MVs;y0i` z-+sY3o#4EzfJ|^coY8ZY^aQfYEE(|F_G?@6u}fmgVfdX%)0DZT0DjwQQedl{Nt#4} zt)@dxO_{I;)SvW}X-C9O$&{ z3aVlDz-&BdRWq}@B zX6-pB!aoXFB!44)aqL>%`H-Cwf6BKZr*Bws*n^~9Ej&5xQS%bYi(K`-XK2B$=Rokz zUA`+LRsVr>JC}Osg}0{LmG$~wkp7yc78ArpYoGO98zq$c^d{vOt9^hNM5@0V@Eqay2AX*2!1s)AB6g}m56 zZS`=iAY2ZC{C!&y{dUe>-CMiL>prHmpC@IV4K8Y~igNtbh_opcVg4wNG9sMcXfs!j z4O zFtkz-b!|@)Hwb|rXNt;tVnnAwpfXA*q5Y~KRf+<>j#!tudjCaY^XY$VO-3HzXKV7@ zrd-m*@qHkn46?x0t1kfvL ztp=4;H;D0RYOa%;sj`fh-vu&PR}lmk5K*b!)PA_+5$%f7oJNm8OU5yo_76U`r|DrC z@O3TgUwR}(lPy9(o+B@-J&ER>@Yi#N9e?JB?G~!9+Emlan~w`pMnPBfFBZqcCFM*i<<7a1ZK8JmGG~bqW+1%F^fJp?{ zc0s^(ds+Q31S~)cI4A#4Z%Mr)Nxm)|p47{UJ5|b(0%CqY6^}k^dc5kZ)gZd+Cj4H7 zPDdITuA8sn&7#mv5}H z#j-n-f=FC~94!W$aSfwKBgE#Pdg_?17G+4KUMRN&r^LhWmyn$w*`6z9P|)uM6wGfz zZa-{|j-_3^dlJM5a43oq-g>M(9FJ))l#NxOfMRAE$W?Q!RXyiw)qZguQ8_CYd7~$a z14B`P77V=-ra+Qt75Xl2Ox~PSAEJ7-ty}Y&0MdvVFNY@Mf=$rqo^U$g>y3&YA*M#H z!Us@I6Jaj=_>01Ni|4yi*JLbvzBCjp3^uyM?>ZwX$hXP@mFB!XHz10-KVBrq!vl=S z=p*mjTwCo(v?BCeiLGwFqfh_nf>2iKUiPHYjVd8+GL|=lyK+(_Mm*R_M{o7!*JAA3 z!^np9RH`W#$vZo>GLZw#isE{`3_~L(X()-D zaqftS8FzvSH0tz5H!gsWqU9E;my-gB36j0Kb#d2Hty2}S-;1f#m~zqW@kpv%vB_9X z6RV>x83;3k=((PhDws)Mtr#z~BG29L$?>B0DOQV9*8Ko}x_{d39E%aC*R0EZ_H5dL zCoy=F2QcTIL2vT*0YbXSG$IVyFPf%phVPg)&tJ!wCH`3c>fJbcr}k43I4?ix0KaB> zmQbEw1_>gF1XFHt2UyR#Qg3lvZd^8jr=)nf8YnoOz`H zM(F|ZK!cRg7;|cc308#;CY|&q`)wRAO3w`|Mcv*ih)_35g_)8Pd9Km*7p1Tr65~1g zcB5JYn!1>Z-Ku$XL;rGmSX5BL= z=)|I0m*o`YG+&yxTF1cTjrIH!{_U|x9ZQXiAbm(8Ur-<5zxRx}^*jBdU3BQuW|e49 z8(^o=)eD0NomfVG!Bx<+C&3TadqywY<~Kq?Kfm1-S7Dl}O*6T{BlaC6p%Ry*X4|!t z0G?mAQd4wK1bD)Ek*+LTA{1Ipj}a0o2*99+`dv`8_a zEkz1NkmLr}CKSnUErE!nW+F>o0~gss;AZ~6jzZ;-1fuIz`7-I1+_X#PcX1A5F@%$G z20@Sv2o`*f6dBqgIS6QKgfNiF6=bK8rYq8w><{89(+RUPjIU7B7G|I=`W}b-Zyu;i zv%}r4)QJ`Spz7db4`BgU)Fm?%{c#HM3=oK6K*aFsXVTTySX}r461lz8nA5~8LJ0Kt zjxR_De12~gja+~dsg~Q$jnjf0_!$DW8h2G?$)EXX!+~2;a9gbtI)gUkVj_-SB4X4D ze2>I|g{Gko6Zlx{*p;Q!I-J_nZ7=?`Lq(CbAqtp=V|Fo>?glCC;Nu`=cI(--5z4yX z75&S%1+~K{kzzbaPh0k^G3EO3+q+B2g#+zs^zYIrs z^NkAuzLj_9)_4d#m+k%88I87)84tv-49a< zGF_xC_Pv%p4qe%1c^^rY?%~E_qc~}m>0>ve8^f-wJe&fyWU}Zt`PS%kIxRiiybV*hT_oNSQxRyL>Pk{nSIt?w!dU?MJTRBwQTH0U%T z@)3K#H~l!_P1JzoxzVz-zN=QK!r??<@gd<|aZV7>(n37rH_|xI!_OPPzXxOxjF6Lz zvLULLq^zYX{D*sNMvIWL`H|)%C$$+vv7Zb+OlT#Oc~>Asl4^;mOh-Ahk3B+;~0E_tc6(=aZD()b)TJ7gOG@qIA~>*s>6&=n$Gu!`9sGxUSjhZJ`d znffo#8mhYL!22y&z6ZJi0)l-S(3G%#cK5NZSONrcA0AC(?zW0dx~G3R-1k4_2U~=( z-mMA&4}8GxI(^mIM8~ETn6ZAVA;E)wk>an!`_IHO`ylMzSoKqp{t3bh4w*2WgfSS# z_D-?svk!Qa0ubuTBB_W+EKpv@`s*xD(HAS;z&6CVQEvEmCHvS35V8Q6Q zk*27-nLXZ=ekOjiH6g+y9e7a|tyn>A&Iokoy}%M0Gkj~2$?$578VErVt8DJeG{ zmTKNx{gr8gyoErH7T-oYLonF_tqxvkxIif49?GEeKkN97_YstJt z0Raiz1LS%%9|~QhE%}(-4>PozHYyX*1xuSrJC&&S_0?lSx~Bpk0LzVN3oL!_6XiGr znRh>uPFFy&6Qid$S53GEZQ@X3=2Eq=t}5iC%!D{T5h$D=w(1*^ey0K!3)LuNzSkuD zR6{LXK78KL@eg>Syx32jvO0h2h{X~gb&(v*$h;bkl#~Sd%v`h3^4DZSg`YxdT)%_G zO7@HJs8r}if&^TL$uk(z;PX|iDA2}T`Y__7l^0a58jxZ9I;KBKZb9ySB0+ZvoFcoG z{c`=l96+#83~4fd6h{X7rM3e>nA+0 z0aneSax3<$ z6!u(oAo@`B1nTyDrM1?_U1@5;n>zkL2@6AF@}A)Ee3Fc{b$+ItV^ONCGCjUk4-SuH zi_T5jk5Lka)5F^GHq@|E`<9(z??Ytkf6(t7q-P`-w21_CTuGkM!e)Q@dye;#qMb$H zqC4+e=Y>SnP#TVNpdc*5F!a!TB;a|Z3tb9|cf%HZIfSZWUaRgedGgJ z|KB7!$yRGbCa>cu5rs2YJJFQS{sQHjJ5>@%;Z7kPYzjGWB7UT6{!3^af|g^+7uNQMg%r!A+4RUKR@hI#nL3nHn?d3xh@c@vDFm(SU@KQ0wYbS7g&-n7L?_ZW+j57sOXQi{BU=hQ{RR>+Z( zm75-qu=3V8i{_&}%(*Oa%^@h(HzPtO(`AAXXnwNsyh)dMu;}hup@fei6pFmWL*mbQ zVAd@)HHC6&UbSrgh%U18>?I}8TH9{-#VKr|*YC4EVr(&9yG47ctHvUj0ef-OMJpS5 zRdM{!peob^s)3$5lzRrGuQ*P$PciIqY%3U_k^&RTMg?0th!vIzxNul*&W+?iX@w1W zxNqQs0&uDP#bp363I*~3UIc_T{!Nvcy3?Xz;J1lk*&wrY1~Uq)+s|q%S!)RcR1J2X z<6IDXifM0ZFLYSXu5ey$EUy_%U{`}HA>E_7wY5F|X(Su0eyXqn`qVM&@dE*&Y3Xrd zc_b13V%?+<-fD%;2zGv9No|%Fi+QqdwPj8AqxQ#0mEy?5=%(JU;K13j`t`>-12{Qr zgtC+vgK6{G37?lgM10XP;AjfVZ01<89P_dK!1j%Eo9cQ^^;0 zTu)=Xr_X76rvXdkrQbfhxC6ByT248|U?4AP18PKnLxNA_fmUwgg{@IZbpe4pBgYJ4 zpHY8e93+Y6m`)~k3zp=J7YoGg6?5ScqzsSW8_CkJ%R-oGdpKKaMpA29W&BPKOx|d(8}3%4<3amhiuwWdVwdZIz6F~UhgJe2#;vjY z8?N*JsDXn>kwkngm0U;e9h=>k2e)06b9uKDgY|5Wu9?(uLw7nH@#WKl>X3^1NY$%p z?9N2~#)f*@GRQSi$bke9V!SmDG}Q6s>J^4be`5V)kr8p$W~BTkg5ZtO2`Gj zvna|MKKZk3&?l$-6u5mM2eN1s>Lz-9bGgcZjtw)K_GrV-mWYPB0E4o8}_(yx$44} zA9a=A7-vsDiM}vzSz}w)$g@3i#Mp3rn4UK?B%NBv*a{b8V@(2+dN40Z9+G>jAV)$9 zwtI?UEIX#mS2xbbc9DSt*;C_9R6a)>9+1@<bN8(XtoXQJeQtvy?J~9*r|C9wOdUb4YDN&`xn~;SCe5VLq1qWi*kgMuX zshV}RQL_xYkBR$D90Ze=3hKt*jrotrTg}!MvO1e^%VD89ZSwSRIXzMoBS|+sh}fRz z__SEauZLn$?cdJnP5)eP&5=O6;tT}IXteTCJMEKx7e_2u0J66K^M@dy~DV$=$2t@ z|8nDGD{$a^lD`{6whdTgJl| zEEms2yug8@I2L?!u)$%Ff=^3;Ybb()Afb6U)1hIXqIlrX(Xay_nn{2>^!EoDmBU4j zgVz_Ndo`c1F0teXdD`Xu^EShS4fg4Q+#P65DQ>bpT;qpKCxlNOc9?5zUL8zv+3Scj zNAy@+&VGP04rmzh7>Z92I9F$sQ|XTMf1@rRs(7K~18vFZ`sIt^(5O_luG7Z@UN|IP4McUmiMGp0#f(B?Po7P z^Pex;m?$;vNOcOzaksgLo+BREm5kO@zyajd(mZb=8VJG_#=Q2ePk;x^r{96FtgdP_ zZQaWDt7PfTCIm*JKV8{q%J^nI8e){D`|jxXiD26vE>`=#?7!L&hjPu^t9mrLk#UaA@tQ}h?geNg%^3Kd5}h%jwUli;59#UFToYWovnX-C zd^pk{q_7mcT>5OErCoz2DZqgqkIQ-nog;BS(!`gY{lfYW-SW`#KiOni*oM7++B+38 z?4k88kyYuPNF|mir4uydJPoiDL~#8fC9?D?hlJB>6~l3G{0GfN(=WjYJ6>L#+C?Clb*_2>}7#rneUnzS?=%2Us54&HhzAdGIfT?<3*H@7&8N z`2Q%K86@rqejgyy*`vTAxy!a^KftqPzaPeW&QM;Yqo-r1|8I zTDoe@Q6JpP)*sR_& zFjgcEvr`z!HY4Y-J*z20l`0KX5;_1ger(nGov@;zh+QyC(AeNnTrw4lHXE|TOUnWhVfC~5dL<(u;WQIeR^ zYVvP4<$(VNQ#`s}GMu{Zt=d-aW&J`Oh{e;s=eZL*DMgTo9c?#zfM9g5f=gc862B&$Q~_JwMp<&b_}Tov|blaPA8> z;Z&YIL_N(Fm{n((f&kZ6nRyG$qFq~Z$p&y*|D2uh?!*|ZSpABes*G+PjipI~To41< zO-T`vXeB>u#z~HRMc;@j1%`n6aurTq)x98ObASsw-KJ5E5TfEyUsLLh4t?7AUbq)j zo5<_4e8Q4|*NRlz5OIVFhUC|A#80B!n~^!Ok#K4DdlXe&#YLf9f}J(~F(u(g++aV! z^n^{7Tq6M5Ad#^e_eKmPeexwK=9Tp$67F4gyKa!nT* z>E)wyI`wVA#Xf0+n1|((irey%IKjcLxM^ASG{LH^dRH%hZHyOZfFS_{GR#K^@J3nY z@(wiPSF#%)9l`EB7gZC(P{G!boOf^_ncsn6kbl5AE405Rg@DEw1uBHvNoma_L@JNV zlg_$brxQ`NJq06rz!xBe-`M#V6UG?xTMJa>b# zk>)5%Bz0cU#@7raBM!LGKZ(hj`wqXIs0{l{Qe#t2cJGkJHTLN}$v z`hLaWh?=N_JZJITn8YH<8tW(e)KG!=BP^KcT~uotliBY;zAZzpi3yR&N#6IE3)Ov- zL^8@SR4Ku6lXG8@IW~SK4Y5i1Pl|K4_Y;3PpBC7STm4FXXA4v@e~aPQfzL?k>c;TY z@rWU#feC9*WG+ee=)F*b7yeQhy@T>w=cpOF+oMHUK3*=LrCAxl`C6)>!pFhpV3~kr zT_YxmuJIDolvsBvG7wXC?t4;g1)sJpK`$0NIWOU!djGIMdJs*L+C{`B$Sy9R^pb#E zaL(@w#<1Y+XZ`-&Xx?eW44AOr@}2(O6y7CltlB54g!K-IB~)K!Eniz&-?eI`1b@?y zBV76pQUXs~52o)8fJD$1+hN=CAz|Og&cqb9!&V!q3(#_ zk8?SMOUi~Ie=v;9m$5!U5Rpz8uNM5x;Gfw;1F@VKp-`H7qaF|P#_$+q6v0mcWLd4x z=#wC@;5Ufi@qKO{GtF2ORt{*m?STELK;|jGZ-!-(ma6qsdw;M+&7^=^EdVJv%xsPq zv7y{M6iwdKi>N$CtO6p=1l72o8l=$pRBlDOA%zm~n#1%$CehDUw!b&lK(J#Mi?@p+ zR!J1~g7gv4*}%E8;e&u05}02bbdy!JB7YKl;ZC8Db=#qZLIFlPX8Dz7KiezsOH-Sqi=%uyCUwlFXg)xT?N%fH4^DLA-B7!4tWQ{=`KQJ2(92e*b=x zsRg<1w<-!)5CH-$R8v~2yWnthA}`x6-NC zPF&hXN5{GKXNN}EutiZp&18t^qpL7Ou?!kSCu(R{gSPz{vM$6%DCR&}R(bzVS%J%e z1y(n)_EEfuU>RR1B$DqlIH31sLRlr6JkmZ$e#e;6c(ab0J^O^Vx+k%C%RWs4$y~(V zav+)jsz&61wrmvZMMYv{ul^sV1E_Gk(>aP1{u+rC5P{1`R>vQO3Exc+sY-8lAy1or zQoodUUIiu*g9s~hptuvhHvKYCoLCbZc=nE~OH+r^)nK(|=~b=oh$fP`g*m%7TjAI{ zvPOilI|wDe$e!YW$@|IjHN-B&c835d9oBSo7@J7O70eXveM)NF92qa`8~Z?9_d%;Q zufyvyRqv1vS?+k`RpBq3`n!Bh2qu%x`}Epd_5L&=E{HE#9C%Jtwj>wG$mj7jj=Zdk zRF@NYakTXP zHqx4qng*<5l|85+fVUVkGQfebS*eyiA1X z25MRk5^k?W@hY$>J`$WNHIrCElXQ(czGMWB^E=4s^9+p`WCAKTF^^));gnx7UCM-} z1m5$+yhxbMd*Rd(?;yZE9*xV!X-GFkhJV#;IF0;TwUk6bz&K;zLodv`Yn(>jKC>6l zy1-x~=7ajIF9^@NLVa|ms#}5-sEYJEZT?dqGiG=jHA$9yHj>HjSvYI>wK+F_jbOoO zv54BsKR3ci5J}pTk{;71Au2JOR?O)Y!lB%hz`w2m4%v{_6+>ueCz259`y2jMlD2~j}64x;oeuDiU$7C?~lslMt1s&YClc_js!+XzXNwOc>5pI z33(JZ3sxZz2%e4(3Tc^@Z5?+E9D9VIsB4BQT|eNnWqN#w6$!rCdd4%}j2S`bJfJAC z0;a?ekupWTXc#b^0!^_hEh2i0KT>JVACiCiX-nL;a*mw#%EDODHGCY$E9uOioMYcv zn(wB}=M~>Dc-q4#x0v<)*NYNoHi$S6n+p2=?sW(66KT2x0O&*SCnu zD0Y{B>07E^)f30JWbKAobA#W_e3VCFq^mQfq2bcC9zTpw`UT&etsO#@?cAy^S;1Vz zKu-hrtgcDn79eEm_-D54Gjb-p#`NBYMtq{DjGUcz&X9IbXZ90~XrFeENfg@BHq0di z_W(-=1%v)CurA!>{Lkvn!S?P%Mdfk!MytUJ_p+`8 zKiM$?T`Yzq zqMkb@w9hsoY|lBakk|y99;My|TDnH%w zJ0_Zs4Bl;eOM3yr9nD9l(cl0v~+iaba#VDOG$TkgLF&E zp}Rx6yX$)q{oQ-N|5_}U??Lz3d(X@>&&-}Z?`tSwA~TkocE34DDs=yOL=7_DI)M%u zv?3yiG^pe*z6y!gqsl>JMm6~yIyK5m#3_1_+Sv&eo~TE{I)aFDYyKCyRz08w>EEJQ z$0VTHz(gv-*@w?R@~6gnukH_~j@N?e(V_D5Ughm#wOqNR&f~;L2b^*ZTA&`q@%1}@ z6xUdHVmr18@;pDh>NdiP3gl*NtvU{$cIxDg))8%7$`AnF7jxUdiXY-hQ#I1E0wkg5 zDLUL*+Y2>Ajye`E5ztC@(*n`dq&@}+3&=`mo*wjT`<5B9{jBBGJ}b^NSI;s=uyhOU zm&f;xJ$Dt@EIZ(?{B<--yKqN@NN}4&T~@DKv+XNc)%jp*0!BWwV3%E#j&&&sH55HI zEubswL(8WP6~UFxRhvB2j~~4-yu7Y8!QlA1JWD2JjyXSuBc z?k~>^w+u%~EclH|+J3jxmERN6U|LNZ2o7-vH1x;4~?*PBeuaP${ zJx*S4DNyCRC0gQc z_>~|W=%RC--q3DQCJ;xK1>SQ%K8>3GaoABg#9YBe%^~C0Mz+idf+S_J%`MI+5zV?9ZP z@6j?iD6YOL%~;324cbx*scP*1PFLy1W;J$3TM5;SjzHi*)(0|5v!r10 zH=%S#Yo2tVceM4?0##?Fw*1l&h9>^{$CvA;JV||?LoH;a*~(YN`nOaeGAVc-VG^{2 zqX+(#o@PRlqa6I3Q1BLBXgmuitWC|j!ZSiYEQjWx^E$5-46^Op(6m#@Nc&o^8r}|e z-D6}b-nX#1n`3HeV>Zh6Rmt0A!4BtZ*sBHZ(g_w_`NVs3HDzyx20{H4)7f=ehg8^I zXu+7i&*J*ZL7Q9>xX{o`=QxC~PM`@jnBse5i_4(qiCHTOPVBSLP_wU=NyQP@2(n?{ z!nfETy~ywA#sf9j&e{2+{WRsg4I1mjmz;WcIm^3JiT2=oRcBj}aY~{1O^fr$!~ShR zn2{VrUVFq%ZR$7q*@Zx(Rvu}ng)&2+o22E!<(cGIL!ZZI-g;gsO3uYfExOfIE&WXD zy8}6tpM-Jvr}M^%v)HY;O%k0(Wt9uAOLIBD-V1!-tv;c~L-jGsU8EVrgpm|K%_=ul zA^WJvwZwJ>$G(`kS9c-cS0TKr`2c2Q@Hl$2s*;9**=N0ZLaH_XI+u>|zWlP|4d@41 zth(1V-pe@+%eRlNw|8$0C6jJSrO4z7-D}qAhXhcTXdgcdHc?El>TIqkUR}4@>TBhD z+R*YIE24N##lHSzl`|1_9qGQJKCv58D%s-<*4ty0&g7{dfs>(2Kd20gaD`p%Tw*G{ zeJRYN3AwRLL#MOrLkAz}pcXKfu6LMSql^m?$#S46rxZ}bO3#Df*L8iOuqnAWs7)^Y z84-KL&A~#R#!2f^qQ;|+Jy>b~%;l%-P& zl&UMVIwCdW8_Vx>hndW+=qvdGl&-xWg0;Ikb}xiYp*`f}*d-6*0uBZS=^b~k9w!1Ud9+}FM>{X;NFB>EnC4q^l4T=}HCIe-91b?@$DJYm?Noc0e_+#Vq>OhF9o^gL zK6&3IwWrGU(v8%s{su{3KouQGC0i|(_O0x)dXKOwEevrhzK+mVB%2q2RDx~dKnTd5 z|4_9Di=8;f!2jj4E31G?Yx2@Dz2F<(MB`)b3GS?d2Eu`bOB|sr_o84JY2v!TF|7lE zG95m<-Yr;Wx>mo+#r3lCw`(iTx9mls=ALmBk2LVGW-r3GMJvzd)mQqenCGsha%I0K zU#5Zc8$$Er#_=e}6_5-0@a^ujkNRepTr9%EdgUe4+l42H^TXl~B!}n)HCM;OK( zi=qwgrVUx?;@Y*4b=+MmR%eqW=$iVUd9ch`HxHCr8^-FH4DIRxSq|EZJ4M#_>EYjw zhHpCc`i33+$z`*g_Bn6lV?OWxnwS!^C@oO>YRca%icwNu+TZ=?nHW;l$>@4xGZyMW z$@h3s4?jlZ!BJ^C1-`Ob>Z$v@+}LfS$(jEJHU{qS+@Tg{vA+*wCka-jFE~8y#`OA1 zo%4GiG?vGC1103=7*y1oFR(lR^6E@tJ4S^~0MPo|%(~ihL_V7szbPRITC-`5fsQE^311W828k4_GBL_ldCE z16woZTv8^755C7M;Z+s7Wl55+I_OW{{qhl6?pZ~kn!eXX3{xuzF%%_2g$nyc9Odw6H71Z!^q*coZOXQ4OcbU0`gij=w z&GxP7e7G`EU(@O_A~fUBQn|s*BfDTe_v;ZOnuU#iNVSY{+y<9Lv*RKd-bl)eRKO~D zD6bTPmLa!5V20e8OZOgBFXb`zp7~XNg0xVcRyodR)qn>jDay$vQV&$7ovX$X>gJ@Cg|paAmE#Ns%$Ch# zLZPoaXj!=?Db17H2~$%wh#ky50cDHlX_Sx(;O+RUE{1Lm?w>c&qG?hIAcaRAz*ueX z_$SRL@&sIY&cgXWALf0dR7`$!T3u)<48OTYJ=%o4hTxbo&%d*Nn;#!Nw4VA2+a*6u zl78D>uXbMIF!FYLAq1`hi9NZ(>SWbzpSXO(!|4NPU-eM-K(`(@uHJ8cHr$)XNNKMs zcs#Gn>2Q}Ga?1t!sC{Xj;tI!!FKZ)Zo?v{EyANC-dUJ)X6T=1OaIACp5~kAGBIPTK zT)(w9E8|k}(mt>m&kg31Omtd)#rJ4)FckC1O-?T`xQWwp^NjffFSxzwWl!sf5u7c@!Vf9%B)z>gY03~~gSMN5)&w7>R|yWfuz%`Q z-KuPGuM3o29=_0eUy?P(w_L2g;BrOxy7P@s(>5EC#YCJJrG%bvZ(2;ZNkLBPq}<5P z(NSG+NiJ=GdS_og8G6k*(q6m6c_Ma)T^4j`9{HJaed}3_Jz`#1@2P}H2vRVk=mVS5 zCyrYCif`$s)zLxqg+e)c!VgGZeO*Ex+7Xd~$HO%)xU#jHXWn@)9Oi6(ZS`pAUlq7A z-`4bZ&fzxpdNkZsoQe_#Ms?aR$2W>=bX_fN+RhsxMvc|Pw1Iz#4Zy)ETXE- zvz0)%TECxB#uwK_7Zfu7{;g?milL~lev9hzfFo0W>{GCem$h=5h~HM0HatBfR46iC z4$r&c;|+n?dkl|6ZEvq)EK&jG)}lqxoGhFbuzGYE$cSH6>NG^DCd*#G42oS2r4Yiy zog4;-tTxg;CDa|~3MedtwV>`mY$A>O(W^Xab3yr$QEu+GPj256LCe5@WwPy4i=5|( zB{A}M+R!YY7_KN+9&I$~vO%|rj6v4ED_4Z%><;fuFE`nkkdZ3$A6A^+gyI@l>6xgU z6`Y31ZClTO)5z5pmY6MH z%CyZ?*vId}gk}xC^}Y_W53#3>lp+ z?91#fH6i4~Dk80k=JJ?QLdX?8)|jD-oF-`U6|gN zT%<^*a{JnhKNax8#VHnU7?C{C$+_tRonzHZ1>%(h+)*j)#j#g3AF1u*#eKkcFK!tk z-W(Hb)Z=yaZOwug6P}il0k2;UP0`rkd#rd)0k7edcCkFEgBJ3LQKEHnCM*!Wq!L?Y+rQE7tm#M z1FvHtSRM^b*NMhmI(IKDlb36<&^M^vvF0c|H}1%U*frt!OInzkJPry<%^QYip9;;D zm!2b2&zNsloOrQ-O$1bmcT_`vDW7(*g-QVg#QxGGvwO9iTBzX}w_p?DRfO#JTyQIg zm5IYee|0)jxbO;%t>=Ol+++? zZSHT;s5eA*3GuXgcJub>9=x^4iX-`%#=RscMW8y$yb5`tGsFmb2MeFbsW?LJA9Kz3 z)y|wDVfD3&(}>AL(vw2=X)aajusf0HWxY!yADsxi9uS~g?F8Av>!5oRN%yxkswew; z9ANir4i~QF)@KYiB(SBpI=SIl-4N3^rN?e^EHeu6LP;`Dd(aOC24V;Xwp~W+e+Y>z zd{MBg#`Gz@NC1wB^a9#a?+^?_CvFwB#a^!TKh$L%kFG-(iRb!uE2q**cI)(&pZjCV zzRHKj9_8Z(*F0$EdU~iS^(Xn}?WKn!Xf||7P}&P5*mo|kYQr-_k`$1uia_n4`pE}a zFvT*#3okHNjC&Wl(v_J+K)q*(LKIO7&7ksjRzKx2|=9xDa$NGaEhGc2q zUj+fLZ1VV~Su?eS5I}!C*T0Z+{GOO?Byqu3xf?i{li7vMpO!<)j>P*N%Gul$vJ<0# zl##mw`S#a6G)MB*pwYwJ_1oG@1>U0cz4tv1Jb5kSc6BU6`?`vUNeRIr^XhK3iEi_t z{5v#jx!l|J`Z9C%*)lB+Mr{5dgEZIMe)b%o`?hrf`J@DbdLQz5vn7h7XetGZsRhCD zJpsJuuM=-AzpApe{j1nAupJ>?CKqkXot<3at-CqpJ@#;YrmL1>d|{l+Se&I*5u(dv z_NG+8?;UGRZq`?~^Ver8ak@8jPp5U%DATJb{1Ww9%TDe&pviwnkO25b5{sAE-a=hEE3Hp zzicI-QiQ5E?MQajm=blzL{GWzdVN{b*&}Nd!~WJOMeqe83=EY=5jseLvqk757a;|Q zPV$xiuM-u{FJhSf6V^kCEeny#h8}~~OeGC@A`R}}5)%8G6pBC!Qz-kB3f*;LazmwB zSE?SVT^(R(;q_^|sv&oB2-I?6orRutNXf2Qd1bfSEXXc(Z_vO(w&nzE?KbwhV}E^? z8Bc@`-y?Dcxq#k4^y2h9CDNXaweodKV&+`z=3D8{#3*Vpo)G6%Z|G9=3%1)3FB>1s z5MRf7 zC)2%?Gnp-5?G%ERz3e^ra_3im5toDxP*xB~1Ic@E-y4rq*kL1W;*~`o{lk`MQ0GGQ zW$;czHjbH4*0^bLoDLa!#pjc=lf>JP=PivoG+I|IV62o`|paTOy!Y;Wz_3rL`p^xWAj*-=6>b4D_lgw&3mda+FF z+Nx<9{AJmq*jG5XM=n$P!UWbqKa8sEXqe0VpeQ*S!FXAFvkO*z)?ejAnPtzBz@+zD z{=#qcfw|{vLQ$4k6r56WYo|6(Fbg!3+F_M!RuBjD@tUlup0_mf%+c^UJmi;hyPnIp zAq>TBGHxnvH%T!n)|3dCNS5U$^tdaBn@dvU@VkVys~+T+;apdr6!D*)xKCKT@=X&i zFQIkeTr8G%39>K)*ACDOd2sM5$;Zj@n+T6jc?5HANk@g)m2@~3NJ^D9Ab}e-bvsy&I^#$;NNg;K`JDb6tTyf)QXh9v(wIBYK=35 z8z4mmj%U&r4H*+Iw{GeiZtWRcTMuy!LgmgW%|@98n6<*l90omFZyoeGaUO#xP-cN_ zv>e*nadjM1amQh(gU?xns}OO51_SN91UALVbgSGav0lF28wL}x@3f3)n@6jb`QH(= z2?xYh#a^RY^n)wu)uhrUisB`W2q4phITO6DlD&5Lvd0+UYN@Tv(>bZx zcVAYzVlpVrbm@kOX+a`jJ)v)$E8807+FieCv2nM!0HJHQf*TUYT95+1TAT%LfDCUxsYMinP25=EM35uyjB=R(H#5sl1(PoK@ z*cD4IXhNh+Lvsp82D4EJQEUj}1$`+ma zbLewS$npS6)FgY(0;#j}vaN+X_n?glv=p3_uv5d7zN!Y%uh3!~)b=xzzvk&-u2@5XKr5$E{2+hS=+SZoS;VE)dHYK=HSKXe^D;Fwdv z5dpW<99*(FOB_g%nw*u;d15&EWdmoP)E;q-B4hwPQsZy4Q;jS=A#yvDRMgftNE>8l z8e`>Y$|&0mPQRDR96EsDaT(+;mIvp@%LDIkR36Cwq9`;L>D%TOp5)fmK_TGD?C_n} z42MQ33K9Q)i2IefU+SYe+ zy88-usqBQQIk)mFvE(4rsM^(muqe^af%z``t)nEn1&BbV@5e=#u`t>kmg)0YF>jPY z1LuUBk`ntj3pX^hQmS_AwQJ8g+n=RxujpIL+Nbezlw4saz0R;#nN1$(Sfyrwb2*nn zde9RXHfA$7d}vtDf30spZ?Q-x{gcFmXe0yQk`~Qo$ABYiOtVt2gBMJ_i%PXPI_$`b zjrcZHuP6+!IehYVy6C1=F;-K*>njI>g;y2I*R5zr*36bb(`Id+(h=+A`HP#%LguaS z3HAf>@_`rVrCA*JVh7E0xQJ<`l7p4Z!PZ?hM&C~3h{O-#A#yWa-j_1GPK6sJ!OL{Q zUoC+MU>Y4W!G9#G-YU`?&10i;=M;;Dol!gHi|5^r+iRA*)REZa%`&`Syy_TMtlZik zIw2Kj1R0^ThyPwpBZvq&#JK+1{i}RZ^=fKr6BjLLpG7gt0Kv>FVix)&7^$fa7nb+ShILa9xLe@h}B4tIiZ-|hX->oCdk?h45c1= z3_-ED|6Z(8`+{ZiR(UA<6anb~YGdi3eyZ8yrux8}ZAVf4lJ}BFlSAkw7ci8k4r6A2FZ8;Q%b3*f+kz*9@{ zqJla-`psLlaoZ~ji^1ab5BRyErO=n^YBgTBzr)_ zwhu>IJC=jM-Qc)EeN*(pj2@+-IFpvfHcVoF)VPtuRr)Z*s!|iTWJZaI@Us-}Rj>_Df<{Q{o*>J)N+8ycHtSfSQNZ90#*D zQbv!4by@whjBC{;ZS6~#j_WMqcOp)Mk;V;9V6C__ZUV2{I|K7-wnTYsOYWmrvqtcf zrtlTFt1aQ7`^NX=pSd!2W%|-OvZ&Ee5w<;0GOT*#vCxwf z8RTYzVI|{U>GMX>&FD0>JRmunb=Y^{&Qx93YfY{E7)1&VJRk3?60~lSSDJ5PbAa_T zQmeYNx0n4)*g^8@Na57_|!$#Z)Pl;S&+k#?Ay&pL1>Vd@P!$8pX5 zHpy4x{99runPALY0P{|}p0%V0c6l7!*OlUWIL??SUIWiIWHEpx-0-t`VP9c6v$c+{{r*;jR*_hoB{uQze)BR4J_oF$dM8CZ#J!?0os&~V5y@Lc5=vst$8j!R zHDN#dL;5zw)A^1N?)_+*iOg#QEX&=4brH$3{DTb4AIWgFqC~N;jY%e8cy`H&=ThTd zMDa$u2oLaVbMIdWEUny2GTJPWFH{(er40ypwwf_b^|W#!s9$(lM$6JqN_;UO;2}dy zCOiQ0fi0&lK-%FvKbSDT5mU2I$S=tyTwNuZ4)Q`@pqJ+M+Gc3)-Hc-|Wt_NY zFU4+vlHQx)jvmLL=_W~GxJ9*`1I0#+AhfN~%8@|IrfR~z$%F3JK-qsA5McLmu3!vDC%7;g>2bpvWRZ3&zI`pZGV7|Lty`ELcY2^b#UjvrbcZUz!BdW z`JMGrLZaAL(U-!PY{MTe69dbE{>A;r3uN%P>NMRt_on`Pki8wcY+eKso? zbJBl1;%#WcO5NwfiLpFwVUMnla=HZqavtQEp64oAznogy!tpSU3+Nudi~9@Q=``Je zhcupwxOBq#B*BUKmjiY*(q5)x4h-rJ=j2AqE(-9qI~dOO#?&j~S0eiAGf$4Z7a2Im ze^1aejIcc(5BjWuV#Q5uPw=J1JtwORh8jux+;FVV(_%HO19gkYTj4h)s@35XM07s8 zS=A^OBL|Y=OMhXvV0ILoeZCoVN*^yMp~u==DVq&GS>N2NwXiTNtmU3`lhpuGBQO(1eWH;hyjo`1fDE`GTQf z8#YX8bwe$m7~FHY6R>9CPn@bZ6|K60aZ}f>J?_-s(=#|6=rK@*jH#1JYa*eZ0{wVz z=oXW4GJnpt+~>+1!Dzln>UVX=Q79M{J6HU{%+OI3vW%uDOH(4WJ#RLrsemXiI? zGQGKhI?s!%3a0<(H$decP>AZ?Rs45B|Ejov#|v1Qud}N_&GG;G?D_8yKy5F=3yiA& zGqXRloF?W&F%2ga|670bzh(}_=M7?dq0;~9-`4#%s!sR|p?q#nYRUiSH^?u1zGl(7 z$N%mOe9`xT>h-u9Bx&1$+8nzhH6^UQDHr4v!rcc03lr6{TID$-r-V=x{5)6g?i3Pa z3n*ME(rvf*7E$7r$%Z@U&i;qSw}JmL)vRdwBoz=7EnAfi0fdg<{F1@_YgIQ zbKb0ROX29a+CpiFSsxsS_l=i|kw(U)AP*-8Ei06trh&{3nT&?Reo>qNy z#W2%ilsnWGCpgBM5$0l7Lg@}W;&3rK{7}zKouHMhJeIZ|U-zW=F2N^$N!=G2$0n$5 z%A(y~T0Loy-T1o00;Dzo3Gqb$f`)!gGVt`^%?b$;oIsnE))HF`v(u(Ve&3B|rR{Li z7g}aS$kU)aWT?6`Z+p}f%E{AOV;*3hO~$M-RW`PDJr#Dd*+W5b{t>NjqUxh%z!&e- zT1gO(7owldEYRYd$frTd8}y^p>*ci9v$d8%G1x`CjSCLJN0f+EMzTr>4$qW#XP5*Z z7mlfGOp|=tX!*YHOa>ev94n>X0#@|&lTM484qFmpl;ZJR)p;aFWXFO-*15vLcI5VJ zyw6Bj=k$(ZE1nyZYQ5*I)AVGY!e%A%ryf^>Gqh(386ncMHAnEJ5XqeXvPjut;S~F; zcIdB@d0{M~@alx1a^4Kzpoc5nlSD!BvR*WlOPnZ^e@yCTKiXV_l0a*}`EmOGVCa_E zv)DA!?cc;Kk(W_v99Fri;+}1U2i9OAWD$gzCKfiv?p?P)2w@s8cT6Gnak8!}AL(|( zvLb~H^0NEtt$FWY*Oz`(EFFSX&cFIKP6~PUbuy=Fpyk=ZqEUMLyR5HKk|{G!lPgnv z%%khNj55&D$X^|aDt|R@8c&Fc@TNw|I*f_dIS3g$mLgJQr29KoCj-PajRBCEc!Efu z27V$&wBzHK^bj`jl)T=)#TT>dA1mxQaiXE|Y$Xyyn)RqU{F|-1J|7Gct6b?K2A9n0 zvj0Wbz*hi>R9MGO3qOCL4#i+4WU4`>LCfn}`>T9<)dD-}CkQK!E?$jRuPf_1Y#LvG zj<#HJ*4fL3C$%5kuG@H_lKMxC-cXs5F9KGDS5WZH0tDDLcW@7nJ%IH=<~|HI9n>UFK*P^q`#6rZ$r;snPvKOmHeg}t#jQvCw+SI3BgVF4g|n*z_` z_vZ`Ryq=2W_2O0|X?{v`r+NH5S5Q<)L+m;yd$4caPZ#k-#@Yz%_z=|Fs02we;DE8b zgxZ-u9B$kH@vq9u@&ogulzNE^P+b6Fmg#)#Hq$*#@k&ufuH9Trh|-S8V1V+)Vf^GR zP+8B%s;WlVqdbzp5pD*nvv+xL(rc1EKEnz^MgOG01xTS?UrhY(_K`#Y_Y6+5973GA z%i>I_%YK>E28H$U2TAx!9Y=I05dOSb{KS=2g{D!*NQHSCtl@0^bj@iJqnzQd#=&~? z)d|zR8YcNn`d31>kO_9fG(1xN^N>;NB2@Kvq6Rev2t#BHcMewTvy1?sr1=Z{PLU{d zr|wS+F(^NBk0x?@urEfmTPP0uD-}}`erYvi{h-&zA=lyFf1-du0Kgs_3Q$BQHnTiqCK+5R}J(*B?lYE#(PU%IzPN+v( z#tAN;^AP@M(nX&_qFExWp;k$NVUAoN0vM1=`ey6(E zCW8os-jHlJTAcmq@o!C`029eNE-sgR{(u>E_m7Dlz*}IYw8Oj<1Oqe=;tZOp!V`>xO1gC zy+{#I428d5sNbSBzvyY9B`3B}-IUg%a_MwT`U zIrQ*6bKH|2z96qtez3%=#MWXR9b6kku2^O`Ur4I1`G0~_1?(YmeQ5T1$EcC9IN?I^ zf}Dk#($A^iA{7V1kzU(eF+_ai3|rw$5K%)G#TIwL7Yz#O$$ZUjY+1ks_s=Oq0Yrpy zHu>j^7d&sp1o;(&UvLyXaR=bof}al(g;})C9Xrgm34zKJ(}BHwW5R@O%?aYylbg!% zpn6H6&tCS@B5Let)Y~l5cm^8oFFobM0xlK#qp%b`JHc|`l_v=Gdg~HEA$84B{SrGj zLohj0x{fe`(96ZQ?QW)VgU)QrU#OqSA@H?NTLZCG{QY@raD2q-<4I15tg4h|H9vu- zHhOpCUYM((W>m@<6Rlr}CLP2wRIU|NlYBh8Db~=iC}9%+x4bK0dH&(dW`F0wuQRG~ zhMsYejev$6L!3I=zg&_j^S_Ly={m~Q%y-A6&6#ck;%_jQza zH??~SSbMu{m;Z|VnVzoN`B$~EJS&WXI1wXEs8De;2t-vf`)z{I-RJ@3nB`L<@B+`$ zC2vN9-aWJl^hgdQ{a2F;aXnKNubgy?@a7J7>ki^>w5=^WhJ~DTt30*X(5t#|Zz?U- z#ZdDpuwSc018A<)R|Wr1*?j?JXM10T_-uh-CT7C~@Ry~iI3YvMN7+qYL$LL0kEr(7~XmMNp(=W-rGP6^sjqy@TO)oFy4 zoy=qQ*1(=m9wRnU|Kz5KuA|Yrr3BT;NX#mb)MH$=CZSu}!l#Y&q4kjRndo^c8Yi>73 zG)5CNEse4=$S0)O`yaCZ-+gr^jo`+_40_Srwad@z82rK^DDU*-daR*PhjY$N}~8ga@&xR?ufd3ShUt>iDH6X@u=cfM7ET(vBa=&x&cD>r^O?Ch-2 zKkpuDAQWJ|qB2G)|G=*|)ens{YJpGC%f+V9%Z?gzA!wiBt>f1ssK*bu${HH7f{C~H z9!>-uwkWCH)>_AwwEaUp6fRz!;SbSXHXMs65ZnuCO)V#Z$HIB5-mT1Yyixj= zzE0;sLEl4Od&d1B?u3KTF>?r>nQ)r+r|a?;lW8S8rH$udiF}69(#e_5zNx-^c9Sol zgx82Nv~aY(gB+o@ZeCt9L{AYT%f68GhOHrg6Jl*jWaz$5L*1!T1gT(g@T&AJ|9vRk zt>>PVA6rm7m`}(vT@T~Wlzpt991e`(l+JY{_-!NV(d ze;+Z#!NnSJ*K#sjeY!-^>qO86g_4;mQbKy1jTp-(6s|@!AomZ$G@hu-x|rGXGYhQ2 z`foWD305`ccOM57f47mM8gIE3v1C>x4-&_~w$6<5g27#Z<*gi?n5`k*4;6oXmO}`i z`@#;b+q1SJmp~a1Y_ze#U<9drDj(_QrAhwc!qO*VF`_75UND8FY6;CYW=d>`Ffzg| z{M(mVyu7OrA@(mBg=mCYtNjMlL;P_=AaR|Kh>7IlM4)R~Q=@89RF4bA*Cb3@eIXs? z7i7nuTZr^*zfTan$xD$^4I?|Gz0;e32LHt!Z+u9+FtY#x@G~AxlYLlLBOwEZ;)GPd zskq>SppCp#qv1n$^n@sl)$%NvQ~C}&jdA=-#r;B3v_>w~@GnhjqlE~TX@pXJrtobT zKC~&v<*>xyNaL<5r}S*t%^3tigCaCt>N_>WHyTnB5tNrIB}!0_>HhmiUnfqeMgF>2 zM*s$AGkPkYx4H(lcX*Ua$Xm;5A`lw%eQM2*v3#*GRE4^B8S(C`o`KjB4GbmfK^&NC zK1sjg(HyvK_{6%*(R6e&r+K{0xu2dpseJ%U{N}8WRl3e6h)RDUr8;)H9q*@;wNkX$ z5#B7k>K*v2$~+!>!H$GStr;rGVp>P`uK@rrt@+lg`!7H`LhbDr5_fsi@DWPEtcW=T z(|Y6!MV(9^`169~Z%I!nHmMqZow!{rXxchB_5U@>Q4$c(+=K>WKRdxMK%5052EdoO zxOYSOY(h9Vo7+Lnd*|BPvZtDX%q1t(ZfyRYPcUE;UH%9df9WkmE8waA&miw4AyFXm zrpv3BwmYwy@bGGMvnJmTxV;cbCXFZhXkQyDgZAj=uQl#45H*y~1Jm3z-lHsSW;bz`j|`zYfYcZ5%_Cuj5C5|ct@ERZiH>+y3 z)b)P?*M@|bSfiuZcvi`BAtE;xsmOPAaBGIwwRxR>Ca8g1+BG0|O}^8f0oRlM3f@nx~h! ze&Ud=bD_aS%zMD0J!v{?Cjp60(-4AUR|0A99E1AoJY{0e-JPSK(Vxtx%IYJ-@ABT< z7xxFsFs0=)(HyewB0QyW#**5zTjpaVuN9(QsO0{)oW#6A3jNxb>@Up6R=_$Eg1(h0MlUX#wrx9R-7nvQ;f+PvS_g@b}_45y>rb9 zl1@|Yr%xnz&+^#B;mJF1Lx1Y$d$cls-dyRQ|LO8FDw`MLELo4PpNfU$L?t|VN~;bf zZ6zS!d8jqrHhGX>*^mAXKkblkU71C~<19nmzr*lB5eO9SGG^|M9Xc8EV z5K=It!X)JVZ2iK`2d&%Yd~OUx^DYF2;%hrFw{5zZ`x}gKFn{9z8R-2-ZXN5typrb$ zt3PsCy|hr7wJSzmee(IZEGj>u+tkCCvF3e(9aa259l19-c`GJG>#F}EHZH&{Tt}B# z{wR|-4P?~D?@N#hEkHB;modOzsv23Nl>CS523q?pD4CeBOCJx}O&6(EdA^Vt+ER`G13fV{Xo&C}P#7r=}?C{_DX(%dbY z4pKoF*}JPYc`*Lyv0!gKobU;oCzM|7%` z9sEMmT*UiLE;{FljRk%i`Y|mA=JIEhyEHY+z?Y$G8m$bUX-8@1boig0B8McPjZ$fz zj(^a34^7nA&nxi#wpGU{?G(Ml)dv`nCnMl}y|$iSUXF~}cQ&?>NdhfOp8lhZ6YPf%vJ31}dlXMXgNpGM{De>OE-)3nEPfp(5fRf; zjSZjfRkbeD2By!1_>tYA>@=;`p4~m=GHFzDD@wEQZ}E!i5wPyxJtBfA10vpDDvX2Y zcpyK_Tkz-pR$!0qG0=;dRvV$`f!*s(Y4JNGq=0yD>Y~$p1smT`(+kW7>rt?LMf-mk z3p972=4~z&_%^NhjQ>ag{fY30YSIIM`l#!H=Bgw>)lxXH?FQkhS)qAc_#7Lczwa4g z&dvY13?B{vVSgZ|m;*1h1&|VedTs=1@P|ZCC5vAgohe;@Y>y=(*8PAhha>PoKb3(( zP=LDA>o6U$e3+~_tFlbHou|4uXd*s8XM4Jo7`$>(P8*C8z3-08g zlGNE;VJLGEw!amPQ?m;|{P7cCh-)p0h{GPj=9$fZM;}KdZEb?7Om)v0^yMgT*l+YF zYbPeh4yqT9^`ER)f10T5RiP3DV);xQ>Z%`(T&Z@hIfFy3Mbk4!R4Q&Xi)+n&Up?Ep zDK`X|6}37Qp7SrH2=hU09fSB<-(Jt;n;WI$(%yMIC;%;{MJ*96Yt)M1 zYk8LKb*&efimOb0|9TKpRtUtNl7-E<=06@}4ywCfsKL312{{rU;zrfR$3a)W5shh? z)d&B&4uRHv&B^Bz_Il&!hj)l&$tk|jF^I>+fBgp*pp@4_&!C+j%csALef;gbltum= zrk0TEy9;nz^}(NF=e2vh@9*Dw27q>#u;9!!>*AzFo#5FRz66ClaUVVlpc#Rh`VjT4 z(36t#6GfziGzto77*9mQ4^UH6v|ooDbUGHXsLHr zQ=u8MeYT3~xEC+cCKm0$sZSgWKqNFfn=r0|Wbzhj@FMkbt|pyNpjPz+GDaAz5A&g} z8=&1FJ=-r6&~yyKLDf5@JXnrPces1S`oJMA-Li$<*@!qfzMwtVFI z4B(;Gun49^S190mH0+U)t*-7tf>zb@IGBhr06e*19XJ06w=zXf3BnSNl)XPe+BCWK z0p2(!#7(eDB{g5p-Y5dD1uW+6Lb@Lj%CxL^0R#2rOYjC>{$T{4P(szna3uo9U+Yo? zw%crC+ra3F3AC|5h%~%>M5W>`QTzU`u2mE^6Ap;!A{B9|F}b0bWJXpXC89AbYGx_1 z9VeK12!_D$vWreCe!YJ-OC|*Ns3#pGMZ)n<(sB-U4<0K-?^!M0zaz*@*a`anP|K!y)+(=;uFeKF)2Cr z?(nJb5dJ*@aQL}kgxF_q5ae0#UD|2aKr|DB2#Z2t{b#_E=fQbpWi|L>}trB=g10KQ{iO#o| z3&l13v3fId8s8)kbo2po--YPY#Kjf@7eShONLW&%_SJYZnyKF?9aIGMJd|D%fG*}~ zyfuG3f;mbAGtnbRfXYjMCJ;_U=$k}uQhP>BEOxg}dlTLs`a5XAQepaFRDll9|W zLEVt83&uAW;0#JxymMZ`C4#=+)M$;lz*)dFM(@KOR#lx56~qv)9}KHd{#~ zpyz3?(`!(GbloQoEo6wUH7>+LS z^vr$<>#wnffzPC=LI9xMj)Umxf1thQwJ{MEHit`%q5w{yaP;QaF2T$#Lr%KbZPQQ7-v7>Z=pD5qHE>`=(j7rZc_I0Zu){V+1STL~V-ADVp|6m1q z0$_uzh0;j4e}=kvZQO>5a0iMVEm4cotvif$H-Sp8F#WQ=1MBnZ)m0W$sVmsjNvJJ1 zsOKT(HRDL9?zQbf;hl?2C6O4U5S@H!nf6JJM z_6rk#?93;{wofCh`9lxQQ346S&e#g@cKb)({YL#Z1a8-QqLg7M z&ZIE6At9tPk#GUUj+!mC#Ot@5Q>HF|1MEPpr{mP-O@Ep{zDXUb@-s0fE#l=IJ?NAV z2(&Yi_NH}a8pL!4+@=@Y;3q#Fl79kTx50Zv^3h(jOm_PHsX!}S6O;HeZ4#scG~1Y) z(H!%SW)l%X{z&UDa#OVYjaHHv>5|TfDb+>%CwM4!#WW+;vutul3Q*5oX`5*tYD)rQ zR75a~%JZ=iQQ;GaE2JQ@_&?E$J+xtOTAXPi%e%2+wXrq^9G&hzYlv$5VsaJ5H9Xwg zAN##dMjj+|%Z>qnZCXk(nlnhzWVHRc}?X9HoH)+14F- z!R7w8!5si1{`}asLQ)lS&?qUAxKtU7J~PZZr)N#zj8o@ijcFy&zT`=od~eZ!W5D-3 z?B+PY!vawRknY|)W3nSl{3Ur$Wowse+A(N$ng{9;b^!wSR#?%_^+&)tJRVb3)k+{{ z_g;D&f5l_EbevxDa4h5W+SPrgX+DL3K#CYv2%dBwmI^`y@*#yA^YME+6@`ydv5y)m zY6Lmt56yH#@QHX^{QY%2b(#YkE!yj^iQYwRe&T|Go>uiBnAu!_vq4NjA6ySLp%={5 zJJZraJXUJ({!MmI1L;vyI6}Cm)G$!=v1%9b4x=G3uqj)SBlhzWT70~Dr2nlU)`2x{ z$p^@8<7F(sRHA{`o*eBk;n|9%gHSSIASqSoPfavAR(*mS3m<_y#QUr2x&_X1MGdL&z&iRugdtWsIdtrp9+n5)ZUyh^0q)DGBvh1cU+Wa=EP= zQ@gb62;|rDu`W+~2XJk10MqKYOhTU?AVJ8O;ZaoR#sruFmt+I&2r(DRSH}HDcFRF; zIA1;-`SR3{t=+y!_IKNYV=?>hBj=6eovyWPjd_wtthm4z78|N@?gV?+ z33JcDx$+K3F%t|H|KTg{DEYQ%)U4gF13R)R*ZRhudPw!XB3}r>0BK~YQ^tpI= zIRADiUl|lvx2%nZfg~^t1cC*3 zx8UwhLeKy~LU4DN-~@N~;O-LK-ANcMxI2LayL-qv-+9ly|G%kSQ?+YN?diRw`|0jy z^^({g7l9`1US|@vj3nytNMi0iMSsi|2!mj$5HOPhY<>tkBC#1~*AQB3>Ab>BrS7F3 z@b&knI~SoR$Kz8iAQZxd?ceqpV) zJU&hVW<=Uj>V=+VZ7A{`V!q*#8KQooM;^7$0$1}dQqpP|Kty6@TYMabY-(^@>IYn1 zU6vQ@Vk?w$Fg_GedGW{E^U)gP-iImlL~nA)ZNos?-y!701k$-7;K;|01%Y0^vDVqa zMVMBV4@GfMwS{M@u=C*qI6$bkdr@^lI@e5uqp|zz{y_P-fLx}3N3wxgFl3DY(?Q}s z!ry-FFn2!@^IqTY5=52c&_!$)&q@W7Pz;#YK!gfc7Fesn?u>Y#tSH?e$nAa)Tqr=0 zf~hY6%wlK1hOg#^-D!|Ufxmkw)2r?vFyIH%ybt8M*BgbCiv|uj6RJICz3S2d3u`IVeLxJ%l6R1()cx4*iAp|kpL?!nmz8J`5GwrLeY32M zl=K&Od`JIHw~Nx~$Hk*u0L-ts-Rbc$pdZAegrz@2LI!kbq9yIgzYD1;Ac&Aw+Td@6 zW+RKRZgmHjQsYnUPfFD-KKiYW3BaGG3tS!mf&(@8jaH#KqOq~`yiV&<9gm+LW3Id&WGC_;K zc6XCF$xzrX<3f-OQaga;U{XflP0#xuOmPA-O4(c-aUQ0rFop8ZFq8jhAMy*X7Q1dQ z`w|X@!L*Q1KX)KNNj_z8k^{4r!=c~RX6KN#n0~w8Xyb$+VX-pq-@hdSlDAs@?QjQz zfDGawf8+aP#O`&hN`89WJMt+xHCY58w=mj8q7G3D4$*%-g7lVyf8L0M(Lfhw77Fu$ z>1@6mV$=YLlM^S5zSdL1PNQ8mo!xY%#cRq$6;z@)Cx+mJ`f+>BP6|N~C(?_crT&)j zgPV?Qy!^kFNrzB~M_QC)2%JE+Qg|qTB#;v&zm#mjN^@o*^v0 zogM-vN-jMYH_rXA@um`)VwDQ0YlygpWprOKVi>^^8?2j&dBj%&>{KHGj;R3>Jp^Wd zDHBh8W64KF{7s~!2nExV!k&^X;mj!^u!(ybi1H_dqHyHNRbET*K8r57Gecp2gcgjj zp%_+TX#r3xSm1B$!OcxVB+RMFho>CD1B`5`kG03`OD=V~0zo;7<{9QAi|7gv1mZzB zr%sjdu>MmXMg`(nj6$6UuYmEq00TU%6C@PDFB)Om-<1*gD{J5$Zh_bu`5d~MJc2F= zI29*`puXG4&hQMyRbzToE}`-O!s@XAq(8YC;;SH!W1Ncw*)wVXVflyy_zmO`&o^bzkqwJ#i~5l6n$V6dIr{&nD> z$7Io&(S107JL))j8jEfLJd3{Cqg?-V`g78R2gt35l)FUrPyA{$dbXdHhS8=9_U-508b;0wd0n>#Mp#g|E@?pcJEep+|nS` zVs6HQJ}!Ki{U`@AivMO1`ksKL#id*6+2l`gF*7iT`2E`3<3E4?x%3`^WMsbIOmj`|Qy$mzEM2>bhO7Cy((-#a&B*MQ-fdxH$>5%;R$L0AmOY$Rb zUR(5>({E)L%6bBu0P2!DN~Ue3o1IsH@}|o3>$cJK=nY{cheSP0`(osw6C?=vw6Oj= zp%BL0lllCp2f<_rY!1o-^qtZV_B_!odc4cnhWH-AAxL`+@Sbor2rg=H{Cq%Rz8*xC zUUpu7kgzD59VzHA2la6vULbK(*Q%vSI{5)MEd})pfW(J649C8HA*qeh?yIH?3)2}G zO$}sb(9pw#6tV-Ji+&V2yoyms;<-drpXP2O{}Ru^EdrjatBn@zINM*f@K3&xk0rv7 z8k?L&?22bIQHvmA%n#c&XI?8^Mc1SgXCq1oj>$kMG9ocG4Jsb2%XD0Xqj8ZcguJ3f z(@*yIVVp3c96sFXUq7C5HMD;aUf|UJr+V`F@1*GvsgKF=OHTX7u8_hoBw6CcMZnSb z4vp+rV57Kyf|;jqYeb6S?I6u0QP|#*6#0ndNc2pTWZ`}#KZaEI_~Dy=Uz&R*FJ9c! zWQ?a`;AVYesIlWhHc5Zj0}F#eJC1qtvAXJRo_O=g-cEPTR(I^Zj~sMigbbt|)XTii zzpD!#{?t#+ot<8HJidBlXffu+z&c?qI1LUMa?%qZOSy3vJtEJK!v&Wd!7U$+B7W!D zWsA(1lr&D~@`yM0=zM)7gJ)Hx54aw!DJ6>_3Lfq5cY1G&Z+G83rFsioOqXlSGGXz~ zQmlcD&(U^9iZ57szlOd*1tlDdO77aRsG{5dWs3PY0<2CgE8M%U z$q;}upm41HM4$+hm>n>hsBwaDZp=U##oiF(cY%AK2J3d)xMBvHcNL$-JZUAKlAL(7 z?K2{P2djmm#M>hOh_vj)%%Anfr$W2Qgo%Yoc^Q*JfpQXW4zf=B`a&Y$euqYpAgyp~>*YUw`fOxDTdPAwz%I<`A202xZS<&rQ zk`4Y8gheaQplgt4%qB!Tjt-8Kn;k({=^;W=>Ltb2&nwpVe*6D1kzs|PA4=}8k>zoi zMqXx@NB%0I28YetTnengk}l}|o6bH+W(4pkDPdQZ;+FerrN0VfP?DeRc`eSVh)%*AJ{?u=4BK{+l6Fsj32?mqNao}sp!*fn|t@H$`dB4Le@K< zJ_q9OE|9_!B$qqjxWI8Q01x=50C1uZbsDh4kKo`x-@n$_AnET)pa7czE&n#R|7;LOFanJEu(Hi(*aZKti)0395mQL| z^*7}H&$c1JU?ya*{+rbG52*`b&^?9a@GsW+PXin(aEvL4QLOX%|CL%9AQ?XA)L(#^ z)BnC`0HVHDpdZ1%-r^sEXH!Li_s~o-VIs>T*hU~x^v_GZ_27kU#)A|6fnpD6paDSAa zpqn}v>`X!0ZW4KW2gUN8tCWSkNiYL2Py=VihZM>MAfQwsyVvw#h3<{^nw`P?&*Rhf z(W{>&hvV)qFL4vG4V3QP z4!B#<%$*)kRt_{)Yh>CiOyIr8VwJi9w+;uCNBO!pQ*k9+eg6JcriE*v)>HMg>nAZ- zd+SxqibbzKP@v!2_Mqol#_tO{SWr8?(hau^Y6A+MHDd&!HK>nujl=pHD2Me7R34>PuG$u-t_{1Do$?bC?9Fc&Glm0B@mXjaz#BF=|=0?VFX~jR*efZ-k0e zI}XEV&#`T!G)c@rfKBw9Ob}#bDy>(LVH?U(S;w5XwTd*U`;oh)cejDFv$*jrZSE8U z*pE4TmUS3YXi7lvNCD0!6IL=Q9GGGPkVGAlu9hd~+_5W)1f`UC2d521UT?0RRkOu&nGL%ZFEV>Y5Fva89Y)!Pp04&J7F3pa zG+gS3ecjp>=-q#?eCK4h5CInQL{&u>}V{QT7xQ6N?00FSMp1-p!q z`xAj*pjxZdwzf4w+nrehgak{^5p!xc4g&T1cFWJcZ|Frah@hkxr5Uv5B~95G%clab z*k3^|tvylI4z<3_o*6s}yC4ndFY|RxqfBWMNmLUXvUJ@1BBO81XBFMww}h5a9R^;~ zag{Q*5?@rQf6i(>u}m02^_(1-{Ms4va^aqexar&Rmh0JAfV-HO@U~8BQAH$a6EAfO zu$yhf!USrg_rs~w!lU7wa^$^N;7E%TupznOrW^9HRJW8P^Y9Vu_<$y`@WD)&-0Hp@$O9zjAO|KoZ)*^;+WC4o$fOoWi0Y)I_i-LK9G$?SY_eaO7j) zh&fh=+bYy*(nHugbbdBH;c@>d&^HHaF`tm%+5B&XJ2W4 zxPBTNFmFG#dzTrXmn1*_DHD3UAyCfbFh8_&@x#hPltZXTN(+Gos9=84rmjPiSY z@U&PbCYePyun3urT|Z5WG**EvhNwTWQ1p6bj5UK^nWI0?`@NeZ@u$TEdb6;P-apKE zvb?XtHE2~Y^UsFnitL(yRc=p&S+%!3pyIiW5~T=m$y)6`)KYo;j#n}-YzSF!%WNGT zd<_6TpJ(F&7n(Iw%%;^n5|<45#v)ql;(^*=-bDM$OH+D^^~_kX-gz@>~r? z)$QTs*yC&Y?*Z4l&?omJ@=awQv+pC~2J+g@FPuO54}S4H7}?%JTnNlb%sRUsf}1c8?4v9*n) zTWnlJu9U08LBfK^Z!@1^#qd#`(^kTE$ zSon4&HS3KdYL==Xr0h@byEJIY$x;}*23a+MK8cZ6>BaQ3Xc*8SD2$sWj1tJFA3e?G z>z9+Yh)63PJ2y|MU#4zH?^2}5LYiuRK8Zm&T0tlM1zG)0q4%gA<1&C!3kd zbG>M?Is7(v_~8K=dZ{C9v|lSenzrr^EME;g$9o{F@N{dTCD>djDiI{F=YLn+r+ziH z@|3_k)Jx!!Kl;L4yN}TI{M>pddwO3x2&Fb6(nwj9bazP#i6bE$W`RlGfjAzjgNE?V z`GMOmKR9gSnNcyziR&GQ-AsOouUoehtu~_iH(9H{eJZfj{0j#(imXfDmh{olmMs*h zg|M5Iy`=4)I%2E(+VJ6qzK-I4Y1Gc8(jvUd|C*Nsa@{dxcjS@F@aI<7l>Y<0)VK!~ zu^CT*ZZY@vSH*MAncDSD$GlIi{+v$Rr}>h$J;Rsbnk(#-1bs&QCI&s|>J~O*_Mf%h zQ2+9ru^<+x5nDI(Heh(Y%0YP}W7ydV8ZEA@GjI0X&_!UCd=hcK zo%!MPRIf8ALoZK$d=M=|4h@F6{K?_zvVXd$yDCtC5s=zgZ%1*Af7p~uGWtah^yMH^ z<(h|~mvMA#I&QZ0f0><}7I1VP-#ok2w@xxRJOBr26+E_iJy(*hE4S1<#6bzg25(r~08KS?td|)?Q5wM7^dfP=guZr4x zKbsQ&`gmjIW0=gdnS-Ntt<1w9sR~>Q5&RP30JZIh0b?Qq{^0@q!<3Py#35EADHY>~ zN~0(_%;+v^%inEr(eWxd(kgIf&?R;gUfCjYvbCuww+F!!Cn3Vn;VK(M&mdXxyQ(>%Czo80-T{Ak_K?uIa4v z#0tLzUr*QL?Xk5?X`^C~`R_anA8elQv5eR~XB?c&Id6yFV(ba3V{gn#N{h}rcw!z@ z-~dbR;E)SPTmvU&_Hc~2$?(KF+?{CA{Y_29W>|qx&8p-4UBJBKJJ(TCRpA@C965tL z@$fe_Yi4D&mKtQNvaP!69`noPqwr;htuNAoxJ(~zHVz9d$l<}kET~PDPkJ$Mx|WKA z$$mkOYNS^YtMWCS5}9xt7ZZTOMhq#qa*F=8C*0B@gRJAkYxVmx3uef3v$tn;PSeMm zTjcHYJ%)zz4;*Exp;IcIAg>w{nr(zD@4ril0JfBzq+A6kz+_0^r)J1jj-1BBMyksm zJn!zSM)0dfD29>+SNp47wi%$K6_ncZfmgRFgV$)AW2q!9pUl~E<^{&YDd^AJa4&{G zn)eqK3z;^Hm(0$}P9mG5)$7Eovs_tMt@TLxxBQZi>(R)|%{no!&2=2>5;BT;3YmJ*pPp!h^IVoD zhFY}Q;Q5n?$X3Q+r*9JC?+jZ32KGElo76W9Y%{6>Fzli{@y9i%mG!ITyPtD$IW~uM zibLoUo=;XQTjX znRz%r#9DFQd_s3sEmz8<47RQ z@=v+VB}64zi9<}43}oAysOG*JDsw9qY6wAm4xWHDr`Ve{%9mkgl~=S$=bK3-kvCI22W@+1cjIp&nG?km2LE&Kn zYt?0|xg+h0Of0G8!~q9h-U+rkZ_}5G35L@4eM?CgI?>DsRbYCHgb-C3^?85G5O5L< zm>N9IGwF(4?#P)QuB(>yE?m)>=8Mhn`?jk$%-^)EQ&P?mRm|W8Ru=BE#-AmZjnCA- zIwtZMic4&A>^T3LK)n~zj21)0xcqxuq<=(lAkAEMPH)VT^67gD1v`O+ic7tt*ovm_ zia-RYbeKn5d%tcb=0)R%C@N9XABk`*$B*Il<&8Tk(=J0U5DK$2SGQXlzLA*|e|KiV z=GJDD6K?xGQJ|uJ#W8`2650U66S>HMTl3=;jZVB6$B(OWlU!_>q^TF{NT$YE34K zCo9TOp|;HLydAC-%$exc!3-%^g=~wd4eItq7TNJ*MkwI~1m`Pi@(Qb=GS0R@UG)>X z`{e@RZ6j|;B~p|&Xsw9h6?dJM&i=6clilxyAJ;R?b!-nc%%3}sMulw?yvMhwJ-2++}T&8-;oMn zF6O+H?wHm&5Jcw3Ly7W1fu~rh9V98@!I@h@{(3g^ z^!iT3Q$@p>G}u@nO%hyj zqJsG5BO?FQL9Rn6`!d%i&F|0l+;}gtShTr!hA_a7M{r%AmgYgatgxll-QRZb^IH6z zWfWdSo?uZ`lU-NukXwHia*RVZ$51)eaGv+sJNH7IVUKX~5-S@D9b2&hhH>N=jGyb%KSQ5u z?O#1`NIJyzNL?kMdJ?Nf`l-*1UX>Qf@C%+}CMQgnUSzk1GABo!5}n)b=>10dld4i+ zlvd4Yc)Vn1-|lUXt{H#TXuo8SYm^y9~(VbHr%rsV#Z}p-NX?HhiX*%yvdj3u; znHy8I+xOYtxKujkn2U=3)@C{+`2DlOjNtV-<y`Rw=FUv@&Y1h85!n3Oi-w;<8RVVh)3M^QnTx!-OXMfU0BYhIkBSc$;IASg8 zd)jKk!5R9-9s8&yNU@?pg*Wj!%6?_w1yw!vX@2@;Zqy5ZaOwA(By3fyY6Zpw?)$vb zey!SE_;%~ztFe59y?As2mMtcIffEJ8Y8WgkC8aJD{pVjo~rY> zn@Z#O%a>!Z5j$$3AGC26`*$^?Z8FfYX_W7Orm#Ve)F;+1%joPtzxTsomwGMKTjpmf(`d^yv2Q|+x0zA?` zp?E^~o7O4u*0a8oOCI{)AM&@oD`VNzz4=he8g=7P<=PmDq(!b5Kk=uJE|aipn2%zbYNx{-XF(H+I5=-b5}XHRp?Y3U=F4+MiS8R$1@n5(c)`goaiy zim!a8Xcpje=@-!Vq+zLWz&KrHIs(u*0*(m$_fBk5Vnvaa3|)qWPix0cAbruBxD8mP-g zyFAqjx?8dOZEcEp!mUN~qGPWTg%4JvR>nG9RPD|}nVeD?IpK?UfjigbyM1*;Xe8CU zZnB93M}0*14wl(J;_>;E_rrK0&&$4ZDb1CjAgpnpih%UtmE9~ying^Et@wPEJp!R# zvuVCaU8!)4**`tKT^%XF^+;n}oGx`ytA5!UIO@FE9LI1eS~TU$H#)m&9@I1Q%*=0d z)NmUKx)-3FY|C;OucL)su(;oQGweJqEv26m|IjPgpG)z@T=S(hkvRca zb^yq2{!Uq`ohpyqS8(u`9fP&G7|z|P^v}2h7xzM}2}{%X7@_i+-!3pTeMo!bp`6Ks zqBUdQFi7zCBEbd$z#yQd#~z4BH5P-i(AQFKjW@RrsNel;Rb1oHl(pL9;CiBr;{2# z{_gLu|Mh(m3m2~}BvM82fAbMg03=!RiH-du2>-9J{372WYIKOUsf z4{0ncPPelg#nD8tUqpXBi-i&#I&PD+jHU~y@&A=W`Tu_F&*J-l^7~d_9v!c>8V>l# MNGOVzi|YCP4~CMj0RR91 diff --git a/img/goai_logo.png b/img/goai_logo.png deleted file mode 100644 index afcd5ed99d3eae7de65ecffd8a9a7030f923e78f..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 30695 zcmeGEc{r8p8$OIbW|>PQLm^ElN+Oauks&hALn4wPvkXxvn~IQ(nOUYKQ&^@9)yj~u z$f6L7STe0LzSpDA-oNkhd*A=wf8UNA+uqxHp8LM;YdEj-I&Y7zX{#|aa5EqX!mNHp zMGrx!u?RwuLr)9;g{J975d340+ZB_$2*S7z{i5g{4cbHyeneeGQQtS^=V-v~8`FV5 z{^S|g8(_cOusTwC`&q14mg{JY4UI{Ma(nWx&k|gii7)2IzKC9ZD5M>#sdBts==$q} z@$9V8MX#^%I5&koj3}f%tIb(9m4F>Svv0&!r-Jt94Ht&4TFQ<`pOvTs{KO8GFHEOw zWwDt3y!p#Ck!;>@IOA0|tLmkNT;sf;>A-ZUAGIa0QuL4f+!?;V-yzaW%A$Y2{e9^F z`|kg}4FucEi=Nw?T+an>K#-2TcP)lIX)}(T{jdp*tlt~Yd1)#<`(a&>tBtqq zover5dk_mdlR*ykAt5dQ#?|`lq;KE8soK3l#$TNMt{d=ZBlca^iLZ2piMiyGp-2dec_7 zLxarF;>#}&a9L`RmIV&SdAj)ezIb&JvBFRGWyKC^O*@7i_fux@v^b-NYrNF6yC!*4wC7vTDv$LDmp~N8#alz&3G9_3PE3U*zKxp~ zi=5i*N$J%Q<^;F71a319*B#p_&JaW61TTZ=0VB07KN&%y3BFt}wj7?f%=efUDW=D> zIg7^W==>6AM`9ubIBUD+3Sailg<;BAiel^>IEL6|8=t+;^Px3cj{V_q{Kvw=1Z9Yv zD4yGBg1DW=@k;I7NunygyaaRRqv>s(vxsB#&_nye^O@}BPG1?RkeEmTPMJGHJxXB7Tp=_%{#=E>RuunIccRmx?(n{UuJ{j+ z*h))Wym&9GGyO0U(>kZYJRxUqsQ1B*7NKX-KCWWsdz^NGsW?q~G3?VxF%0(M46V*e zh+~`#uyT$8~ z2?{pBiAUv*5R)`7bt*8@pD88JzO3EpvDWLY#x)dywtkqbl$d&XzxFnx77QR z?yVUZcw-x(&U5{4Tw(%jDU5Ty-8Mrz5UTs zCtk=$rN>sxxCB#&7@8Z{VJ!4j?5*jLG%w*!5Th8{`+_{%v$GH2Y&dl& zDX2v8wT11Av$tqAqDLaL9Zum*GsGAb7|V01JiiFB?xz^H_Ow+sK9;Us;C{EsTk+@l zf_AO`Q|bi}?0#?kLV8|{&OPRq^Rfwe>Shw%P&9qCvnP>eJimUU{{$f`SE!Sx3qcy% zG>YhlE0<-%s}Q)9i#=cRzi$q@E(QRffkMXHk8^voWl0rlvDPl=K`R-7S$ z*_w3XqHQ=%zeDkkr~w;a?g}IGknOUvH3> zU_bTQmfkv|I{vh#Onn@hx^Y@_}2IiF`z_NsT^SR&G_mxqM; z;f#h_yC=6aubNhfia5*ABUaYFeNWFqh1`r3Fnk9;H`?)3G7&oe_#Hn-iFn&Ya_pPbTHF z=~c^l&%KwiDT&6T8@kcZH%V zNkai%%AR|xoFGb1;cG8jak8Cca#`?98m2{R*p-Wrk;--9!sF)6>E4F+aD)_|(L(|8 z23h3UJ#XDxCQm|C`Pwl)NU8RU8Kn;RwamJ^v};8Tq1|Ba_;^V^>X~z;cUdulyr+318Y!hI*B zr4=Fs;k6p{3<-I$H`owq(h!wjCj1n8Ou^Y=Z`yki|9vktd<@c%W$%GD1mArc^C>a@ zwW#w~PcB4lVwBNbGW7`m<$88r3jSU$+dt3Af~zjrK(=q5CcU9j`84;4l|z;rH51Zf zoUHpoZwz~lJ{%YrNI%a<-Sq5ZWWeYI6bb~%J=41MO7l``+)5FVS0N(004o@!$a(Kp zK)}XvBSS;|(}`MY&3g?n1{_gs&opwDvT|P|();|=VjQ6Iz4l;M$c6_z)4U}267s#O zjd-S=N^iBvAy#=u@7`3G5mPIxq1o9&r^=U)jHM|@stiOq4T=q?YD6y4!hu3VSj>t2lhWE`<;;^E z?)t*<@ON!1;McbHg*LXu^dcfcA^Bd|8_{iN!_#?l)L5P8r4fIm$XN?vt&1-&kG1?X z$I6zLP(fS~{vuW#;h^YIy6JT!K9w92Kp zxcHWrm)F>Uf@69r0i%)ATR|4)jJjoQt?>Sao}8RqMDgN_go?7UHC-fig4rs++mR20 zMO5}Fiz?Sk5m9gJVG0wih(}{PI=!xAxHk2Qkmt7+8fI2jOr8S?f3wnb0Jd>nN{Wb| zMZof>6Ip-aX6NRvpH0%YvKn*0v>M4R@{4l@PgyJMETeEGyI$J@iQFGuBL; zJpQA&#N?V?rIwi?u6@EiIwoco45nqYB^{Hc^eQ7ho@2e-^crRW^&iZ@%R78OiNsIi z;~D!NsffJA)M-IM7)orLrOlAg%cXk2WX3?D!lGPs+BSqH<=tz`U|u&NOVU6+x4f&Y z$GSmqaAQ}xeCel8pK{HZoGV`r4h~i{G-z85o`JXN6=cI{c+Odia$biYjD4T&9}pnR zVP_i6F8n%mol;@HpV=5!ZN(55>jt8t)v(o(D?_?xykAX+a4xtEk8hqy<(?FJ^MDMH!XI| zQ&L};-JU)$T&d)^Sz=|>8f{v)zrW@e&DvD)%cN)bcNVgPZW47&2@!ZE$;Y34OybWWGS1Sd z(sY#E+6+;*N+$bwz~?kI{M*Wlbkskvu(GO zxq5bE&Q=6tai{PXrJUF;=#fr8brq~1Fq3UO#(epgEUoe9M0=A7ry{r2`Cw&m;$isX zb0tpmnAdFo$+n4Kr;khE-FXd!;{)qjY_Tua*Vp@phK7oZ!#WJE***2@^km~d0coZF z-sZB5(4m>Nx90rx&}bRc9y*+iwplFS=o)B1K`6q&_rUzo`P=K`>1%+Y<-dR)?Ni_@zz;*<=kiWulf_+3zzlN zj?HsBb^2n+C57(n`vypdX^kBrs~(iR(@VAN#?Ki=Tp{*xb948$w}%%F%<>KKbv&~- zfiwhO$Kql&*tO`?I4}muh?|Rxez>p#@3dksSK2P_{K@T5Q>pZ`GT9QqM4vJi^pTJ( z+sFZmqQ)1>)2@J_cGKd5d!9OC(E5$D;-0nAryRY@GmAjb5YuJySipLsjf?EU;aWVzRh2(5Z z20Mf-vcwpDQ+>qxD%tIh+);_XZ!_q|g3RPAj7hHR0Q+-2ozC`kL zZ2a@pD&E}v3#3!D#x-G(HMODz6i+a*usFH7U4x1Okx7vFT40DVH8IgWn*`wB&AOWF z&|tt(^{uU~paq=Fqr9V?uzF~;QY83j7xx(2cpU)7}uA$r@jB9IX6JwGnd z1b-K5`B8##pLBgJpc>Gr{c);3`T6kyE^3-j_}G#~90k5SR}^phbfrje?R7KUUq@Tttv#!I3a%3kc(1xUw`of zS%U*oT67m)opJS{M;`1G!49MO?m7|^Kq{Y$Enl(l(m{-F;)>cil{B6n`8bFBApAAZ zRSg)bs_N?GHT(I`Uu+#~#pX`ni~<7#A*Zr#3^r%g<^VL1(xju1>8v~MH|QC|UW0oK zeYI8mit1~3?Vmp&!96FkAw7^_P4_PzsL}Sd>}7f@a<_hJ!{kqZBqLQ(-mdYhCBwGB;z%X4mefj zQWZf8WoBU!es)E?X?QF<+tn-~-Y!a@<=w~!&Q^-SuJ>Q+ES#6z75Dbfryy?p3*u(V ztcDXbJLoEnzyp40MJT=n!?eTKKiNL-tWsH+$^<`_7Xjq%#9-Tm@^z}}3eofKP9H+g z`_)r8Z!@#L^>u%>jq=+gh6Vb1IeNEhJI`=N6~NB{zQN&eeE7X*`(imy;lD>8YdMxb z)+hDyYx-n=%@A|uEzP#H#F6jQKOPC~+hbAhs|z%bRw>liQ>RYda&~qeJ6#;-W??ktVFGIy6ELx| zN{5oZzOgZ~=5^$TO0NW)c~0RIT)J*M+DOL9 zaN~|imm~GJWg;j5@cDz69ijXlv@Joh=Ezj1=c{_CByx*7;e6HF8eA@Sj8fTPBeh%+G*{2v{M?YIQVg{b|M)}gQ zcLN7a`Y4c$R4!Y-u<8_PXX)@gvSE=aM1DusQjxjGG6A9Aw9Yg=WjxwQ4 zqKB-ktgKafAWF5ksOaZNBJ?G~@sA(3eQ$VMLLgWI_(c5-Yg|ZD_u)*ml`n=3k}p+AgG2P3Jm_5-2d^oW^qkPJHKtSCtF3jPv|!t%+?5F zX}9)ASOh)Evr<*k@JUmMv1l&j(z(q)%UXaTN?vTt80VJH~sznpI;F= zCoTPM@k0uJmre=?4 zOJ_i>bk3TKSp6m0Gl@=p)1glQDKN&S>8{XjHUJ#SLl6C9l7!Yqp|{4j#KgqE4O;S= zWt;68G}@CidjqBp>*Pc0nIpz{>6{&G-9o-S@Q(Hm9sf%O@1bw8hrXqPO1bFB;$z{k z&6TMd@AcC1+^`JU`grf_PW#E#q2Ox1I3p=QrH^gZGDz^;+3#&AIJh&MXi%WP20p#p zPGplmG(v+0Hr{OEiJ@Mt9Dhe96Us#1(edPH(jg(<){$Q4Bp4R<5hsKrw=q{H(2J~S)m$CCNwR4Z( z88$O|QPJXZ6}w{Xjo~>mO^AYBq!&FYVyiFruIKwxpMp93EOiFsIyQ&Eahq6!7Q=js zOG++T6zFTJH3Q1;t2pi~HBv<$HKpJ2;B_e17sfft&}HK5bXc55FE7;Wv_Q7Y< zHW);60EcqXTUC3WZ&2yPn?i;BM1IAX{_a9Vl%T_f?g2X7wb2;b;4&Ap96e%FaabL| z*;8U-kbf|Y#Gc1lDJS-Lssm-fUOUiI62{*sRWmz~jVM^3X_X17?vQ6NyJ=}?N5e_8 zP~&3sZ9_3kl-77}ThrS{z<+PdA%PD2H2os4^WY4F?Yl{&6x*7LifoTJUy1qB_f3nv z#~;NiRQtYtdaLh(_6eWf;7@h(wF3<&EF?3>6BemLgu;y@-UGW#!&FvQYVnZNNoxWb zJm!AH&d;K*f}uIB(w&|_11paoD_&bha@DIQxU#L|!=8iWdnx}`ZefwYdo5Kiuqy*$ zj;g>GvCxZ5x4a|`o*LQhp0J^K>u^LhgGheB<`Za5St`5vQRp)ww`lDZd2rQ#Yi%~= zCRL1hRaEH#YJGVkW*? znskg&+&h-9Tr}rWL+&Ix#xn+;OL(7mJy#iNQa{r=jM6YHB3CrdG_~mZh84k7;F3I` z%uaob0i*`;#TmiBlmoOp+et=W%NumJ3%C2%D=fWs?u>OW8x0&-J-T5MIJTwTH}_vZ zd)v&L_jVq+P+xeyuAD`@jVJV<7hUd3r%E%I7SdfS1fou&X=^;&R#bgVudb}M)})<%v%mk?mcctthN%+HHc38b_*{&8Z8yU zruWEUH$+HG?Ze6=26i<@~-tJCRxeAbg$VJcB1by7Np;bCY(LV`ZHeaD`7 zr%Hco(tz`-(aP6Bht>G+2V&%!V!3f9R^FHL5>pOT4BV*xXSZ}`iD9T>0y2fHL&;@Q zkwD}}TNpylj*7igHMe~5sU$F)+Mxo27J@wwSA}#e5Kl&`Cu?6$uEDCc0!oGNm{^?* zQtWmKtRlV->sSpW%dW@D<un+8+oSwVzXO3c=y%x*cAt3LHNNHdF2`_6{bThG)ZIt&Pn& z@o8y`f701IK&CB9*8&TJ5-(5@8d0G2 z3ECx~wLV|+SBXeHvFEq`H6t7KXm%P|fp-{eE;P#d=wa|OVHcjC<;zfJhy`+&0@+9P zBrAWwshWseLAZeig6V#-=CFz0e_S8 zx@*pX6^FsRk=_6=y)myH)AVZjEK+Uoz>5DH^T7`v7)I25E`0bP(=-j7Nd^Bh*#M^fSjyi7VN1P zP`gHow_0ddXh)x_!Lht+xwy6AQd;m0fN}Des&|hu26|B^LY27Mp)PPa6oSrfanpbaO}vcp zydAHhy6PV|W=k5HEuCqMx9oYm6?XdgT#=-Yx|O$l-i96P+^@8Br{ToPhK7>u?I09X zLrWJPo5qVXqL*;Z-`k&-H%F`S^sbZVAa!M~(9v52n0rcJn>4aLc~7}$y3UP~5%C=D zYO?e4k`%{S%&xA_4_xS^HBRoTZP42cO&eLiJwt0;hnw9_ZMf0R_c^h$ygakBv$KgQ zYzC*K0&^rs#gY-Lo`A@&tED5CVP@$66+NOW*%8K;_ie7FKmQrA^%GHYaTxzL9lDrG zq4vi7QW<$xr&vb>(*6Ahqecd*=pn^!{cH#{`LE3U^r+1m0aC-!>;xMvLELT_`q5hZ z#aDNpR`1uSPejJ=K09Pv;4W__wZ65zu9V*r#FQA85Mc2BF8gt|;x1|DK zZ%P`Bx~(YkbtEwOzxPvwd|vHpW%#!pBx&b=4a}SjGE=_Kwr5HFi1RFC0Qq}&^7^{V zQy8aZp_Md<%l4s~Vsr2k71HrN@4gyPPFQ1ooTaOO;Fs3Fvg;Vrp4ziQm?z?*EN1Iu1n1I`4?|%7`>)xCXbS4me(9A#_Ri6 z{Qt7pP&8^UDGZyKbPOu0^!yVVke(tDsPogszTW<pn~w!td!#dua3o zy}SmN3}Iq$XF2W}N*$#(6#d_8#jHlIMvlk$2O0iJGzT_`kz1j*n~MJrV<<>4VO#DS z7$aqPZABQOAChe&(5+f-DbLqIM4T;$++CRy+VJCI9tvp^7y?RUcgak^9-g;1W^5qd zZKj#UcEfKvu%9(R*=Z?ERQj_XeN&>~g;s_yhDwGu6bkIAkhWJW(RO`ED%IeR>lrWre9Xa2K#+sneySI#EU)cj#U);^TTQTkdG!EL) zEDSo!ga}-RE$s!9j3csCPnw$*(b)OryZJf&ZPh&0Rh-Bjld%)Q24Y4H1Sj*KOBUqP z)xZte+TLY$|GPge{tuE`BuTR{Z~Axl4{yUktP9)I`d4lCZ_jAYCJd=g9Ul^cPC|w8 zrch@xRJl4;kE01KmGR;ZG9YtOgh=pK{H^O>$|-DR1Si2%CX-$GxPf@j0IfyfP3O0` z@Z}Zd%YRFTcJEri&`a(ho_&mb|0Bvf(!rdHVymed-pn z0(THLdbNYTygB6A%5{vMhU_+G{pORwexhV}C03xS<~l=>P(|5gB)pea{BIfVMC zDnV|EJ@;(8$u@YdNRzPIZIHwrS$^%p@A}(ZHc?5jn#0 zb^jA1-WMt3dgY9Zj2h!+!|+KFqDh~Csd$z(>jjDj zSN0p(Y9+Ds919A@(uOph;Ed8TTKM@6QpMLF#}8Wn(KO(|P1a|KUAU{Y0loSgT;OvG zPPYtx9~-&DA^y=kdU1h0&3h+FtcU=6fq>z$^ajCOt)4OGf6X-nBlOW|l73U%JuA{C z))O#g=e!Wx@bjLw*afk3^LR?Pe16uG9=TM!?IW`V7V0ppvj*{(~E|0ARsw3jErBFCfQ-$$(mD%fWqMh3JbMph1Z$QKAi=>)|Rvtwt zs5^Ajxg~ffvKbf%oVX^Y(Y3vkE_sNO;90a>%9V^aTp_L1TRmZ;4Jj*kj(O~CA75N( z9&z%KRZR@W^i7HXmmee==_D9zo~qef5nxNY@oU!`dgr)k5j#PNRy|)s8weHz(KQeo zFfbRv7#1eS>~ks6Fa2B?KkMNVZq%gi%j$a^R8@*-&xnFyREQw{Sio?xES?Qv^58S# zz!c68#QqcxG#R})#A|R>vdV|D7={l$p~)1NCTGI;&)Dhb;V4cZTzhi_2A>lV7dMe6 zvCjybaLIgO#itp(t>m@oH&r7uxWeHTONvuV;cU=5SMZmO>Jd2Kr4V^f*-COEcA4i| zxuy6rn2jIOAeZchcwI^p^k0c#j(Gc*I!(z0Hz;hN^|84a5LCt(JDoTEw7UHgQMXS3 zQMnF%CxP{g`19IsbhC~nRegyDxnh`8=I={-186FeeXMvV%#A*NexAu(yNzA@fHAUWPBYRhD z%%7qk{71mht)92S?g0gVVM=Y^3|?5@Sk;pWnuk>eoRI<@8lcV5WnWo=(&@ZfYyUq7 z`2f}Y6MMXXA=lb%u@|sZ=|gq4&HQb*vpoyLx-oDs84_>p+^|Sab>B}Oa1cAzkYRR(>GHOww7rAKI6m$>T zTAb(0Fo0f+35)Pk|1^b8?Y6;@FNssbB~DWyhJZJg0ISp!igIjHnh`TkGgzn7-c9yR zK%I|_>0!mz+dT+du!sH?ncww3F|}MVyqEca2nf@tBYZ;C2~TYRiEvJ8(#`C5jhb8r zr#Rs(gM`e`Z?RqdCY|K=364c!<4w%S4Bk{sN_&#=WZf{A740g!&UkuC!^iJuI2;9$D4rj1K1I0lG zI2W^HXxZJh&!Sx<2Vz44!~X29w1__sjAxMf=;ZKeoKZ+HkR*SDB!$Rk!MRc#2P!as z)bGA=J`^zQp?Nd87wI|X-T4Xf8b|ty{XC$74_;Q26mG{|U+j7Vq@|n&{#%2H>=3(xG-9fzKj148nF5lr_L zgD7c`rb8tYg|-Hee0(a!@Ff=t6P)0K$UUGlo?P{Bf%7+i;*^R~{`H}a>_=v(^pQ)u zE7hH4K^#96$bvnUaUhWQ!5}$A90ioH{L$*KJOY# z2VaX+RT*6^yfZXBtTQE3GHahq5aCw<+A}{J0~Q>I(TjQi;qQc}`r>?Sb6Am&V-PD2kQa_jtDKcD?Ne^#&mgoy=*nOUeZtmyo85^c4~1LMFmm1tJJ*KlPH%P z)^r$EwuIss=4z|g>C+(n`l(k#6rn0RN{J0m8OBE6Vhk;lHAUfjLzT-U55G-e@?VEhdgexO|Npbci6H%J{>NgIh5mb!^@~oVOF5DZ;KKX*OElKhM8H z(R0jZOq(8YjBDH)X_TwOVt>_VADD76dGwbm`|{dhU~J6wJU<(%m(dW~2Y+XYZvyAm zDzZUQr}rbaS#8(wGow1GoGbF_#HXGdREzh^Sh2n`dk~l2zP{xy0EVmy8aZ-9w6wII zFpUP`$wE*Za|R6zouSRtI2QM~V04;&s5c{JCr_F0JKq5;;3|X9+pErOhGthl$ z&M_1D`oh2^ymvn_6>$?Xg>H}0L9Y}e-{4X3I*V1>5~L^g8Y|oxNY(upbn=f>k!l9} z`m%u7Aflu^FYmeMZ1lM>54v%-0I+_WtxsCg_`idZw4DVuEy6Ya8ReoB5L`Wz;LA{j z1WOi14qHdkAktW%NI+xb;^H#@h7;!UKRzBzl?yNg3()pbG=GBkG&eWb1(e>w3^>79 zsc^uFwzs*XMELW=u_2tnW$6b-_jbymf};}+{@;MH5YzAf3aLI!Ev3!D;mZVXUko%+ zV25%HVi+7QfHk`6f3slK68xjs&qn~fMM0jU*6gXaj6tIic(l-HANK7If;)2GdZcP_ zXsEKbcG$D>9f);+g}Uz0_vua_9i1njg(;3YAO>^Hv1>|AMPNo1!B?ogzd^-)&&d`P|^nK7Fi!{hwe0(=}SFjnQ4G| z(f#`rNB2F-9V0^nm!tQHAi^uWJ5am@2!w%9D9CHoO<-m zcI8hwIY)PW2=qCgFm{Xn=bxU$hu7wIhcCDN-`#f^0a*d?)wlUtb!h@Q4B=@ROOF>z z!AmpIQ4dJeLae{@h3(M1nL~i%A!L{C>dz{PoCV-6pnW@$q{)4S&Qts1A)~WNU#rQ5 z=BF|G!dNdhwI8@Y>E~0kR29eeIhg*p%ongshII{5`=whSMkyXB{Ux z_5zE4Ct)xQds%X}NM9IT<*1_(HSR9yqRE#nF1ySDhTv4xnqs1{uW?@V&9xQD6Vk(l4ZOb;)q=D^<%gLCp}2oj{8=xvPG&P~-M-!5NG5Vkx^<6Qy8|hidNB*g2Y19w~qo$Y4Wjs zH2(k*%rHOgeCw7vNKpa|3qd1=TJh(X*DLnTQY>>g;1R*SRxDqgLhp1mUs}xLCx#Q* z1mYl)nIX#|^a%i701P;WxkEO}FcZZOSBp7=2Jf#7<3yl|_>VGz$iGrfB2Gyb1tq$7 zKp(}#ni%+2D+D>W3z`j zKD5c8H1IF&F$6@T(VP8^tC(SZP^xh^F)EH7W6X$Bl5bu53QV{Gp_36&(EwNZJ1tJT z^a`DlgX-&mm}NhR`CdJUd<^9d!Y&tQ1SBm%#8dABpb5Fdv2^%`36818+pxl8ex+37 z`GB8I&%xej+CbO|4ggc{jWv5zScZlkO!;3upvU}`J}is`7=UsRWL1~}rflW?Ajo!C zN1L5U1LIVVaQQLlj*w6RVStP&*WR%)Dr8@XIkL-2WO}@*pI&cj3i($q_?C8D`mZMJ zx#qEfk&(LTb!`&JP+yT5QWGR0YOCu9lxChnIZokj`=fY%p7}P#xPtB*Jgf4fN%S@2 zbU>N~23nxjG-^a83pH+vFyn-nU;>3#JdBOS#mlC-aYi$_5V*o8XL5$1ya)<1VX!mp z&|>E^VSk{fd?rmBX`1^_V7tb|4zcD5yehaf_7;Uyk0yk&;Ip3b@$s~Qdn=aBPY2}LJC_JS^Y z@2-ynX9k9>jbl}Qv$1PPhyU0|t?GM6fF>K zDW0uu73m570)ZUuM`eyULPuTp59BcIlhh%~Ff>F(BX{rK9cw6p88m4mabReuW9==% zqsjDhgFFOf5UPTTAl8)x6Ba$3CF!z$MgbhtQDQmj_sXAq*&crDmYab`^8sPun~P4r zGd>dtgJ4q7C#1pA2kx3Q<-*;7acqDWODLwg+6eW3xNS182xI&{a5qpA^-&ms2ol~> zY-l&fzsW!2`rQ=(lJ3l%CVO5X zva8_7G&DN}5Bh!2oAKAnOj6L#FUn<^ve9TWDVjY+iE0F+|-Lik}QYbfHVTWWXbHJqY0#(BG? z|J8J7lNlP|?$}?TmB|C}ohMgWCCMlsg+#lBvMCrYn9Y>?smFm+D{6c*KE{N!8z#8C;J*}j~6Fnb`m)mtfW%*s|_6H>5*LzR&woo;_T&R04cjRYwTl{ z3yijDpD(_?cKa0&<0gFvGjw4JgpzpO;`?dJ%=SkZ>m;*$qD@4*Dc2|=Wc&f|?kwfL(hZ zHpE#vjPQr@GbtBklnRA4evITWLA)RJd$*u)zVG1h2{a-fs6xUV95n2^P9!|#9TICK zX05;ZWi@bR%!G>Q_H=`tm6%A{vDjLEEgu{ZQfUV9+ulw~CHfog;%TdMp8*|rk~ot$ zMTKs5xN%xCR9AJhN1(53Xt)Gn1M1qjd55u_Ck<2VFYBkNxM#Fp#Bll@33OPMAa6fT ztKX^iva;38%F03=2y{PMYE;^w7e~gGj}Sr*AVR8lBPCsd+gzmLLuB5$^RBL~5!?$c z08@MNDAYYIic)e!Rs zV<>*_N-nndwh0U}!vXv?0>~g>VyV%0ovE9}yr*0KH^zVI+mo5?S5gNxxa3@Ug3KGw zuaQ;=%ZKajYEDRo$LtEe-d~f+IEte4!#r@WETxM@L5~`SHV6Hol%G#pN&8BH$TYe0 znk3`^_>IR;$qgOm4tjlQbB08Gv=OYHx6#FQg^Ssh_g-A?m~ah%L_iX+4db&#d(QfBKwDyLY1-#PcQaP}AbUisy&v=?N)R5y>mG4T#|B zLggYqFoh#ZD5EU9?zR8mt3#Acxz80+-HO0gm)XtLMPfC#2G=3L*zz*#3vxsM$+4En z2=yZ$Rw)-5IyN?@dlV`#lF5xX-Y3U^^ceg%H5|S>|FelE=a9dK&ZW)07F;zq_$|uh zF0L4E4N?VXkkfWD_WhdV8~Q6t*}ZxZzK_x+9`X-8+aUr~Pj9MvL{8ZN6`5gQw^9DGn8kmyCHvC)51pBg)0$rL<_J?OyZ=}lD{-h}ClE5* zF2DhF!(|bF&Q^w+Oz(V>t{X?dZ|+}NwU1s&9oSh3 z&s&+?Y8Jj$UN;rQAafFZ(tdt?v*7n-w6g#E{I zSczG#+Ye5sZ=Inc);fR_qA@AK8f4DOBv-qOXDUr|$PBNh<=j%v@lw&(qUH8uu1xvRG+n$!T=!qHY#Rt{8%%{g=gn>4OnIrcoz#8rN)Ce44GbrquV>)mLA zh>AYCHtNs!DOL1gXqpD5<9nJj$%kxql#7B)3ObGXH*4+JS!dop*(?dVvuULKDbXwV zyTI0m&m>|VNef6j-}((y0TJ?ayjhM-@!4TpD29JI+~nvJDQ_ zyMR?i55`p-W_#{?0~e2E)N$ELqD<;jULBw!$DmSxaO5T6RKA*w^x#FK)ycPqKW!CfpmKE;|5csJaY`ClzSi9eMsvbUM- zYBeJPZS*K^i%^Gv>_jcW{+7Lc&Mm7!<+rYb5r?77St06H$Ev2tA@*WUM4+F4Qy(rvu7O_JQrV{m6`H;=UsDk^|IpbAN(QqhCJ zRN;Ym>yj@GNlI&a=4sylSV=_f`u~y5z5z@m)!Q3B`KP{k?hE1W!j}K+efj<;F%1FT zk{#~AH;$(AL`oBGP^C(IXMRkff+!cs0?3%j@LP+bQYhR1ukg;fyYe;Ffpc5AW4&># z!72x6PkUX&zg2b}aFtu61DVplISimi%~fdWV?DpPG4cOId#M?gaN|1Qa;Ah(CR+LO z-_fRY8c_8BS^ldNHwjB;a68O9g+~HSpk7LXZMRx@g6woi@j`$lKr9Eu=oXfNai;!^ zt2Gn6hvrmN*+)B(&-_XW`Ko`!-B%_L!*OCjVT);xt+qnx=0cscM75sd&J^0CIcC2Ck$vMAnI+mY!ZlWwE9k&AY zeS7TAZ~&A}S64o2w;K{pp0SiOX}2JY1M%$~3}wC6I!%!ip-$fwt&LQR%>^c2%0?|% zdvC2{;rAAJCg}n>sPp1)cP%DT;02IpLZusUiYJJS|IT+sK3)O3D&sC3%Mn2g_Hrvf zBcfynx79$k3E|Tnvr`^4f!#52y}0_65VYh8C3Q?T6i=BvoYX?x}|f41NPo_1qICiqLuR zPVcx|kmZxQ?oSj9IQ>t1-#1_YC;&}L7{tSZtfT4BbO?n)j{o$1zVBum=@q$48h1nL zIC>K+6}+uLA#%#r+;K@dyOQ&^zH}= z$mMkPb+tY}H_w=fcaZ(T3xNEzEKQDiU z+yRA1@m85oShTA8kmmGExdW8ZE*PO~OP)FNYx>E>rFT4h9e;4fqE)wmKrgW@eeES4 zVm`E0GgoW0fFl!jZp$#PEomz!pHHP~;sXK-aTE(TRs&^VV!S8bGv;)17hEsr)ny## zeLN1m_4uMgewPQ%86nN%v>hAoHexO^YsycHtw2_mK` z8!pbts3ck>RJ?#v4x=ni%?OHd9n!sNS)-#{En}OVW9D-tQt;*8Ujy(2kwDt9a@j`V1}gU^mc>ZkOhUjlYAtUyb{3(5Or7 zHc@{E2jSTcz=?)ofUNMMs;LkBtZ4-2Ci4p%uJY^IDO&aL|}Jh3s!Gm2&njz$=}n?roPQkS7!uFTnKjs zH#@y7roTlT?rNR*b=b>-+^soHac@BkN5BLtlMKpy$tpHqJZQfDz>(fVyMGKF{8Z2> z7qO%AZ~@}=+`C-o1c_1bFj@@!4uK!dbcZ)L)D8!(SY7o`pR2WkQ&@~M7F{cxke2sC zw;Z>zP)9nfIoyb6R{&nUL)k%3z!i3Tg$BK0HePDLLW-gfeY zumns#Q``Rj37+dC%DGWCSBs8dtUu?@)gEq~P77#dqcWC@FVKiOy0Z=7?;)8?Rl zMrZxEA1Kb z=|I8sCr;_rViW32V{q3%l*#q6t6v>c$kqe~U;oQ8yia%POBjoSzST=bp0TB*WTsG#JK2vS15>mH$s{$Dn3} zqI3O5DqyE!eo-n0&q#9GaZ%Z8Xq|IpRxaWM9A$rW^jzv&wX~<`@1N+}o!{GJqQCI~ z>J(3`rM`5PG=J-%dyCDD!SKZGH`(M1G0QtIyQn-vhXPjR;f}!A*;G0g_HYAZnumEc zJgI1v(>mri&O*(jT(nkG=*BBQq7AowB>s3N*ckfwij;c9TI$-ppr>v>0)cw~-H2I$ zzkiSvXOs&(x8gl&`)Zes4}s)SGEJXN&m{s*3lyN|YrYxWKqB6}TC5*KbLcqK)XMwh z@4*9}EQh(DtyE-tx!Mobe7pFn!g*voilL*5R^_uYtK3TC=c+9QI9wgL2Vd+4>5#X> zujBs=W|u@g>8jb+@oZ|}abvr_*mv;;U1*emVX|*Po4l<|#1kzHwsWz4XN<7bKKLw^ z%(r8q{vQAH>{-C`CV$=nr z`j@iGOR*AnyHfrHAuY~Z0*5j3#)RyuTIr&d0J>>S*9y*bDPf$%E2_3t^+-eirH@3^qMXOgVkT z%{SRP;C*x=u|bkC6X|Yz>^)*J6QynLn;W-zZvT^9LUp1k0`{Xdq?nOlX~b1&lz#ne z1HqG*E<^yk4`c1!_}k#&_1|R-r5fw$u&Bc$!zR(CVmFPtnZgfF))>1O!)FZYQ&TXwBij1_|d*WGP6Ho1jC=jkgB3O;G z-Y>sDvv@3={3kQrz)xFUX>I@xq(h@*%q(`{%M`x39G_G=TskLcLv!P0S{hhzdS}gb zXKnTH&acBex!LJneC#jt2AsAOdJl}upC6%3KK@)rMquE)(rQBd^UD0XWmxgh7iP+S zM)jiDBG2=OK69w<;Jrj8S5snhKBUKG!|QzFg(?@+{jZU;4Xt`-%eThZZhGDSaBnFt zeRC|8RV{@f1mXE0#Tq z@KWPjH2X8*u-*rdaMgipKfjFaUoF3uX^2@Z1}7Q2Ibxl4s*WXtAC5}__r>CPa{0v4 z%K&>g^}^9(dopOR*v8+^d5wKjA#wcKJKi z)hjpxM&Y5wi@NoM!PGc)82&II`*G4eR&}; z6s6+w^5nsVg$3N_=Sx;DHD;l|WI6l^q$xKkr@Be&T#x&ABha$2As>o!I5dIc{8m$gT9Q$8KF^2|?Q%5`d`}@|728UZHcnAPJ0oqfBlvFPS&iAu3cz?`ZEs!X z5?ztUufELqj~_o;DLntxN6YTQNqz^oe*q|>H47)_jL$VrC)l_T#aUBdc8V7Jx>D?# zEFwIp$Bl<-_GXd^M^!Y^+CZArO_?jn&(F`Hqpgi;s;eXF0)Bge9>tNofT9| zK$wztjhP5PlM~qV-^3w?&VQgg@91;!FS7yM9C-4=g%>3UBc@nGrAa+uPcb~p;o2*w z`n7a)LTV@1Zgcp=aQ_{2stNhN=Ic=G_-d%CC|Q097~AC9nX1I+5}l&Q|H7(fGKSLs z6>nM;u$i;cTe3p)u#{l(A|}}9L$zmWIo(G4&d$yXi2nsoHi!2%=R)NGxtqc6Mqo%# zP*(X%>2e^)jJ$Oy?f&Tu+ebBT% z8jKIQcM}e5yVAK?*D&=ba&vRjy&iFPw+>Li5*GjC~I9EIYNOoR@{=&v(i15I8 z`1)}7!>(x~bnr$tLFq+n(+*BER1uJN3mDg&c2WnGX4ZWh2gk-(1T^y1^O`=r`ou~l zVN+Lyode;PbEsoJy>SkLWHniKkuzI%u}IsYu7FXV-!Z|t18uLU4!-PzTlPL$09G!U zmyhC0(xLL^#Fg~PwY{Rw&d%crAGvd#xc(n@kji>Q0YyC1s1wWZz=@HcR>5w;z>R1q z4^WIr8#s5%2BNT6=sm4{jJ~EPn;*!}^lZ6+yC51zV?J`Uo4lZg)=aKl075zA!+tTG zvF8@DE`u#x)7K|huu;8#(zSy^?g4$7{pw!J1&t&MrKuewYsBuIdx@Q`9pE; zuc4;JbVk zh$5;Z;IkDQI2Xk=C@5&a=gzZ~6K2~2;Z;&>JatS9c=;P3L6OjQKBhd?#}YRnG~RO7 zaLp2;*o+&m?;m}@C3d2tK+e0R5uqPOoa8;(crNm+R1oIX%%`QCBPmy_;bPL4R`5wUZH4(x8&_fSSn>9`0AOXYiQ_kq+1=l z(?>YlAGhXv?AkaL65U|zabnf^Eh?d$IiCBsgYe+q0XlPX;x8s{L^D-kMP7x@N;5{3 zr8hM^(aoduc)h&59RIPX=$L#`Z&G%QE9O0ct-)o$dmxs_>j2N^p!>8;-HDmM6Vq}_ zcwe}14-)401t@%2wfRMN&A-w}W0qYs7DliNZuQ|( zQiL~*X#-lEaO8{U*yA$gTdwpPH`t2VlT2tGl1=fd--OxxkqdWTFaveP%eY=H9tDM- zTvAe!J^;3%m3NWpyX0zzfk%MH7ltw6MQ%becEV>sn?~`f>Vz7>{cf66dI0!@9j$N3 zD-a;{?aZD*7?GC)K&#y*uZdG2n)ZLbF$Y;JJL6P?wTcQft^G41S~t0-dpxX59zN*+ zq}xFM%(y^%=rkPto+bE%A!eX{d{^N`ZCr%Eo!bI+Cjes2=u`cWn0Lv@5@mOJLyN6* zs=K@U1F-I>8bqP9V!r$TmMsq!c z*A?U)IU}^sPP$8;uaPBX_PZ0($jBT7KynhT0iHu^zM=LmE~_=t2pIk=MdxJ=_IG~c zf8u4)3?K8F9?KPZWZzVB&3X=YEgTrw>iAAZhsFxqYe!2K9T9o*X*+3!btfhF^ChY; z-Ty~p2WALoAQ(lKZJ^XxXUjbQB?PuSXP$3vhb31PMF;nIvZ&HHN?drm4QmhWF9C8)9A=qwKiYnncd3zzz|Rklh0he{0d;QD~(VDFAQO=HseG7XYNuv zbwPr$CnmbnwK+YU!wl%0Pf&OVWOyh% zv@HkpV%D6XD-zi!oq_@$D1}K{Qd2O_`a~wT$u@%;oJwPx%PO)`m9KR&eEKAO4{Ue6 z_uw~Q`JbB*GkFhIT8j0*Q#0hLa?maQbz|#POe8w&5sABhx79d`ccj3xz$m5f=Nm4P z!W7{?iN2(?b~4RAmR_e`A3Kej%>+HM9zBP(<+6)#VGYs8 z6ge#0(+H+Gc?g}l8Kra&SS%Qw2BS^+2^l+qwl1ET`3VLo$GXX8V{exiJbA+{Nn;wX zt6Xkn@s0`stC`1`C{yBjuimmtAD9 z{GC@knmhgz6Zw5A zFSEj(Oa#lY=>YbWCy#N)7N8$DPWSzlW*T=kuPs>lrUf@4BzieBxm(&&P471;(kJdDXc2XhyMOdf*$WCf*cs zHun-gXjo1VcL&43F>B}M4A2w2Mdz86cAXG;+!IbKHTUtM+^}4bCMS#}l_mv#=B@`_ z&_(}F&}RoM$WmB)$a4R8I(3s@htGdk6kqh!!536dK40Lx;9IpPpIP}NlmD(ZI3WLX zptV1U-GBeqq(*&;_`o+4Js}rki>J7<< zi{*kbZiYMC*WokrO?#0?jd5b`8gI49TQ-cUWvb*@Zkx>19s3bQzjL4Q?N3Ru-rwnn zhgPD8Phowraj=#U$beOWl?xaQ;lYFor_y5I%Q$gj{Y>W~i(piI4iR@)&|q==xVFe# zpRt^ob&-Yz_D9`JoxD|pB*w@#UB5>OL0+s)&Vk~T1Q7f4nWrgXvAnJ~ujRU^>VJ~j zH-68$zp>l7+IsqICyctUdXG@N$GS`Q3x!XOe8q^$Wt!4Yc+{SS^|`$~CPOm^_G?UG zeGvuBSpF%$nTp`St3C;5mQC+;AdO|x?sUXH+gQAK%>`X4<;y1XGoEvlj@S_!x!WpP zJB`;rguxG&NFC?Qq1u5XO3B%uyO;w*4GTqR5mQTiGUSWn*xSLJe$G1z!|E;qtFYKi zsSdf6PvwPToT-tnF?*gDEJ z_Op9z_eu&OTvB)n)vcZKB}8FM!Tt86#RhR|-r1Uk^fNho%{Ng`*0q`b9<1-RnEhA8 z^%=vvZgRK6d@^WA^59JepVE*q7V$~uNl$tcwf0)BMNgowF?Q=(D|hq<2iZf3RPiSSJO6e3?+<|wAxlBc{M~gOcRKpU7DnFonD(*lUN+$YBQ7}u+q>jR?j7#iXf-q~8P(T)f6(+*qefDRVZjW;&(sfueH zT`-ocVxtkdsN+c_fuk=zqAh zj%vY$#!c%y;{@BKfvixbCD~9k_5QFXTDL)})_r!c z>>zczAs0K->>m5MgWjvo1#bC$u=>7V>ZkD|R#);`jbA0)Rl@zUY@Y2Se<)-v%AgiC zTuzIBrg#{Ul%CP!^j+eCwRUD}U++6^wks4A%swZ-5ilieRYk;5;jz!&Xc)aBHu5ZM zoK?J!S#j1qyK}mNV4FR9dA_4jmPP8*$QmKlGH{@`Hj8~c4m z^paTiTDznW{kyLU&vi-TNYWDd4s=H?%<9;OZ(WQ$aNl84bnSkC?6~5NeA!^*gFCG6 z+R)XpQQr*KZXM*hvhU`hn~(c)P^ZI8X+!irynGN{-KICAFjMrBj`ScK|HJ^Mf|z%0 zfL+vza;zku?uU{vz*)j}6_;ZN%k$PEGWpXSEK{uGM@hrJ%PbDEV8srOZE5NuFn@>$ zoFRCuU0_fhw1o9FWRr=i9T|~#-!x4^A1yolwmFC2xwl&HX?|c%6?f@yPPuXDL z<+?U{Wrn0}7ti3E%F9h>Y3b4B@7vLlb~KGNQja(UOnbDerW&39Y+%oMnIKhZ`}3od zY1tS5iV-EAI)*)36|~iuU+P8r+mG}=1>(~81Y>;HxVH1riYa7)5r6KOe35r~&3P}A zp9cm)DAGEN`fBk}J)9ZYr5NxZm{kV@8C986e~^&tgXZ5BpzpK`XdTR|MAf7}VADT` zlxY}A?hyE1S)O0P;#Vh_9ra%hpd-<>>Pw8z-M1cGOG0QpS-U0Ez7n&hsk_rMfe*DD znl&<5eqA0`YoTg%%4HC7MaP?N&EHo(CSORhwuq+wmG*gdzG8>lTN3q(w={mGwUl-) z{hH++58-8Z&i8~Md8lrbB~>uWm{ik*L-~4-N%!HR2*z*ps)K1sVH!sCPs$HpzwPx+ zf!5+h-u3moU0ml2YtpfWe44xH>{IV9k&rF(W9TcyW?4i+VOa6vG8Vr!!Q67`;}FAB zyQ=}s1z(ST89@}wt|C~Z_UKh2T}$5N|E|zx*`4t#;MMmX>xCLMHbWOm<*%EI4*ptX z8AT-kdBS%c6$|>d$dmte$#Pe@m@0~eEU*V-&A_%9i%K39EVBebyfj>x;sTtrt;>JVG>#CH!TXtUq44+x|6k{qwE}J<&H5 z!n5@Icj=^QF%feKk99G%qlpPI)v5f`^H#+DZEHu|2j+Wqllus>Fh0BV>rYf496SvE zb%T?vPr14+j$!}Y6Jb3n?-&?rEEET^VK)%Ub&Sl8s;P)g6~2ASF$sIB((w(n)|ga{ zrgE&7C?$Wdmvo_~HWkHG@0q^l-g$p2^=l&LHn2g4<^Y+wz<^rg zd4JB)Ta1|HeB(gIO{?=W;c4twn$65NK|5-Zy&9~k3JL?C;~SB;%hMn6qP$qk3_Ty_ zB33DUbM?6C=C>9YM^lsb%npRVa=-j!CY!)MBFm*)UshM!vdU2HCIWMBy8Fv~YETDf+@5O+{LV@Bfz5ddg z%=0Tf-<&oS&|Wq6tch&~d&*00nKc(rCD+(yx1=s+;%9HFblAab6z@-6vS55FSMsrT zPSq~QOLcR6&UQqbrM^P)>061@<~<#q?%zEKBtIypLc;pW&+jDg|SE$NBiK*MfzReVRtXwg<*s zm2$jbfsI|xAZl?-+pTqPMrU(8`Fg3_!~Vg+542r~*9Cuwm^3{0lxcTLGFGr?vgLN! zu(OO?->rlnuNC^4!{nYP#vc`*k>OJ;+;W21i7-B4GNP5>j!uOVUoFmVq^_??2&LsZ zYev5KSuFrFF1gq@#y($ZB*WDUT_RZ#ovryvUMbkhfU+pBamfU1heD-fY0Ns=BMli{ zft+|NsQ*6aABGKm?>ST}pL~iY?`%&IY0fV^dq^)tu9|JTx-UFsMGU%C;S8JmJB}EM!UuA5V|>?);^Fs!Vq@ z?O|n|$oE${;GtrRePgtbrMR7yeywmtfa8b1R0Lyze@!Hu!HAb%oJCR+eg&NVKiOzo5=<7jT5@*xYQa;*0a~=cOo3l&IQ7VeH!% zNV&0+!Z8DvmJNLvybuWAyB_6p!F0bS(!2Y&iYj~i^e59+of3wcQX-MJ0eVooPFbF1 zteCxs03<4WWaD4nF;$UGja-}avrSFQ8cZUgusi`2S~j`^Q-76w4CDwcQ~;O&sQz)IhULiHne*EB5mFyv}V!N)mC}P?hm(j3MvtIEg5W5tT=7e8#6s zcWX9$)0~j(ko!baZklv+v84AOptp>BoRC=z4s;@YbjRJJ#LteJ6FIq zSI}N}g!pH4&2J)Uufu*m?(oWRhRPM>3pO~%nhwwQt|G?yCG6Un#owxmJ!iXO#q58kwBxIFvP5>OnOhnkOx^S3yZS2OdEnsH zmstiEdBcO6`enod-nHe^Iuz`?NZyATt4~E_6%_EKn+FDA*Ym7N=rQDR4 ztNWg8Vufm#AT>8xX}lGmP0f%|(*FJ`US)Iq!8U79qP#Y|TTb?nai;PdcACgp6+uK`S}w5x$+o>E2cT6^IOLH+S8V=; z)G16l7zf^fW0o1lxw)dD<(b3jb>bnuocnL7U`@SAkA^8-_n$H5d`ro;kcR$q)WJ)q z4;j%(l;3x@zDm;4q&q6GO7y$$n3lb3ZOg@}@UE|cNp(>7*1-=4rnd@x$fSr-Wtf&5 zMq*)m|DTta#s)O(ZxAA8(8VHo8(P05P*k#xu)cFIYg@+`%iWze7DIK1`vo71OgQ&H z=`FjlY*#xc7-Lp+K*FLV4BgZxbdJ-Kc%K+Oo!}a0CWl>XJcZqDpdt(jcg<}N*O8_R zovrg#X-AJt?sWVZnZ1Q<+#42_>lI90-9~XL+$7Os`Q+dK&eY%f{BCft=%TKmBRJY4 z!<-XZu8EZ2D>SFv0^-Xk)hKZHGlQG--KEvKCqwaF-{K$ziR_k(A2(uJ55C4 zolh{Ueo$?A1C8qKc+pt!_DBgHd(qctM1S%FGvwEj(`x=eA`9k=_QXa zmC=nL|C8PzQ3`Ej$xf0TZv%jb^hBS@_nf3*8hcW>z$5=*fpvOzCtt(X62>_RghIfL z^V8GH9|m{*HD~4}!2CIjuNrQkIr08F%t(H$bm0&ybuf|sRKI(rOt>+hqw8rnrxZsz z`NwsbmbT*k)XZ zi>x#&Jv6`Us#)*V(|?cbS7v3J!olT#|4y+!T)Xqr((;S4V#yyYz7#&1v_cBb&c#rD zDQ@}8{jx2MUbolbLzThTR=00E_a>zv)fZK7s>L17^^7r@flk`qu7=$RMq5%6T)zH0 zzUg`Xh{@h|@-%&eCVN*^1eRV?16@I;4QXQ$@2$ozkfd`*g|T(Hx~V9?GP97#%U04fNTnnn zT7Kq>LiUJAPZ$*`w*KQ+Xh@IXa9mXFbj=IdtSpNp~35C|63=N)!?DOsOrmb<|;7$&&#vT)D zSlN~{Q*Tx}RVCM){X(#*Rj9*f6*v)f?%%h3DfvSkEPhPDv$HdsV9ve~TMw67qxCLq z*}i~o1}{yt-d7;DW{=@!9qJN&T4TQFVn_XU)Ixg9pjqb1$nciG8!F!=P=~hPY>`k7 zRq@ZlMpm7vCI-ILzm&g4BUv)|E>4Qrh1!pL8Oe^ul)sy8b<|_1y&5H|(Y$T}_eXfb z5Lu}^!ydBlb4g}1gPaGw{5h|{5k*eG37&m$qd$<#3>_Kw^%>b~*{&4Xre_y9+;fN% z>V7N!JT>kDHsgbXR(z-ZP)~Hm=S;GAyC|$FIMm_ql6klxW4^@@0a3baeOZ_|j4^ ze!d_T94{~Snw_`Y^C6b_4(Y8b@s(kwQN6Z8b`Ftf@aq>n(9>gW>*dH91F|u{v<#-> zQZer3@Uqab7jmx8WIaTE=RI?JW$}=cB^NP}e5rK~ z8$55ge_>5-Ltjm8Udp0=1aYB~cUF9EDz53=^E<(*KI>i?2BF?ytqQtr9?U08$&bkK zhiCfS?7nGOOzX?PIh2(LZX$65g$FU#kKzq+Cl9`zg}4runUZSb6T8fS=+}vQ*}-a* zqff1#eLZaT+$`<+@TS@~}_#u>sXcN(xQon$oWX1sL_SNYspI zneXGZ^9hrCftm+vyE@Tpy9ji07~JOi)8s{_oN}nXHECo6fREE*@Lvx*e~oQr`-j9^lIgR5X3`HSp2hnr^)ibx)H4aO)#^HHR1kY!Gw&<$lGL* zJdYxMi16}qX|8h*M?zoeEaeL^`?lMZSnFXTk_SE|&Z(fvOSk4TZRE=zQd4nS*qn#E z5_7wF7HZm&0ao;s?fkOgNfjf7(h3?LZ5J%x+m57_1y#YQO_m*oTvW?$BFbk;{WWUk zY=d+IW&55m_r9^5Pb++A^u*Hp;T&`E4i{!qMFGH`#mab2@x!!ijjQDhef3=&H^P|iHLA$E>pou(q}%=S zN0VLwOti0jj0Fyv7GLkl<~d9fAV!a3Szy0hw^p=i`!%|cR)u)0Jk&SG&8n!G_PqnY z-BDA!|G4cWk`*qb=KQ3NI1DOwEoxm5-u}&||9m2HU?8?o>Hn=IhxYxTw$Gd9q z=L%o7J!W2>?)sfoNW);8&2vy>G;S|_c^RNWjfZP?D-8;oUON1w46!7rk5>!I6Z9F+ ziczqjG1DUt;8R0S;P3UlXu(R8&;i~Z+U+E1e}?;JiT0CNQ^Sz(X$S0ki-w8?@u_26 z0~ejK6r%@7_F}M_>>ITj1)t9t+G~wMTwAMmG9J(AJ<13a2wIz{PdUn#yuG{+6IKx7 zD-HNB1{1i!oG6iL(-!;~Z`#hig2rdm*Q1!UVsXt53}OO<`w?HRoW6bY+P&QzY0o6i zYLRXA5eqiQqS(uRuken)hBN>Y#Bez>d#*==3}O0&x7u2{A*v=xenh;MpFnD9`H4$>1O+6 zV+HSD8$YfHU<2Dn-)Ii}Set;odygJKdb2U8UbVlCt0KO4j@KIx+0RdE8+7at!#z+8 z8bsLXq`XGc)H-W8Io@8`Nw!?BXBF$%g1kRGDJmm5<;>~Cpat>=ZdT0uJMm->7Cy}m z`0n{f-c>1e#>o)b;$%K~bJ7bY;m%>>rmZx~yHk_c@k!?C=7;I$M9i+3z8`t?6ns1iYsOqRxJbV)VzW`o!TQ&dy diff --git a/img/rapids_arrow.png b/img/rapids_arrow.png deleted file mode 100644 index a0543ff677c8f5fce6db6d746d37085efc46fb84..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 192477 zcmeFYcT^K^yEcjl(nP=l3ep6mgG%oRN)@Tni;74G=_Mo)m7+*dnus)MA|*(bl0-#B zY9N$Q69MTZlt6$Ga)#f#x46Hv*LvS`{{CExBu-{B&vTdSZr40Fxvk60$jL}WMa8VI zcf*W|ijF`f{-BZEX{MZEZdif8U4hUM^Hrde2j>=&a4#*|Y7; z=outgkBR1m6j(*}f?l1H*E-FYLeFr$H9K_!xpLUuNB^>TLH;nFXIW>H}8J7%6MQt@47Ee^Q(PBoV0*ke!s zji@LpR^?o*+rxVZ8WycGsb@6D_@++36K#2W`khqEkXh!AH8j?SDyhV)?#dY|AqL-v zY{F-Ln7)2gbLP=5b*#emVxepn%?+YolG_Leyut|Njh)rZf=fCH-e7ZxOBVL>S zOT}g?iKc-}rTd%3^Tc^7gol(0-^Kgfu@|P^(Ksmu-(@O9^zy_{aYX;<^lN;iQ>LnJ zEBG`1MN*s1+g2v-)>>85yW<~>PH2dny+7aEndjD&u{j*rppwev|$< zx>lP5(|ma?zp`i5SHJtpJiGqs8rLluW7)SH4?FJ}VxGVse`2hj(td!jzGj?q@oGn8 zClh^&-cAZMkZfmcu1b-{uk4@by!bvNP%CO}~7ET2$rN2SHzWT&i2{ z|MrA;O22%b?$=_iH_<3T#Co=8Wj3sL3o$5DpErm(8N#(rs zM_htWgwZ1Yr3?F~Ct12O_)ZhPxn18_Mk7(D6r#$q{4~ayHe$y8PE;^^JraKNTzd6* zKtS{;{=$p4&*0LT#3g-eXilO%>_N$I3D;NZ+Yk5K${;VNC=T(T*Df673)P)fiOp)e zeUIZ)zBSX``@)r$Wye)6e|64UyD`d;U)pNWF!M!m-;}HAL=&V5rBT+M29qvH|B$3m zYO-;42lX>Y^0({5bJxI=K3&eyjEZ$bx9C(r3RHriCaIcKZMa{gGro5bUfj%7cl zrM`YX_QX0L?X9QhP}E^bo42lTNKm`vQkPLLebK3;^`dL~s{V+o=5hY7vx;<5Po*Al zk4J&~=o}wkUpsYDfL=?>jh?$ti@vPTfug(fGx8IT*+o_Tb^Hocg43||+7*6zUh%e%mC9voo=@?+ zVncDvcX|CT-EPbH%)H#qs(Nhh7wt|&gc5=c7q2kwx>mZQ9l=U5$Yl$pw*=f=gg+xm z|Bb$gfCXFd=~HKVzTt}!*Anyuuf;yRVYTB>#9IA4Ahz>)=5wR($=`jZoTpfy-#}lo zdDX1fCn#T%WT7*vIclhUeUX3bm3q>-_Pqym{;W!=5`5`Nr9VMG+uGHByZ*jMdye)D zElZReZN)dQs8mLkv!iD#Lz%49I5uAG0wozAX+Qc68VuHzW?6!6wax1ULX~3>ZxBx*Y1rc1 z77FGpcTkov6I?O2MEyfv(PY7^4_6KIb*A6wNFR58b4}|)!^_l{$6s>2G@cG@+fA5` zx0A2%vqya1|HU8{0@*CELCv>&~gV`L0zF^4m_!K^81OglgiA|%-tu?oHRU{c@ocM`f~i`wUpm^b7V26-gGMx@Ehyx?gm2cCBXgcY(W@ zx|X_NHms8468;kMc7ZnU(|Xcu(&W;t%1p#L#1VrJOMN__)Hc% z@^u0QjwG296M3X~qIl?34>oGnyVl32cO%}9+`xY4KEn&T5dH~X7q%5H1MhFH4=3)O z->C#IM(})kX>C7rdq~EK3Fn1_Bc#d8`_CXTWGo4^gWA-%@R*i{E`=`e#7El8G(IOI zSRb5ExiG*gdw%(%s_@POdVWFG2SGo(O(slkyMgELzVRGxu9TS^So^#dKbd$*Lf{#j z_bayt*1zBW1esnoeexmAjg^}det8e_(&pCnF5l$XF@k;WWdH;`gVVIn##&>;LO7E0$o47yS_UJtzMl!611F0=p`5}JiQ3* z3GZQ&ifc+6Lv2po*~oSYb+LG;iXwE_Cy(UJB;QGu`$C-a{f+E6)sB5lRu@)wwd&FH zG4nC!QIeL#$7&iS2`I{F`~E1qFtKa7>#_@{<)K|-JP|edE$h+^**gQaA5JAIao-PE zH=c^QiM=UzllS4}9#uudptQS|ZC;%toewp{Pl+A4oKlaGDF#=39lsgC+q0Qf@2!4K z0}LTi=hGH4?XjPb9*4(PiMaORg@T1rg?u@LC4IyPB^|8yc8Y)HaCW40JabV<50x6J zI_uQoG}_SRHxljQ)Y&kTF19T8QA}NWP6A1&CZh&563c#C_cfqhTvo)U#J+yyDg|5O zD=f+-Dk7RM`x|ZA&W>)>`eKZnLOykrcyxw$gs<;m7h@NPn166tdhaw}^_q+f-MllO9otdv@N4lR%+DO97Maw62za`Yvht+W(bZ^%?=rSh^UF=t^ z+fbYdU08p0rd2rl_H52Y#niAxKN4ZPsW{o#NkwMi*4)TiHAfF^mI@e4+<(sG$dTYk z_#FEC^;ltAr&M~T$~%=)l+}IG zt!j;=H=8@HXZn&+;6h_}T!Q3fEqv*sBV4kDWXvD{%ZVrim)I@JqEe2h~}jkqJYw*ky#s&0iY#M##6s zq^OqfA4{raeM}#!Qxereb6+gNAOX8B@mN_ofw>{}+C}aMpFXUe(B8RnZ9|8H%8ZN3 z>6w5)OQx<4V)^j{-{X&}G=t}o>DjXJ^`izRRyf18{QWb_q24s_4_aRXRRY5!J(~b3 zDt4j6zhnAlf@@S%$8y{)tOKo$ZYewadQ07V;CtUiD%ATCa5fc{YN#^s(c2~P9$%=p zmrsClsM@7}&QJzEA0C#z#P`oBfu3rYtc^_gw0-?u_~fKyrLJ63XXNAKQ}uuFP}%H; z&VOAF{7>zYTVUWLWohYMLfwz<r4P68*OvYzuYZmL0#j#HmHs1qb;jFYU%&x_VC9#tfS(U-CK*}q@JW2#7X)PHgswdh&CUgJrl$&`O{_kR}1*E@INKUk?IUlcXX zo@le^yVL)P2x#_HkN-QV9tys|rj9tnr zlO?duZ=U>*j6mxcHTf*|>}`gBt5!gs)Zpy@%88Sdr?TsLm5l@cJ7z{M0_bOj$;lTu zf(nx=AkMTuTmE0sGt%M^E;s5Pnf(Qq&zM-}W9&@N9Iha@M02UmEJn(V)U)jZJ0LF- zw?~5|-ACRx!U~xR3z-VSc=ebx+#99S{tp}#6bJ03Oc^994dDzHo1l}^q4SEtZj<2y znz=#@H9o>y!#bmG9Hdd3jhy1zA<^4+Gona2x07f0FeG%$bV1@}J@NE&^~m1xI(3W1 zQQ?~4$#C)Ii1d-UCL#VSP=2MQ zhuBS5Zu*n(KjoZT2*{bW@)p)+|UQWF`x8X^SyuHAfw z!bXY3e)&DLS#UF8$~SX>ooAN&ka`G%IaR5-1D=KAX1!SIRn z(lTqYu$G(!=Ry`8LmAwy1HBlE*I?1fbcK$SFS|C2c~o9mXLH|Ja}#@VIq0-4U;?ayP1pZhT!1-c z05VkAyL znof6gBiuZ2i2WHUXF7sL2L#0nM{~ zy&r39Ca$7;o)+9zLx=q+wKUL1iK(XE%Bwau%*|0k6*5T-P++o?h_2TU92p47eh;Mo zawypl;2yIOIwk0zwM3o(^uVn6wwb@zR||cgBZYHbWqh8iz!Y;QJ4q9{F?;dq1T}Y~ z9-V`fce7n1Y7;(Js0+2-Q!0||McMZOK2i)7?(q`@UhO%5#9AYlj#ZY5CjoT#O|d^(CnoS%v!s14Cm$q#zy^iUrBYs;Scco z{cQtIX5&P6W&+uRz3cTVIW1mm&D0NL2^FU%?{P*X9wnrjo#y}@HW0JLk6gnH$nGfK zLcJ(_|LdX4Z1YVRy8J>qsnW8Gkq{TSl5IF`os!ll>2!%v6CvAqwrlqGE1xF5&59a~ zEZrW&1*;8x6Sc?WnjV}*vsbwLN$!uvsHkf%Y34eKhrC`oN?xrc0b8S72R8Z7?as|N zjhG|TE~}X;sZoQzlUk6BJ6T1RDUjp8Xc(Z2&Kn4^pahJ7-eqS=Gl3&Z`jhhN#& zrYRF&a*7AyJuz~ci=P~Z>qa~pHp}NSVa#8hZt?d@>%@Eig;pQb0LK|^Dz<0)qkYXM z)LcH&+{C#`?h0!$2k-;y2|-V`dEt@~|LxSkT%<=~AilmJ<4lLIUak|wIU(tGq@#+y zF8Q{{MBV4j&7A&v@9>wev9R zQNiq4ppmVM{40SkoCb8C64#`{eJd7gf*iZ0J+cQJZ7&35D4 ze)SuhTX%7wW{s`P&vtV#r*=(^fXnKqze3V=@>w6GMIo@t)*J53Q|#2_&-VxZcpZSS z@<9IYxry+1lY zbeuy6$fLV5S$^Q;%k3-#dPA6aFHEk97hvW0G~ZpD;zfT}nf78E6zWa%{lR{82l|2c z0J2JA?^GOBNPJf#-!c)P(nb-AS9yX5UeanEC%3m0B1kb#^Hn#|D;&d*ure+6&+U>m zH7d?;Wj1~Ban)~PNr0Z~WY=kd8*`n5j>-C(mTTgSB5Whi99#|>Q?+*+6-PhCvA}Xl z^mB7w{W|xT2GRfsz*XTjtd4-mL^F4Yn%w5igKL-1k*1C8)ubfE6+VRsme;#~9s!{Y z6AxsYA>(gp_WW8u;NH=W2Bq|T*7v|EVed%15mI}g(Ufy5q23|jgG09_4lIu~QeF(1 zG!swL$7JzII$cD>V+4B>=Wcrcf%RiF9(REJ%@QwxiTVSNNwJ#cw0l_XUJxU@m}k)I z@KwI@Y6LhaUAZF6d+fUHMh3V$J#BR9eFM7d`SzC1e(w8Tr4(my>CkIje+9g-785u7 z-f?U;e;a05z8tQe%I=C-9(`z3CEbMZM=>}hcuLlwdNrZL^#YmB2e4iqt_tXHVRQ-y zrk1pN?O|;B1))DhskK;#t4b8o+=hle^Os+9LHV5I{L#gtgmKJ1s9lf#6EP7KZ zg+rl}Pp_J4)8oSsZnxi+_uDDR_7bzVArmC1@nVSe&el#6sosYELcOlXPc}2IdAs~1 zo-w`5QPe99-%oQYPA(d{Lz$yo6Su%NZ+(tIPrKj^%XHT(?twN$l^o%6Wj=1@ksKNA z99o`ZL1U<~0IMhrOG!h2J6c{hY9Iys88xed6fg5>@C+09DC1N2L)T%7dm%$Z&sY7d zZ|1Cp3Ueev!yE)rvR^5|=_T{9!0>Vmgy}F-uD@*47Ko3qL6y#+qf%e+l8I;Jf%lw# zK(Ty*`-INMZbpdWRmkwfBjdiN)Xu{CoRlB7B8v4d;Myv|)sE%Ca63Et6}8P|5$saX zw<}F`d8V*M!EKvz@Us(p#$%OeUp0@&_i`^4w?;$1PI=ZFN*{>tir;e1#f3oKe?^(i zIkr3C+kEu}Vs1=AL_rgiGz&~lOk+jv(VX{eW#Qbh=~+bdaq4wgOym9LiUY44SAaKy2QH ze%E-qJW+%V4sIy1`d{_6I*CTM$Csp=zs^BTJ-YLA&pSD5-v&tsNRM%Qw$MV`wS1Z4J8U9iCtWGlA>;l$oiQpJ zzg^MFOX-^;L8NRs)9~MI%v|~d%Q*$vDg$t7LRL&K5mIeZ z9=g=eI#2q*LbAW)4GL@ykgls!NC=&T$!k4hrgOTbMAqR_%XJ-2V{wv#Q|?<&Xw|;R zpiB^zMtedAlC(f9c6jsWp;Q;mikk0g^lLDb*c>Q>iOHuD!lYWD(z`VbwT@`k`KI+S z`TKDS=NUxLNbd9cnDc)Z1rW97L);$-e5|YCdUrrvqZg88 zHG2GLO9~6R9Zqo7Uy&q1!injp$YHzTGp-sRYwo5IDx>zQaXohgvN{2kW4vPLIhM96 zSL+%A-@{oJR^ByixR$eh;UcviNR-U+~5zn#cTNg6$OXzoV#EWMt%RK3}qakTM z4dxqtC-0}<*$xY;Qj^K~fvsvkAxzQTh!){@nYTf^%4m6{A$v+N1`Y}g8a+Qf3$C(X zsekr5O%u8X@2lM5pPerA8;A{$&)2Hs?`1K4f9kJo*GdD)I*Is3{3uBkxt^OwO%ChN z&qEp4VBx8puMz5SmJ903W9VC%EUN8QgJ03d($8Jsyc~KdPNEF4X(+#cuuKzT`@K|9sr6X!P-((iN zV&92Q3-*L%=&Bd=8EV~YRH#v!-{_U`vcHBJo1ctp6@(fAsspFF?h*0k4%szSQ z(9pS?3|%KUAs2lu0;k>4=-tUN*}1pwKU2V8)gVcX3nkw=JjfZG2Q4`$ydTkeS=9pJavU`339kkc&9)oX_mtdBc8b!@5 z-Y(+jMxyXY(k`zd&kxZ08j4FLblBk*Z0f>v0ZA$z6QHc@KSSZ-oHquwu8t5hvvEaE zNvAx|5){e@AGMQlwWcS3$lF_XHF?nP`9+ueMmWUH4%eJa&NpRDlSmr1Gp3x{ZVtkG zq9(%QJ?Wb3d&zl}S49aYKQt9}T%`zajf$2XOk3$C$HWYF-&u;=9`_2;tEi1iHxBRO zjNGq#1doub+j%c1wquIm3|fj%e#&uRNkqKGc?TuMD{r){$*e`QV5#00rpVIlWq}Ne zap$Gf_lfB5=b5P@P&9j!h=C(3(9BIcj0(G}{@zDXL5OL+ul0n`p7H0WAYLka!!09!Kq_CQvU7G`YZz#^?L)re&r}c@nro*c6&TM(8FPLO=JkvXY zeyr^&d#-U@pXMbmSnJ>@juS~xQH>azJCHE?EZqzVfMrB$LNf(`I_r$oLVr~=#ViyI z)6eBa#f=_`ihkfDkaL$~Z7P`l016*lBv6uqt|*R1uR)fswj-r2A{6I=$m(6=nfKNR zgWUQ+AmHChB{n^%uP&Sb8ebhPBv*|CV6lapewPe(%mJUDR=o3p zlXHG9(->~f!WB+=8!909vau(-mdl{&mjT@#Hh1PMmz2tF<2_<_dUlf-_Uy^kg??SH z8pF?t#GS2dB!6o$qvApizg(eNOb!m{Fnq&p@OsYpx5r4LtXB7{#ae;!y=arnFhzzcp!(i9u(xUSV_GiTO7j{x}+tfgg5IgYlA9(SO2E(qm0SdIB!OBzzs$+6`|`-pZbsj$oP@jdCVk zL}@1g^+fqYD=yJ5F{8c{m&ZnnI%lD0nT^$54f=!l4Z~Miw|cg-N&GJRR~EvU0{ZO% z0|Gil^&@IMcY>8|M*Foey;G4TF^QW+F+P0&C?6*ur|6Dud?}iH<*6^4FB|#={Drt2 zN;fkKYMBvF*H_3nhkf+RYe|lG^vdMrr9R$K5^~Z;pNx8r;FLx+6X$+qP^$0El;1N! zY-|k5<9B)^n6)=lcT^H(k^XtES}Fl<+lUOwk1>JoR{ zD`By7dL?fFTR?)?jc$LabPiA41(nCZS-71Hp35{>5x3#HU2X`1S%OWEHxQ4*0Cfy_X6fo-%VL*23mfUP69qNVXr>Xi;lSbmwsYzX}*$GGOWnb$fuyFJZe-kiQ~ z!%khd!dh305?|RgOt)3--7-BkG$!F?obQ{W!^}`>7jFwBWcwLIijyz*T|e*5eP1mx zHf6QV{){ft2om2M%ab6ct78@0hh;T7&;QKLq zddJnLH1$M-!(X0qC=4ijbsaIs-%mgJGSk~FEYpL}YW?0WoVZ&R)!^~lag{lK{A$5? zaJFoShhFyvW}&VDJhU&oujpL%wf(1O%MwPb+vCLE7_az>cl`IM6q5Hq{DSk;e_qPl zG7~sZ@zE4!1}g0<;Lx?>(DP<8AI-f?=Iz=pOo96@w$Dl4UQ`CLBqi(z1u{n+*tvb! z(;7<$f!%WP(E$AIcul>0`wL{JsKVKSd;>Bs(Zt2{Z}a@-9+He8JLacMPF8bI?vxhr zjJeqL2WAk1p>qpGom|ojy*HIf36*Sg5ud@l!5Tgzi%D*Rk0Zv^wwV?tT;KnojVh7} zB^0AUsJSt_?V!2gpz!G@TCH4nPpXxgly{Wl+*}vtHJdo*6VWkm>jqpq-ESGV? zC!k)-+YG43d$oKYlVLrb`%v}r$?Co{5;X~SV&~>#c-m0Vp0f)*adAv5{tlFA_E}zz zc^Rxk4ojJxo8tc7TtH_7x|wBBT6CX$U4G^ZyHl4FO82bw+NZPVv4%H&HzFA6c})ch z1m#}n{`{o1NSH8I-Vx@P%U3SE72mrfqCU~m7@0F2Umw4Tg*w&%WgPR8rE)KRRG4u? z;X?14sumW$rm(Qpr0pSJuAPe;cqVit0{WU8Kn0T)ifMTG2XZ6tM%tbvq)R76S2})@ zeZ&7@3X*Oxc9G5E4lh!2Vr!zam`Sma|5^ z<4;udV0mD9FJKh`$n7_k)l3T+4p3vD(@u=gek=Fs0R0UYb7G}+`wHq^Wa5us`PZJM z_;0=5ZWsXH-qM;$4k=WT60;DuMUw8i*MH$z@GEXxD-rGa1dyw~2w0Usof>`lvj{MI zx%`<{LN2b4vw!*A&V|j`s5x`TpU_jO_Dw5&XatT?W2!~I_6tcm8iTS5fEyK>SM*p3 zWCUS`<@WIJ?;8qHlj~OUb10JA>GC*ZkBQ3%BR#dzO!Z@^Lg>&rb9zFYsq^0x2M=Q) z7;UeLS{<=Gl5#aAfh;!0FwyB1ANsAsN^o5bso@?;@^bE`S{uVcQG5!|%d5@9Lu>??dX`aL}C39sHEesJ0<_*;z<~O%oOvP zhX^(^8AopD(0+!vwWW=AR*+~WZxXkG;pn89sfiLY`6$ORM;nA|c_SdQ#h?!{J>6WFrQkys+T;!pL z?W*j9#7i9LNVaw!o=z(&Bl9$HA%!&<&)BfPos&V>8VOxv^#XT160#ob_54pf2Ff0w zC5b^#{I>n89OR&9D85D)&lT~oPJAnfg4ZXXV!|2D+mB6xUV7{)?KC_~xer2-t_-MO z4AKjm32#TfmcrVNf{bEPGJI`W2(XxevWP5y+C6pIt@fem(U{d!l766it4qC!XS^yZ zO?F0|y+4CNMvXgetm6^+|q_u!!iO&r!Zq6Uw?08HhlK z(r*YrB+qenG_T9=ZsXNTP;g;JkIli&pnj*?1q~@li_)qk@PlK%&1Y?Wnlv}Ef9aE$gPqCzVfyU!%FRIHS zOBSQkWS~p&9d@*;NIKxkXVZEJbH!dNuZGMz9AN06chv-RXJF>bqnmh^r3NL6z{0Fd zoSH#cy|`5_-c{|FDpQcNJrNisc(c=x-YxPMo{mHq^G^2bJxj?rQ#ScSh+l3RGqP*p z{;T!gSZju%+U@D{ydAh00*k0H*?bY$GPU8T~v9l%pZSgkyp5fT!M6{!!r=7OTIhJnk(+4h#xdnq-fH`r2e=I%xU?j)C^_0T z@>N2*Jk;mvQDlb?Dg8~opP0&}U2|~{%iVZZrQnQNHz)Q&a^k5g8-*H>r8jbs$f#Yi z!~l!v^lFlhb&^9u)d&`y$9^*UtOlJ%deZ6Quis*kV~}!PXx`~-;ZZE!8oxhBn@bSJ zuPbjG?-CIa&za+WHTu;#werxoBFa3}C(pKTPq1aJQh>eUTrpb=-g~&d@WM1BUZLPb zBqY}CKpDI9n~lrp@^dy7dcvk%HTz##R~j(0!kMkgVvfAPB;~L#Tzy$*GnFIC^ltts zhk^iVEe)K})yRXxvW|2!GTFL%bdC+W`?j*Nbzl@r-R_eago;fo8i=MLfI+=O_Tc2a~; zp!$Y0MPzIIDlYqd5Rcjpq&CC6|5}7* zPhX^p<&Ub5(Y!?h70V_v#{Sq}Hc#sz*_%Un)4f0%+jv}CuSprdc4E&@yJExMWgm{c z8*oqhEf=)O-`8|i6NjSw;xz;lS&*+$>M1@*lMU?fZU;M^9vv-BS0T`IDhl}-Qjuz} zFsZ)Kpb^8nosx`p};WT@@VE^(h zod@b#(F&{jRltO$q#bdDW~^=qz}CBoSijk$=y{IyL#-xsJd0gNzQk|S{|Ais{qg^ zhkR9nuy)$_}F$8x5rx0a)t{x==kq-CaYFJ$LcsCkIB6%Z`0LAaEAez*jq&RZOR0^~nn;W5b^l|2 zb~rzsR`)-5q&An+fZA{%br8$Ee~Lfd9x1|nboGt@yRb0g2>LkLvb+$OPK6zR(tFMM z-EXyqfYe^cVSqw$*iu3WVEX!-RDE*s0T>p9ikILMpsB%PpGF#`b9`_;JaTt_bESpL z^kkjmtoPWQdNeF-mQtX=+vPBKgXtTlUbER`opsxhbW7iPp(ZczhtZ@fd&q<62+%fa zz5G-MPzyDxmGdVg4zQY4D2)~kRhl#p`|n5>s)N{@f?kc4ztAZO-Q-`$Dh3A1FUE~u zcGDA_@0$KSHRA+`0kFx;&cdT=FhvTOOh8_1%P0&H`d!GJW)Nz4wh6+IyIE3rT8nf#3+sAH<=e?mq0qhT$}N{_4bb9d_Y0ocT+LQWCnN zA-PVeJQ)m{Z1cLYAg4>-P%p*#sW7znQ*TEPkhgtv*z+J$|&m&2wb}g zuEVREzhR-9s z1vnWvuw^*0;UdQME1XzEob%Yie3OI8pVO{$m1uM9s8jUGb((iA^Ac`+*rwhQofDwf zKmnl%SNY{io*t!PQ~*T%53=rWHo(#~oM3aRC$E=waL+^6V_I2gxdV^6e7up%`!)P$ zz+$ago6lsgubP-zmP-QBC(}^a&KIv$IIM~6)IqyrI7UHxLmX%3+b+c4{f;a23z37b zdVB8>z{tVCFtQMn#w#Dy+Jo0cxR~tDD=I@TF_}Kdbmm>ebo-%%szNIQ$)a@w-^Odw zj!Pxismty!^HibQN(5S#((@nM9w%Fs`F$5xNSS`oi9`i+d)GwYn`~>qaNY*X@*k3) zq$30)Mri?!Z;kN&(YU50K;OV@Ow5cP;j;l?_bkE^A-NRVCDW#-TV2ee0mYUY;~nqj zfW-E_Qbx!^81qUL;(kfs)~oh+^~M>$%!O$QTKS7h&B1pOnyfA4PzSoxVo35b9xJPg*i_y5<*saJx=5p zK?TkS3bBwA$*Ulx2IsaB^RZLtU1E%{>BcSY1)E<9ll}bP+4RCw}|Li^TAwaF> zh4cd!>d>IG-{VL0VoAAcZic4V>O*%pLop6W*ts*--PX69>QUW z9e3fkP5Tu zs`hB)a2Dw4N$U~9qJZdR$0Ne+*1`~TmgS;aXaS&DOF*$_U2AAz$cX2^;Yu2$UfjsV z?xFL)zQw%9Z<6a%G@-q$?OeRiQd}v5a5m=;3rS$hD!{jij@#v9?COj88mU+iq5|7f zY;at<0+h3&l7q*#%qC5#kx}rbHPvyJ1I@ScAAJ&=M(BB(a_kN*g%>F_u{CJ#*8W`E=cbHY5VvuB6=_Vf}=n2Bpy zGXIBv`mY&sWSRt!J-^1QV-o&=mkGegM)LK|odc9i1akJFC;hqi;g8_DYp+pHJnDZM z91g&eG+M(KVouCLo<=yj}t%80H&}zmo%#<+;c2r`BJPqW9HuW&=GPdqn5|C^XC)$$BzIIq`42s|y zlNPRB)6LQ7FTz*OWj^_{HK1U+2O#sk+keq~-SU7VlM0CONJPvRb=}~3la+YT=@=*d z&JQ*KIVI|8j2eAnC!(c(MZTUQyTRQug!Fh`9eqhHM8B79(Ed* z<_~Kn;GK_wxWb^Vo`(LS%hvUWk!ZDQL0Y}e{fbs+g{y56*zsVre?@f_2IU$0aOuH| zXC5quIh2o_6)VFI$}OCH&tASu8u4E?@qBboyT&U2P^%MADe|05L76?j1L$SVeRsW; z2TSh=(j_GJd!z`;&4tO%mPNCpX$7EMLEvEt>OX#ajAnx#peTNIulWe)4 zT#io|Fxh7z3b`AS;zbRRQONI3?dJh(>X9BPQdLxSp`GOXRzp9k@y^B<87EOdy>neT z^||*2wCtf~Tj4r%SLU6H0)DK7nDha70zVKAz>w>zD}yr(aGc4uVE z{k-aD``T;Es{1>JbGiZ{JM!Iy-^mx&J)3~eN0)*aGztKxtM~~g<&FAD66ALqI)!a= zEBV8Si!({${C$wexw?7}c~u4#ZI7dLiZ4kDz^A|omYx%T>Onr2f3#U}Iub$&$lt|i z_jnxWqI;|JzepDV&nG1ViIFQ5;rp>y*tDgY06hr*6^eagwsSh2c`EBuPS?7Lc- z-I>$O70iQJ|Lt1P2JL&GP~ymn{+zc?8=-8}2n<>^8sx%1JOH{CiNTq*%hf-11R1!q zW}o(=;Sxuk}18P_XB z@Iu|D4F2uh(lPcbZ?M~B0Kjb;I=&O-7JHzj9f`@goVS*}75VXR{NsMZ_+8pv(agez>S|Q9e^dt;?KhxAXnFVHh~d?%1vHCWd?SRY z!T6HU9`4DnU%fKxNw+hRnP>|>$w&lCUj)GW=NXe#YqK|eX!LR@6h9MDIH0NVK#I86 z2K}dfjd4OZBXCoXzp1T~U+6dr#5X8;*iXlM`JeNR@SAwvXF_aP{B66+%+wrt9dYfz zaUl#8oUmUz0MbOf?Q6`GAP`;OH>4s+`gC>zFFw0;83n_!htJ}oNs(>=!I&vQbig*a%P=Xh#mv>;>B{`ML)bQ6`F-kAaqE&-aS?ym~Xa0LXX>0`O_0AwA$i8OQW!usCrq z!c#f8-Vs?pTgnfX{W=?m{Juzg>)bExO%f})`_m5jd|=i|Au(C#~vm7-g zsHnpL6>t8MwFUaBxVatmiSX>`9^q#HhNu^~w1jSrXRraUxQQ*U{Sh;Ymdg~zNnM7_g&K#+H;VWtF-{NVzDl4bt_II z)kjJAt&?ggUjiVivkcCIO= z!z;Z`6FD(DbbK(HE(BR%0Kiu<587Sf`_cuKNran7D9b>HP4nFLWstQm<`a6Owf$lV zrNN9+y@~gnxwgQcIej3)#I$)XXZ2e#T%_(~^RUjeq8w$22vy+R(|}60FEyA4Ncj~% z{s1kHn;Vg37-IFm`1P=uHZ}_iaPS z)VH<0Wz+t~t5Q|#i9l(9FfMtwwxRiOdOQ@~U(`?_*&9;YO-wzt>+bKNPqPP+_nV8| z_KrBU)CMulo+PRbkNpT=+*VQbpz>=G1xg9X0NRjci2<^ICyKcFgS^W%7%K3Vi|N^V zR6^CdPZSsPS!EWDAHeJ{NZ2~Jp}CjM9D%qq5PB~^D^9~Why8v^)Hl4UP`RGsuyC*R z0pgiY14^NAC4r&b!OFJ>^g7XOefvxyEh!Uned7wUO+JWpioAOGxWJkH2o^P|t)rTU z^#d)iVcXAf_Q#J*VB`?Y+{G`(lNbr19s#DYOGTp>2}t7{J?jUp|ipJ1MFVdG<3F ztNaLMyE1c;3p!1DJilke^nupO%coGU$%%u81opc9b@2j0ofgF@jHz~XW&T6h*7zOt zvwI^Lxwb-{2fPpSHO_8I%r;^RHw&Wz-ap20y_cxU4H?mDFaBm-nhdr%UbMYE16%UC z+o>`;YCh}_gz{pq+WpO6KXE@YPY8}dB1ZNzf>A3FY4R=G`%dCsXXCbmk?`w@(4f@| z&{)ckwvD;@RhVG_&bK+5_?O`LSOFG1uh>nsBe}7Q9u{X~;ym+XC1KKoXDFg!zI|q$ zXj*nJ>$SzNmYKEh`8r!9C>z<|`~a`NSY8wa)|)#vr{KU7q7OWpzzdqYj{Xuim3^~+ zPd%m~#_eEc)3$!8ayG0wX#feskq`?*M(Kjw^CXi4610B7Pbkg7Q9$zv;9;88(K=my5~QgHY~ca;v=p4=RhaO@ z?{7ac0<2-Q8D@X($d9xRGrFWfdIEVlLJH?hIj<4Of~?v(f5tbu zlMC08#(>etC=Prp>E2F&O3jR4S^zS;f4%OXy19b4V~*5}fkH99Du9#1etW z+kCh$0(4`uCQ=%jP5$OhL?0JAToTA)a2asKrIk>#1(N6-&yQ27j)=y0x_5DZsi5X0 zu-Y`goCbp<9b7t0SSEp4IO)Rn_}3>xIW~#pipe><9)KfJxU^lxznZwK?M@Qa&Bezv zH@kWvTMSg{1?z7E>kilq8*TD-9JfiRp@MXEw|ZxKm&$lKwI+O9-_fTkL?CVsPNtR(O6T?i1!5Cv1GiJuj z+;4S#eplc7z908r_n+VY^zKoQxAQvZyv})^=Xt)3=vll!isjyAARpP106~x||J4G} zrM&oeZsr6u2rbb*3Z;? zzf_*2DC6d}7d$mLM`Eh8Oyj~&-Q;Ow&UC{XYP~DGtH{Cm-tk@41&beJV6r9>*0uKj zW#j#i#Mi+1f^aotF|ba*V9$Pbf^S_nq}4&IK!%%=u!PFBa>yGm&~flP)8%oP5XAU&*2T)-k`vpJ;{LHC^JQ zPUh2l4Gaf`DA-5VJAd?y+dW=s(ipU)$=Y`^w6h5)M3H8*jX!J4q0CoF}ry-+A^PI z2ssF<_Lp~*4Fbirp;hu5{j&D*V3)3TPump+YKwJDJm@eKp8e?Sf_5PVaO(t3jB(Et zT}8ux7f|(cP=NE)Ahf#!|3!}uNdw8KC$#|BuNg|&&5UgRA%`7R5uU=s&l-_J$GiJZc6hhUG4h)49!?9 zI`Nb+uF-g5wYGTv=(FLKh{tMHOOh=LRg9`vd#kmt&$;_ZGA$ad?t-}zdlnty3v}1+ zH@d9JCu_!@$GIBsVt~`UrM-Rf2Ijr?#)WyUu<@;!rM~l|NP{qnyL&tJIXzqA-(Lba z&W%l}up<$FQm&nWN9}*tkqd#D@*8iQM1EXwg_2lOJ+@lJ$C^HA7R$QIScI(-tlT>q z%~)2$v`q6HvbVpO$q!@@d8Ae4{XJl{Q3ZJ_2lr{ldgVkf!5Umqux3e*q(=?Uxp}|r zHsM(;abHwSF-5l1d2x#hIwegDTRXy?Wb8Odg;lTy$sVe0nxU;mQm?XVC{;T zXw@=pX7G};@QpMj`xOteQifD_g?(Q!=FC?lxA~`|^0ugpo9#4S>3pZ-i%I^~)7G}d z$v<1B;rx$}@cZjWd4e7N9@4Ff>x=%U{Vb95SCv639Mht`bCelse&CzoRK?^Oykr61IX z1zPtH)7(Q3{ArGVcQ5|uPgm1{g!Ij4XW{>t0{n4V{8=(T1c1gB@jsvZ2~qv|j}re^ zAOo=L|EjtCubRvMs=55H^w6J~@n7=O|NQBHrHB51k{)W?^JX1XWS;PWMqf{248Q!p ziFV+Dwk$=267oOJL|-T>TN@(gBU4JWwJFYH%&NNg9CBZ-z<+bmRCiJan!4VLTY+VD zbs_-NSu-^2Jf5fqzgP8qt4d1|dGyI;`<^$jyn#A5t5w^c)al4078(wvRa!=wbosz4 z@}ymuQ7q*OZU)dhFeP=@NK^l~n7?l56cC@P{m<9AvQ;Z6jmb7Rbjw9MX9pL@H_Fgc zC)%WL^hH%lz0S7o7^v-QvIe*6x|jwc^SPn9{vkQK#zf9hvjrL6qzjYuIx@WT7+uzLLXQGFkvMLHN5>W|BF+lbSL51A4V7u z1!ySf&;LaZ63sse66;%bXqXPm5vejr!hSBHdppFuc^x{jdZwc^3A%U!^qn2NP=)ZcATJoVdi) zdtiOl1V9J%jrT3OYpwfh6-`|Fj`n#|Rq$%a%rk##Z?#BwdHZG1Reqg)oZnYduOlnG zFY!9~M*?MyCxa*p?>)Cp_d<#GT7%gDz2l2c%<$END(%4E{3+_A2uODM|2OC_Y^Fk` zv(^%F-ao%>;|{;Y4@m<0lRm;r!Y=l~78k>vKM#B$O++#it97+@KcCpTRu^OquVZDb z-)~e6M05vxIq+3T_|BV{WOb6lo$+H08mcdLvfiHBTy^F{e}{T2D^1@2NoU$+|A#mj zQLlm4Nhw9Qy44CKwFH|0g^dqX#!5Zc(=o8Q$ytVjV+raOIdk4F6HWWCDq+<1$9`@kvM=TLO*nF!1P7-L6y?7DRh&{l}!A6-C_p}MCO zUZ<2A%y!l@lCC$Lhyd%o{Zxg0N+@)t?FMX24}M%JK#j-q1!sjFB<+zIf4EHVrGF*i z9&I%PUQ^2)^otJwc_}O3^xkn`EZz5v{@vnSYzViWc_j1n z-J6Zx@9e@iOc<HGiVp@J^&q@mQTj*m3jjP%~QWT3p5@@1V{R-qo-cY|M3?zFi0o^7d_DR7--f3|r`1b!#tNn5B=qKuDY~ginTiaKRCMm|%JNrdcx?k@mkmBcKGZo#8 z!nWyjfd}|G;#x|+b?8=PZmyg03iveZ6~0Q)QkJMbCO`kiqAamY6Jk@RQ=$jR4dqmC z!;8$gRZ?(>rvkmY{g&HprJ|~aI_7l+)R(>hyhTB`jHI5p0g%>t3aEpEqdrZwW)glSY?13eaGI{^-Iy6LK>rdyuIqpl7+cu^qe^|+ zo_w|5;817?c2ri_)k6E^$owZB+4RYwPj{CWJ+d#X#z%f~!4JhugPU0)O=B>AgzACa!2kJU zvpX>NUbi@tBKa}*YawQi{z8egHlGhSvg7_FV}DCoz$J#N#iJnGIef3|HDPYeZSHzA zz;7@)oQe_g{P$J&-3Qn))Iu zh<9cBBU>T~z`A}tT^o;0hE#YEZXSD(UCHw`D<0a*nrg~+sJjr>U0}aFYb(367PucG z!S-qPcs(*A;!u8jo7fb4r?$C)s(s6!GpyZ!l||LxZvV&1uk}au;S1bNULto(UEpuL z$*9H_80MF7&E}q%)1ljjXbO9eB{cvjvQMUJ>8#2Pt-MDji!8kjOs;8hxnrG}x0`Vhq^`gtD1b(OISv=ktU|WSv69!?V zk>QBG4bib{&kZEpXVve3H}6zE{OsUN_zW%Aschmz9ejJ0tv1I)JcBZ#c#mE396#o; zr^veYdQ2c8;<}Gh+^2a~f*`XWEW-;vv@MFWIYo$}f@R9RAF87GkHvZIKI{46eDHhS zJm#CS%Gshn6U-q1>Tq~@jQG1T@7=s@Vmn8mcV03{*mU+Y3=nF=qT^72=#3cd??7S~ z^WeViB3dJR;K}xwh8qe)l!M=CUhXV^e*z$uk@knjF8LLvJCG;qKK<_Z+I3VK*mB zd@>FFqq&q2PccsO_vHaJ%EA+%b#T8|3|>c2InC6jBMfC0}?{- zu?D$!%fgd1mYIV}a_SafooR%2hC1tc z>8YrBDo~cNw4hEtX%$Xe)ck-g&w-T_dXGV7J)c6(WGw>X8#)Qz3oCH>LTcM1BwKn2@YSN_MSfdf=> z{lTW^rGDf2W^##7^}$;zE`Vo$bfS8i&KZAkpj-{$JJIBpWvEEP2Kx5ToFSUkLNE&t#`a(2btjR`qY(DjRqJ^R` z`Dw<^Lx7f`Ly0)?-!Vb{#>L;j5d!DK`q2G1%(=tr?FS9 zKT_mVj4l#vp0S#$%9XIKIX_#`fx)QLN?2pTD1hzlfFl%W_OjR`1e!#8FwJ==i%^?c z#|`(Tb`g?rCy`*-(8<{f;k?yu9!Bp<+pI|a<5a3?st9gx=~yE*aM)_DJg$;iug<|{ zC^}IdS1P}PG)E;V9?Q*l)+llg=8je?VfcrG=%Pfowj?%l*b(hdN^^Y~t5?w< zsat(&f^anoBFt*Ibifr0&Xxv}$-mKAI4)e-Bt$_otfpbu^z^TrMe z((TFzvaz9k6vCY6>w{xyriHgxUUml0du#~9>CHC_V?Z|TqbK7U)hmO%2dG1qb3TFv z>V3h2El<-4%|ALPSg_)8+p^8!Su5GP=t)|JB}!1w(+Z;<;ay)O?@1@`;zQ41$_6o_ z^A%3zJe?lF*JHjf<6w&Np)SA_Z-nZ7jHf+}yeZI4V`ot?SRQNiKtlq5j^)=*Wzv7L zyvW@@G*<;=U_ddpQpd5bq71@!aMajz6}SP9+~{8j0X4e{FoPHQ-lF_~8~lM2xF>e=i*T&{qtZei^bXD(w1j^S(0BPFI3fj#a^*_mCddTYE1F*sXbC~o(I z+(p+ptnQr@Ge(jffs;d{3Fq73*zn0@K+b7B#NGW9!Bv5)TSLtti!w{vc^cUoKi?>y z;0P218iY4;zPJ9>%t35yjz(0^AN}1&WE^aH0(~;oJMZAVp%rp+Y0VssjzWR zM|gwDkyCLJ$5tQ9FLW(K4^Nz55OQw?`_C7qRJJVP%iP+9m;C2%@-=mb^pqMeT&j8@ zupj*fli^m-o0gIlNfo{??2sB>OWVbbaB&L){&%YUO{v)Yo-r|nP^Bde$!q+GWG!qj zZ$lZvUl$9jj`#^E1Gqe9d`-CIF*G<#-pqh8(+`<7mOs=uCR_dZ@h8(eN1wbNjvwDN z0V6*@CgN%favQ$0=OD5iPUg&j2G2;Y2{rcWU7BazH=A-6i}d)J_$96cqG|uNscT-4 zT7e1$xsux8r32;bY^}-ZbqZ^Vh_`d-vBQcraVYOu-%Q`GVM6UIa1S*k-`P9ML2DcO z!WMYWD=ZWL5G_iXyXSiedMy13#!)wgI;lF|9@ccfJf>UFc zv|9Vb>_%7O3J5=&4zU_k=ORY}yl0Jd=3$)R!~FD*xE2MW#2~OnhF?CTShjJ8Fe}C; zdM5^l5M2C=mn?^I32J43$*IVE$BhHKs?<4v zbOnMZn2a$9^(&nb6daYrJbEzra&i~b+CkwK<5l`PS%UQ8LsFaAmas+YUboKMdLJp{ zZSX5nq?3^^Ay4LEq!cZSt2y+hDYS0&EL~gxPibndpY~UC_aY)qEmvnuOC4Jo0QDEA z*pyo!Xnv6CcZkKfxf7sMFa+e&Sa!MTKkSM7M}dM&VN%aO<#+#tIMXW*FrW@Db>$AJ zGAArlGVSwFhXi@FPqBI6_|W4PaTA*=SXWn0j8BRnsf@c+Im)F(%i==^XQ1_I ztAFc(Gw2mI6k3Z0nCZU<-36w#QQGxbWs4P!xGyH4xtcmr=)DRSVL}*ZHy0RgHI;k& z>aJaHRk(7%x^IgtXjvswgK%*Gg0McACt0s&7;Yxr+82XgO>8N&>p+a{Fx2HttS@}^ z5~B>p&{lKu4*1t`=`AfV$4$jgd{fToNq0qqWqHqt;&PFnVLtM!+JO@-_==%DTC;uI zAZfCVF_l|6F?Nnp2Z!lXi4#pZO&S4jlB$7~xHPiM=lca~XUPtgZsZsxCCjB((hC9~ zRi0U0{HtMf!hCrGPjZEvn<6W3`Td>skiOHG$PUR;vl9)-KvewU12bIqT7?HU|Zev1%%{dscr>#c$+qKm|O!&^k!$I``9JEW{+dkdem)?{j)w@9tAK)LiFSBQ-`$O03@1{n& z^1e^ss?Lo{md%ZCIn(=kDJjJ^1$Qh8!7yOX)Dlo8)#`HX1f){G~v-wV$uO@Q!TZvo)h@H7xp* z8w|OQa7*W{{Avd`V9uWhfQC}K#qYmia=iD>=19=OU`*s*1}AhKcW1(kH+Zm5w)Zs$;-mN?J1#dtI>)qixhaM%o51T* z)GMz)-QSlRuf#J9TeZkc;6Cs?bWit)yq$%{T~4z^;O%8YrrcE_o^?n|)=Au5fNb(K zSu7s5FIb>=h!HQ-K6bi}qj#+`24xg4{%UmCg=`Mk9XPvT(lD-k34d`VT%(Lm?-!Zj zYPuZTvG>FFI}O>h#;46AUPyz3ip&*Uv~59wbBhiT-Qce3>T{h1FP|BxEZwLbFe^$N z!&R7Jckg8s+H4`m%#$yUaUfWp+~?4#EzlPldQ02AVyCzJ?8_So%e>+UQxkxWuRxz4 zDKz;%#O0imK)af>_Wj835+PM;!{QjJ-y8gEO?8t_dlY112sn%W@M_KjfnG#pRje>2 zB6HQ}FW#m5{(Bi?p=tejnkvQK3t0HKFUv-z`Qz98kS&6WV9!dasEqMK%O z0xp#29RK^-?#*V~^1i>l>FhKD?k?+VwK|}ttbA<$<^B8i4eWgJ?%A`e2WOjPuhN_N zIJ(48FuBxinpCG9&eBBUnj6AdUFqp=KH+p0qF%~@&WjXUCF^z?lqe?Q7m!~iM3>*K zq{}YVat)dt+II}2P^!=3nVh3DSR0)3dv({N=^Y0^Ol%B4u(LuRY)e#%I))n{>Fwvz{~8 zKu|z6oA9#-fufdRuU%BGFv-HV=?l`>yy@s*3yH6DZ;Zinq`*iB_9vx#CW3%IKzET^ zD_xPSci>7AdqV44Ty;#DHro$>#QF+(IH00n*COGg4Q3-N@Ut#z$+*>?A)&0BZwJFb6s-{#G{3$jFjSPK04$mXcwLo?VgrY)vg zocB@ynSC_*jcRI9(J|1ieCg9A)~$XetSTl}A*C=r!XW!}=ql-AH4=6X)d@d8{7Nx` z>m|fUi?D$}&wnJ{8czV5#fj`@orN<#e+y?VB4y@y24oF9P!5^`{|z<5q3}zfe_Ixyjr?KhTG%xqWBALpbUJS13K?~;Fy&|WQA@)naZ$3o^uPW z&flv)7eR}z>`ZkW({A+v}=uAvPkuE%u;JZ+DrovBcTSUZnRqN;O zuFedH60-QFp{qyKt6`P+YrAr=b2H zF>g4t2t{r8kC>oICLuyG;qB;KQu$>{3<%!H%}cZ(o=i~eq7VWe`;Epxdswd;pb@9V zyFRxb&_r0lx)G_ic^t|9hUSL3oAa(u-Ob(}{IW??2R*fR4`PW-#kY}*vWdtLuL)hOrq^PUPG;wY+h4)HrMrB( zX-GyC!bv@*EKPAMrxPrNiWZVjdqH`ng1Kc6_6Sn$MK1g`(PR*y{`}@Q7U`C5LMwGP zsq^R5qvi1S5IbvnTX##xOwL*<__`eF(yGnSSXo9@{&npH-8c9vGWh{ha++Njk7WhR zvGXXmg$(>OyUF#;b90-w{VrV>_u))UIyd~R8SpbQ0Dd~0w5Lw+d>e<_EK3f8B;Fd> zQzpXuXK85e6+YY|PLO62SV|2ViLZ8IB4_dgX|u^n=KBS78~@^UK`sHIByrB~jfTS= zg(&4=fr7yt`}!}J1Ne-$0W7~bhqenyDGxu1?|*F)1$(>&jRTD$Qg=oOk2O>>2UeYQ zr>)XL1D*>*omEGq$)>YqR09EOo^l>Z6P$s6s=Q>)du5}%_a0g*^Y0&eYp-946!IlO zXW%ZZB5LIlVNHKM2OTGD?X2315~N*1BGqjBHNSkOE$`^)8_%sA^>c?{DV~Evtm!kUaO~ulcAh|HM*{A3^<3?4<@ak}pp?ahF2t!l?+w*o z2uC{G% z7YUNNval`4!7bwXLgjRfV6{z158?Y`aslnbM$g7EaaC-Aog-NxONdGK9a<^Qh zibJgZrH4muib;7=Q+G~Iz~HN<{)zrm3|p*cUzpPS&C6;v_93_YUb-fH$~OvKl`#)% z_wTHEN$3N?-|j|2YFP zNsJyru0pOiAzkBb|CcLRO5YsoIy`lMYWo?}r3|y0b*AOuRc1Xgt_(!?e`{`%^@_1S zp1;1bN)m3Rx|~@tD{Q$Jko2-A>0b|Wh=OgUQ6hA~Y17wbq+R;Q%wAltG!I>7N z#l)q~8S&+rEiS8WZ!KQ-kSKfP+MMf+o?|=O$eoT?gy@pBc}9C^Uj6qwVn>IkEgi}h zluQy%yKZ~-0Bo=c>&&mj4MB^p?;G>X0jb<*y~RbDUpAZ(@0TQ}Um7^&5~omn6N$=0 zD{-zv3Mq;2jEW}D-{ku8uJf){2yOEOgLmA9+s4&ry`y#bXfK#{wjiF5Sosalt(N(G zI-t`3hL5c4zu4St0=a`@kHebj=qB&M{dq=)$Zq>AZK?I!?ES4*Daqh-k=V=*o6*0) zU?}j3N5Yj;#pR1MFz5cDoY%N*!z%IG4$jh@130;*Q8OE-y65%gw{XDs6c+Np{4|g1 zShIvzVYpFLn9e|?4sua6G$=jRz@|tLNm&=iEcHwgSC{$cTzj8WzQwK7r6NKE5v^i# zf#OHT7JxMJosz>G1b%m8S>DOu7ar&K!U8fzrO6HxO{Q7PoulK+8szsCZC?T>Lr-bZ z+e_r7rAfsGShOe1!Yn)@>vh~+*7|wHEccln&lRtuP%(16?M4Z*;Q~Cs16bzfw!dqV zXWLwDa~y&>f4pp$)-0;Yu?oDZ8A+5Zog^_p;<5%;la7_V%AFP4!^rEbox0tK3MFhf-*$Q zb?Ti;=EZCBlGJ+U%6Df5EV)NJ@X?I@7T!)%smls-eHW%4vwjBlDJc}`hGv-ro4RAY z?97}7*7l=u=6up^v+!}PKygvh$BeA;pTY#BonwIZ{*SGdlfGbC1`+Uqcs>SlX9u7@_&JH}8&NCVQ zOEozZ8|i|H(48xBhSPB7obebLPhSo(uf-@pM^2$U->@sG&)C{`?0s_3th2Lvqk!Jl zQvK_!9)9ZDvyxMnPIa;P1Jq@h%#N;Ld)xXe4Xblo`+#3X`8{W^VXA3b_)_o+;v1668=Y)f zJo32#&vEf?J5=WevatvbDAkJ3tBI8sZBEwKx>ha3ovH8V-3dOGQVh$$%Yib>295a5 zSA`OV&_mg{mf++qqt#5JWh2JPye^qhJCe+X*wAixTt5BesP(`&&W;$LP++#%E&vLKdv{ zOVH+0tzI}~5>mVmOO!RIP|4xu*1RGQ?@Dv_GPwNlg@DE^AEP(VnHujBb(y;FBGvXX zT9gc4k>S+KW~)BM{=BzSGnbDVA5c>WhwGsRHI#^)cu&8=++arY8U zi(tL)9yS1R4fuTgjvV4o;bh})YK6DWFQyG-&1w-TeFWS2K}=pP<*E2Q2x(RlDhB!_ z{C=L*hOii}jey$OQY(R6X`JIAvlc5%P6+d__>aq5`Tgd2ETc+#s*GBEwezTwGQWaU&;)XqZ^j0EQci{}C2gDE^2ojg6F zE>PneFZGZp%Pv;#dWJnmM^wU7ve7eO5A4+MTW=@B6eya&#tOeqCn{Iu-Krxbbq(L? zmYUadm3bZy;+fmaum;j~r}0hLfr-vC8w>4%L(q=~(^sj$C!w7B$CeS}p)c@Bc-*J_ zd+?^AqJv3@1$l!gyC5h2k2PX9VMcMx@-ZZw7w%5eiCk}!lR#>p3n&`c^;T>Lhuy+j zyHOwTV$y00=bJgFxRQ&ppq8dpa;nyIL&n@$)bsJ;mqtxZd0k(jJhnYW`LM+9=Jgk@ z+_6b>xzg1)%ib74f*fz5WSQeV^pfj%#`MOy6yK63;(3p_buUfM*d3Hmf~3yYuGD)7 zyu%oHX3iEKD%IGXs4;Dw=;!tUA1^OOA#_#}10^k&UdLmD`psmv#z(etR?Fif!JRh5 z{7P(Lr)?uYE%fwy<4@BpGm55)#FDSS89@9U$B{0(}#r_=G>i5 zDd7cLsOt@`$AOuezTDy-o=llh*W-hE->Iotd_~lEOQSjlSB^JNhIphBid6Cj6ZA3! zV0OiF_29{WPGw*}pwiQ4mi9YF95-SFT5SkXUiEc2BRz0lapaQ2vQ66DR#NY*K#F0mAgZ3~AA8YuC~b;&o4G}%&eZC;{* z5DDz7paen1XvZY}6`8ZbOK-?$d=$@4OGQuLtz4*IRJQHR^#64G*t_<>RFw~!SVFo& z4&?nb{uYS2I=cgOrpOxI)5AXI{RTZ)DV?ykJ_4(Y?$GTYdh=@c7U1)D)YfCR1(oCR za=I$P-(TLq_&lHe;SAxt&S*B8^~-^W!F-x#9-nekTU8G7~ZB6oHzH-!c6dX(jS{FI8#(4^b4B+rk2gXEUGzO z;aAb{>CK_n-cfu5?G$wPQBg(-+9G$-tF|^!Zy^5WX=NLf`~DeAzd${m$RA8ZGwiPU zQaaJ6wA6fr9m$SfP2s}-Qr^v8PF|UuPegnvM7kv4e%}1X{%-s&Qui9v9uMo+0xAI@ z)GfG$jJN9-f{mWv997O}AiH!&WHGx4#rb#EXb9M=!Z&5+bLR~j%HtO`>@D4dT>_~+ z$M%~qU?VH8Tdxcce{7d+HD9?(w%VO1Kb0at;KP(oYiBs{pxM(62KR3bc?>?6Gss9o zfMBx0@!C!K(?6&KyBN*^^Y`*<^OC<#tSk@UO(PXBt`PZB%gD9?`ijLIZL~-GZP?!X zBakeXUbJ{F2n}7z0Cgf54Fj`l2e@V0D0Zg5%;VRkZ}Y_kl)9CEX=yd;;B&&`5bxL3 zKjYP6rsCCv3G-Gez`YU=zrlUKk{!^x`gQ#dfgYc-zgn>Ld`hAd#_3qMUyPA5yv2+h zj58n6P&$0WOfp;+HFI1vp#zqxMG3_Y)|u0h41fB_-~WWpnSJ|als+sT{C!?$pM3J# z!*I4ESs`mG-}qZWDB|`u;jUC~Zh`papY9ioAzPjHzf=Q+(IAR?UEy5eWn3s?A^_DP zTQ@r*K0uxf!<>chW0%;iNbc#WTJUX)DWycJBplGzG9_~FkCkExXl*}vgTgl1sepk|gylH1fM3xy=dBhv#u61Vp2+F~% zo-?{1YY!~JYZBDwMmNwoo;zgeSZ`&3S}#<~ud&tfNtLyf!Ig&z#pQdN2{=_L*+ZJA^*C6xL_(T-wPgpI<2*tfP4kN6MxajK%E!ryzTC zc;5_e*qGk>s~OkV{guHOA9XGYBBv-v`lZ@cpI77#Sa;scA!Fn@)#5X0*PyzYTX_3R zkHppS8q94XS6Cy3HVGRT_J+Gf1415%lF%NOr7Tw=v}9 z>9{_W-IehoZfq3ft8N0m<^tvwu%%?CHjs|Y-w<4W#eREnab0IZNU+ypPWT1*Y zm?#T?nR8dI@IlG@Nx&68xwfMUT_44lW(Tq=aj82AXmhutA6}#6T<6}IJFzpY+6(g6 zwStfEZad1PBry#Lp@Ud+8`hINX!<%OgnoeT_f4;fhni!Q}y(y$LtuzpU3AL$3WC9n{AdUDfPkn*`Zr1ume$DoRB+ zzNfFEG`q+rV>r@NkJcmn+eM8%hR_v!rd-p6lnb#Lc{t8=vda=v`4rd8TKH5Or2Ddw{Q6WvM6C!(_~VW=HG!I z8(N)=^D76ILr2C7nCh-^a^YJFxe2^<}DAkDs|VkZpedbT!dE|@HG=& z+s(PB_}3<^Db;m-={d-Tt|fVJl@;Q-fkzBx5_F5K^7zA{hsNraWL1td@a|y@zw%db zg@^L68Hm?U)e4*y9Ir47a_7VtT|d_JSyxpO4m2Fy+)Y4srJ zMJ3a*q?G7ayl->uUw?Ei&#YB{i(R>Pa>+)=?)lL}6Iy}gKw80YgV{~DCGEoLbZ*|z zpuWJ6UAH=iT&0N1cuZZ<6d&~vJm;U|S8Jn9W^JKjaPH-r)_Ha1P7|Z&M^F3oB($kR z9``S0H|4j^EbM$jMO=6??|UtJ#Bxj@(=C|v5>VKYRxmpyxxg(vt*kO9K_CATTNn6r z0*z8o`8nQ2K|m&>0$S=^0tf$_>i^%WBv!n*zT^$qXD4q0QBeu)P&oZI*7C&qY0YN| zXJLGn`28!g{Xw&Vx30hZPu!~3gb=;MF*SkS>^De$?x|df^u%=9E*p!L-Nf9bSZ=FXwLkTD630# z@~ECe-|RG}t+w!IYJ#Tji|ECdZR7Oqk}9JnR zV64#Kl(i2;i-xIEZ)@g4<2GCvTr7sJ$GLHS5w`4Vk5O+>vQ%yrKCce zPbB=^P*lletU5hrY2k2rn#=DNR_zgHFVCevKS$Wi+gY#PC3XKxQ=I#RaoPCU9ULhN zh($56za97Oh43Bj&5W0RqY4Q+@mOo4x*xSVE{pJ#0r7||;0jF!7N3u`MT5B=tPqB+ zSv|ETSr9qzFn7LMGbLQyGBb)lmswEBg2(G}JKzaI2Hpy}XW8~L0^_WzXEIQi5Ae{% z9)u9O&*(v-ApxZimb5H{X`zTKELd9Ec8<)?bf=MHO_EOo3-Zw`{K?W#*Eelx3#{87 zHBCSZZuS@z_-iOq0Mv~>IiOBL>p69Ag!&fVpt_MeQmN8+fL0wk;+Bci8C-bfp7pCh z-}jmZ3_pfp+a$ohvvA@@(Gkq0h%=gU#0wojL2d6Ay@eM9JExabtZS$505M$3qD)^u ziZTyXDj6V?=OS`QWL?}f&%V%!I=sn+AM7s*Q{KrnGxfyM=-6A&tJm}jv`^dwaK0D&YVFABD zmA_1?m1RR^1pBaD+kq{1a%gr0#FX49Z)hGQt{k^QU@8=A&)BF)%wQv_PjxQx@X2QN z)0NBxPvC}22o2BPG?~FHi%~ithlG1`ljx z;CH_>a3*?_H=zCt&2*sht4OO07is+~;r+4KXTH2T4}A9b*YNNA)~%eDpetE`=oIka zzR}(4Z5-&gJYVNX25@z>wA*wyA&!(cnq*H0bm~a3CyZC|X7CybraruslF#@e&S{g) zE<@Cb%p}yJ8|((tQE8sOZ5^Ra^kX2BWXS!hmtVdrUt8sM>1^Y$4NY!4PxCE)#VCDI zRL;AZ(v{`VuUm6;M!+mVT<@ZW*W4cYi{YI9+{))|z7F#NB&M+*sEFqpnTZICJHZdQ zRBZNE&g5Y_I}%b}&6xsUoo%*l7c`M;fB*$A$H|TPI|Fj87sC786|~SPsYycb15?dU zkI|3M?PADJX+^<^osLyUXP}klu+nhB@daz%Sd2aflWimQyMPthxsxGkwf6iU&$_rV zNRuy1#(QG4+B#mc&4Il2E_i@WMsPuFJ++2#_ zm^`XEy^X!33PruLor@GO8_erhKYsKU>Y=I{p1o=2G&W`huZcy`^ub@qJEPoe-Bo%q zS_zDoE|MdsfkW%ll+toXAYQX47J0@VLS;Q-@jd7?d0`(o?|c6pVFsvH7ey@litA3Z zbXf8qRsV_~)(7Lh9m*0LIW1p`#BUPVk~kVdkXb&lpbXS`@JdYWu*Fu+;NjwfGvI*W zY`c3HW?29b&91;Lm4xx4e8YH<2O#8&D_khB63jz^i1Wn zz1+U zfLh~^?~Ow|c$)WWs(5dj9uuI#=}0r4yU>K{Er?P4V!c$euAH+(2=GRS3x5tD=c9CA z0a2r6USjw`bdy6GBn#_p86R%7nb%NeJUc3N$};Y##DtpHJmFx%_2%vzJ|4IdXKD%p z2&*)?l14`e7hcrJOVY-Hms_M*!R;~nq6=gp>HB~B@3yO zRB(u!+822PeCx4^{gOHyeMo{X0z*NABj9IIN6J6)nUloVGpUY1A9Rl|iM;OK?PeNK zz3LPDaAHAhS++&{G(DyQKw@-Vva(uqzQipics5dFzKw{^C_Q+uk`!{U_;o$|8Tc66 zdyWqMYCe^LU{jO@I3aFn>%(K@FzQoQe}Y9HW|x+%l}_mQkOe7rs!qRkH+ zE)QN}&#W32RL;M5!ID#KS`N&RLtqecC;gSH^Q=b>d=xo7yl1&m$U<-PU-$T~ppEvf| zHA#i1V0y02GK>9Q3mOpF5pVBxB2sPiP#O}J7OYpJazA>TI_G+x*eHcrLTqW4^lJ%% z6@*XSJ&w1~*M!e0Lz&I?1LwKDI_FlYE`dAjifQ_Sw$P)|3M@J8;s#izS&0()$ubpM za40JaVoZ!jVfU{B`AK6`8I?v!cZ5K6u#xP;(rP;_Wemiy?7fcf)+G+C)=7?w2? z;@AI#M#oD2o~1<9$sF%Rh7(4Q{&MXogUuTQ5&*ago5lNK)#N1lj~|B(lPUYr4XF#L z_pDZDIlCY;2xV&4$N7+7|2fMoyu){*3yf;b{O?Q?7s_Er z{`1hw0 zr0IbOP{VxdQcW7kOy;%jQd8Z*TMK8tTMYk@rgiOOFvVYDOAMu_*1(-v^NI>Z!%bR& zWhi|Nf3U=J#23C*pY}S=69<!v9FSqD3c!LJ7#*|WBr zd0H1PEVj<99%-{LEUAP+cC16Q6qcVM#muy3xA3gpssN{KQN0K!==>cREkUhAX1aNv z*N$P&c7ze*N ycwPKn#B5ask0>-zsS1{!g;n}KFnLaRp~AYt$V7NS{*&3V4sX< zh7Lt19gsSn5b0*ry)fubEbloAT!ZLpBM5gOWbB3XnXuW*7L-QgxuHm3k-G#h}#I-li- zSrF6+b<+c#(vLp*C##<=0rH(=)5o6JUk@18cY(Q5vCkWPbLiIKtmPKrsUo^5y5b51 z2RDY`LL=Gg$|!y?|2?UaV_qfs`Oek=X%`?HOHZR1?c(!is_%7JTLxa|dS$~_!lzaX z-ya2qVPnYBZ3yl``L+Xxl{+sQ)lK|tY-VGwJ5=hnUMX=I+UO6_rrgX2 zy5Cyx`opA*vEu$}=PVT;AW=5}EUVhy6qiF^7Nyy%4=TRALRj)tQRW_9_2!ggdyVa) z_-pc-U-bp^Ez5q2bt`vmr=;QOAf<3bmA%_g(agP?FC#8A<_hoDuq4;gI0qcQh38^n zG`m;L^}eQ-KyW5?ui1YGkS(W=_kxDxsg1uL=p(neUdq9Aa?M6&T&86BKkR_H{f-r7 z6Xz_y8~nu}wj08GsdZBGa+g%GPiLMz_qIJ#IJ91nmR1$i4xcLO+!HN1 z>+m$Qt!2_v!9bbC!yucdgF;V+KDL&Y9}-UB2d$=LyY%KJR; z`M#O+=lpOc!z4RlueI)S-F2v((L>^d{PS9)3j+b;odH`b!Xv&1npX1So}|) zsTvThO7E2QfTxl~U;Zlm4<#}Aukaqo2X)lrAN;iXq*r1$ty_ed|y$E*%K zuA}Gu6g7J8J+xP>zEm3xE0V^MDSHrkXMej|AyI}MuqZ_w%cCtPsM#a(u7ep4VG;r7@ zAUNMfi{+~<#DRc_zg858HfZNtYO}l4o+!sWTAGv?JpZe_#tPUzRr_h)pNTkjKc15q zemw|6+{vodT$XO^vYc6>!@Q}GnZ7O2OO%kT(~^9+&4eA_nu>A>{6WGuUScdSGz*K-2$$@djcpw z_nBZnTq?eVs?M#`T~&$$mn$fSc7;E}u&@|W{BTYC2;`W-?esRl!xL>e0cruK-zSKf zHfzgg7th?6_g?(f^-;xo4S3uem819pxl+?ffLO;&8Rc8|rnil@-(PI$O*D|`p1ku& zj3k;&O`&~1I|sXW)XGKAqpw_D?{Uw8?11i#kvS-~0qrf{%bwoymG^#RS?xZPJwZN1 zMLn;Ydo_K_EE)n7eZEFRv}kVC@bjK3+ztj%(q-!bN76bqVE1)xDDS$2ehobz&oWz= zu;p~4(q5b7d!$*IsQ*(giI}ouxXxV#>%iS;24w|6!D}j}#(|30P_E!>k3JPHyA)|{ zz*`H4UGN7&mo+<=$xt`&dRcUt2s`8^C*qsy8Z0l@ej@C4?YkIyO@hs-ZArSJkpnE>HLGf-1OKVmFco;n*@dmms>L;B>sso8zSpH` z;oofbBr#I8EAMv(J)@_|5k_Sb7sE5RVbyO3xbsShNxApUqwBRNFp33ldTGPC#j3~% zb&>GhpqBQ8t6L9wO2yWhvTsJ9J!#>C;nvbhelbU?D&Pvi=wmB9EPZ^StbrZEy=Tpg%yNqCR;wg4(h+)E<%*?1sm~+1 z61m$M4ta?68eQbTJu=Cf!8#;O04~xvL@`;HBj(&qj#wXF&}ERmnmYYq!27Iz`6A$+ z*4fOxI`mdU$@4GO_UMx2%5b((>XY!IaLt5-F6>i`m~9W% z66j1IFEAMyLqz;j_tHQ2T%rAugUZ4x;_V<$Zn>NSbhFVw2+GT0Ik%46Pgi^rGoWjp z9=(H2DXMw|n4iff0BoKtLQ4)Nv5m{Nj(*gKC*K?)w|;l-9BtIxcR&?syg!^v=?;H@7r6nbMleR z!@DjWJQs>Xqhkc?;(Cu)UaYs}@OR?cuoCMIHyhq4P#<;KjZDv8G=N$pU)}eeMUCZO zX`jt4lOHik_TYlK5zNaO%FSgbA%H{WF@5(W#*q|B7Xv%T2<1|?4;m^yRu?Ijm#KYTM=qHk9HzANy+q?X+12F^anJU+o1_-Lt)FXi!+)k+T?zuq)(3d&N* z+yYwy;85geqIoCqU!fgGcXcyL*$SxDl~>=&QWeypXIxn^ia^h-5Z&8efgGiN%-$Fz zUbcuZr$6nmhKdpG>ymls?ekiNTAf~>3%T9~efOz(59Z3>9nMGi?qdja-QM2B#V9s2 zy6cyBfLVVuR-(SNEYlCo)L%OZcU72(=QL$Jf3CT7@!k~;%G!=Lk*9kyXr)tI(*3YD zro|&`ytM}5cSU08rmvCUs)_@@I;LQUzWPQ znOMJ$Kh~yipEXA;J7eyy*Z}!k6Mf5toR5U|ai$HS;*Zj%bj81X|2P#uBa+_0e#ZV= z$1S+J;Y{B#y!p@ft4wz06{*$vAiEAue6Abw`bz{lb;;)b@Qm3n4q4kF;-?FTikw$z zK5y`#^cQ4P8H-82E+1iGk^CmP$ju>vk&<886fHm-0NK6wxsZ;!ry2}Xz63jS6kAT| zV)H+2H3m*3E(^PdZa}YX!0J@1VmsttY32fh27%wjvbw z)wzLy^Q|kGqs#=eSQHH}4~RVNNT=o{n0V&+pLq|C6cey5lSLOwJe6e_z~Q6qt807` zGH=Hk-p-gEQdro9Sr2}>MxQ122&l0~-w{H#*u4niZ}J>t>Eg(I2wN(*TgxurIgXjeSpW~ROpin8E*E&w zHq@=URWeV>+H{CU?T)oaUOIn{(Z$0k-KTIPENbF{;*Z2+@Ypo^L+1rSs;~3Qv`|Bd z%L&Ui51G8_#!`1~1MT#$jbeIvZd=58{@93C!In3;80(uFDcDG4;m=%AUWJz(T0zpe zvS(6Idoa}Uu*W_4_z9}{mXpEf8DypSH%qrqHwB7VH>K>`C}<~&FQkMJKF{>21uq~J zvv$KM#9-JVjtLi&4W7%NXiPR0Hv=2ten|Jx>yYZStac$xi3}I?!Dzw!a3+*0t8A>;T$&gs@o?^hX?{tN>2+$>9O(i zmLWbVo_PHw?$_tTV|9?#^3Qfg3kDOCh0oa6a)s;QU8uh7H+*QI06Jn>3p#dtaGSoY zATVwbs*XWL%$_|f@G7i+ABY@N8h>!TkLpl$md(sjTb##XV?5$&@u&*|{ZgzTdU21%p7RAzHY)8F@w3fupYFJ1%Ue zhu_h+=ddX890AR!Jl$$e>C-5uOF%GBKxQ>=XX3|*5sEl;kV7aj0k9q7V{s6IwJJQk zVc(#rHAgKt%qS3vjLWonPjT}xMKz$2JSQ+4Ue4c5D*F)cTvszVlAsz2mUFw03&+)9 z?TZOo1yr;v8m`wRG2FOG|D9i989C6fM!mC6dk{@r;|r3SkOIl9DruVDuz4}9$X)*3 zW?|nsbe`L=!1vRREWr*RC<12iRdfKYVG(G5ytq?A$+)xh?F+*9UWu~A&MtdYa$ciUhX9gZ8)BAkTbdJmqPZGWW308TE#i^o(6Y^Jw zVF*{>RiDTdHmJ7+RWsnF-G{X-Id{Zd+{K=C)2ZlFj{+P9p#r=HfLy5i-U<626X+3B5 z3(s<`GQ%S2~h z^VzdGEmKu!{))L`XT=aME9F#J{dLGYz?3Jt*EkI0E|S9#!DCsR?uN0J&^7{FVTOP?!5(g?%K!9cO_2beKnhey#H9PnriF9-j*dP{ZUo za{rism1{GF}O2K5@nT{5nCYnIT)p+m=%|@D2zEj-(nRP z7glN?bnXnSwE@2}G;@9k-B_~sT}ZMd?0kn~^7&U?DFG%q(OO0z_~UiLrp6djwI-G& zrd5#Hr+PB_b5Ae5Hq!OhR$xC&%<_lZ^dMQZRx9HT*=uy_5wgR>it(L(`1?ru^tIJa zppBWC@U-sbhJ<|x#hvLYE)iaEyZEA*>#bbsl5;?d46p554=%F~t4S9}q8` zt|r^d`EknKkvJB!>~S>vDmA|vUoyGb%&Hn%Qe1Zd$hY5>_hGNwDT_*K=r2dO+(pF& zpMS^S*r<_SGrhr(6scvvhLGLd9k`QCGS!;tetBNi4%7o?CL7x&NAr)YZ&2SIU+c=x zpFM_#Fz(_iGPl-tRb6XBTHMEXr$xh7YXmSJz?tlYzAs*#XO;P=Q~uRv2?bN&SPc+lkxv)$Z#J^R`;KHUSiTzWt2-MoR|$h7yTo$17Uv2?F9&D(qG^s3^! zyoAuj${P$J8+cv|Ox)rIZ#P&63e;QDD3i{otD8n$2ya+|=++v2?Xce1i_iV)4fCiC z!CKY015WsO&kc|aYl4JyDk%QL*kP3eI4PHG8=IZYeJD9!Ns76E0*FEhdOv6aiNpaa}_>N}K9nkYvL)H^s|Y2AdD$`L%op?`nSrZbZTp#Rpy$YW4L*liIv#ty%RdlM8`RN# zxr8WR#xf~r(bn?r3cssE=&j|9ueppTeJYnlQAH^WWh*|ZDXQ0230O$s)dismw9@Ww z1za}uBE7_A{US`gnB;a?K^4Jj8%@wI~W> z$UssS2vQ*QCm7Sy@4XP%S_`S%dgPO`c-!i|FWEMmT73MNg;^@1OanGK{hAIrcoIl` zYk6~TNZuCx+N903QTTPf#RTuUJ&%p_J+)1(^{(9yB3H9G?fGJ#ocfOSrhbh~=(s9% zD`7WS(LcJ@6K`GsEnut~LVZ*Rr-Ef|rmDPs zs;(|nt>^h;^I`#bbIQS~dG%DEaIazUzAzulel}AD@J`}NEjGpbt`_$}_ptGZI*u$({tVnL1+^+$gTX`u7qdpPLcP;9`^Y<%L3k} zdw!`P%RxXdajFBzYvC_p10`-Z;iDvyR!eAG*l{?R#%p{+nG@{hj5DScmg&+I-ZlhZ z)=p(iziKwK73>-OqpZ!v53;<=0@YOxK@9RM^{?7_Ds1SSyOufUvw)bn?c<*Caq#&V zF^X4!hzf$jxWamf?nHwx>%uMI)N0?cb!~ zC4@4!lVI<>o?tVzIb)7_$lHq<$yTcYQve55GRz0caP_RWi5IJLP4@*{?_T?&`Si>| zvd4H2=$vYVzD{B~FzoXoQnAMl-2Eb;?)`mOJPD)hD>Ct3=FTvGh9T%p4v(5fEm8qQ z0L&-U@~$9ucA-dn#G_&DBCfwc3N&iu=^wp=zIQug=RS53RTe1)lH@}y75ih4taU9p z#if9T!#G6Jb9a9%9{ajj9z>v;S6vp^83Yo7&Co2Dsqz+l1_)MwtJ&Y%e|X#X2~>~E z!5k?4I%=IiS96p1(FL1AxXbzZUAN<=t^S6SLW0b;M2#P5;i9dE@bSthwn^`ighqeQSi)j*j44I~@R>Lddo` z!#!%p797-pR5LlYLDoY0tI^KwRd7>2OqoTh{+&axu%IRsnq%RUH3g# zY1O2V5e8$p7` zz-+ix#4H~u{2TQ9&T-^Ln@q>fu2`f>vRxO?`8DtKs3{3p!OqAI zjJUIf(hI7f$TPN+oM{hEg?$c}f|`UQ_a71K5p~?KolXsa+|f0-3&iTJUnL7r@zNg} z8#(;gOk#LyJH>-5L!YOWdBFJ&HfOR)6_fGre9Nxa5~fNG#b~Rh)~0kfCth zY@|)^8UO)6Dm1giEV-)&(bpaWJp231mR1{MB9bD0Xj~MsJq(?0d!jss$fw`oxPlyml${K z>nsUt#TQov-RnhAf&;ls7xBK&X;VtFDLA z^?mYLMsRohO)sxAGQ$RHWS0I&vwEfiDj6O(@Zm!u&`L4k^GLqa%k_s|-S<5I1WzWJ zbBr4puNGm=l?SoklD@eWFu?p*qc66&_8sR)im`kV6t7ohV0r`Pw#Fs)g%%D!>azOr z4y_^pc!v#S`l*#btm=ZYQ#uAgQ*3QbQ&V@)lo%JR;4c-VjnV1Uz65r4-Hq{pLzN*C z!YK8NFh)+1m?2+@g3>w}AY&NTk@VPmfs6kmU?p0>aF5#8DqXZSHwbq14fw`gE&2{P zdUS8viTu#32Tr2`G9Hj8PY#PwX9}g{J;s|CK#5tRF!Yc82Cs8s*ZadFYQsUj&BQ0cT7+E6#rOD$9pCQ4++nP!QllDVv~hYOPYPP zT@*W8&V$ZILwH--@bTE*S+=Z8f)89IzAD=fv@i+I+~zq(jZEtIG|!Dyy|wQb}+X1RoLi3ySDc zW80uz=E5@f8r5KAk7z*w&{Kkl-T&le7}TN@yBfvc4zJ?v2yfhfN4e&iwWuLF__zo$ zO(5;DcE{V?gK$l{p&#E~Lhy8}F3?sAYUpR!_~E~K%A727-Lf$nmz}P)2gSU8HB?Kx z)Vyv% zAJQ_OmWY(Ecri@A0$WY3%Q>_*({XDnU2tKxc_tBARCLTJwC^(VJ{z4D_g z@Ed!<#-J~s_}&rTPRMu%XjU|7i%*)Jo<@U3zQL9-iyI*Ajw-KY{s7+>Sk64KVKsUe zo0UHuiFUn7g3-CooNqo(V63REg0~3pxzE|qB)M2%JOxFyAAF>NW639|w>-Jtt|e@Y z3OIoi2{c{1Em#M`r{K^qmW@%`ByuW}^y7;!>H}BAe6v1y)EPjrPqflKU?oa*Yt3eS z3^p61SHcTCR4*VLD!BsDXq>S)wCiQj4^d64t80w-buYfT%EvQ~Gzi94I45w7KtS3C zP3X`|vHT?ZwJR2lsHdUR$KMMoyNHJwgG5umde?%+wYO}Ver!2q#qWFQ&*O|BhBTfC z0fIMVjd~6Ht0B8y35by_)ev=uA=6L5K4_T`Wq!!k4f=uNm~eNtJU8L@tE>m4ZumX` zK*3b|S>>k`{oH*>OzNa2OC?FchFpi1=;KEZ>!b&1?hZ{r?hI--yAloy$$$A5VqE{5 zxc_A?qFg{-a=ppX>TmZw7|9=~9#M3SCjUxL0l#1eF2#yuKlsaVhg9O5Arb*ruM+;p zsP$aHHfQ*kkhwSiBqxR!7!{U%`7hthD+Buj6H@a2a_J{oUfu^{iW}6u)_)S6<~l$n zN<*crUs>)yDcBA4(jdmWV*aw866R~&4lSBFf4TJYR>exdt=oCx8V|!2;t%*v61~vx z*To+EGz5Pv=!o9}K-oUp?7M$1zup+QrC~gOU5RsPO53-unvBF%4fo{@MF4Up1d1+ir=d@cVPVKLUww%-MKb4I&I$lKfVZU!L1@StwNM2c`>5Uv zTWY|L-*~ah-oMT#{1Z56c-QyRiQVlALdhw;4q@^OF(EBOj)w*OxFsNiTM40fiU$&` zz$r**Km1XVZm63=2(M<+j2T77_Sb1?uP^^B2JtD9UN+>SsCxysNmg1YlvTsv97r+6 zGrb~8L?1ZrD3IDdlRCyi#QhT9^=+Co!f5p1u)vy4&-fS?AH7Cd;B8>}iB(c4KaK{XE+aBg`FKS%?o<^Yq}F%+hatad-^^nbGPTK7r6w?cB((@9$$!JaRhJUO+F@rQP=9m{1ULV__n zP`Q!z)J?iXq&V63#L0gE#RO2vaNnvDzDXQ?J`+gsCk!<`BJO{b3I_f zpR9n&0^&lh1+n(=dxYOq`a|M>%hGRs_>CGwAU(v&-}>E|MsAC|bvBI(>xs04`IK|Gu~a zVG6rBR68)Fu#{CYlB2F?K7BJ6w$c`zzxPOi*a@oE&b=2Lfuk6Y=mj?Qjy?EX@E0Th z%eMv2fpLIqwG$V|=K~n$sp`zGSAg5+d=J>)W3JuqhkHpp$C<7h#2ttN8o3E0-WMH2 zg2muMC2LroH%HRwNoj1an$-^*hgMrBAl<541fAzP;spu)V#wNf5s9r#tZVgzk(F7s z-PP$?&tt`kWK1`h(x_IdUu@T<X3+~ePDu|MKdWu0MzKs=$6?H6{UGoi~gFA%fV>ygNrP^7*B)@wv)y&f<$Yna1(x7=L_` z)_jU2dmvsvPXl8kdGSjPl~b<`3O~ZkMAw6io^_fpE+VuwQR-9@R(WT;5fe`P#iocR zQlPK=2idw{AX$CI;P+KF>xl+cc1yLWhs%HGysG9q$F-TE%WN6 zZv%(r+#}p4Kgxm5>eLX0h%D%4K+y(diLz}tXT}##zc*2ytL<@`i*iGB?WfU**2!Pw z16J}MFTZ!EhSR{#C-*iFSN-}ru~|QX*0ukG zd~1&v$Wwtc1OJvhPjlaaaO*{*!^@`|8f4nMEG?M4e<3cB;`MRFWp#8{nI25T?J^-Jg85R#PD0>>UEYR?a*ZEHccZ1)Xw8mw)~>Qvk$;ziFvW`GBP-LZ8uN zjt!>O(H{({MdBP|?TU|~)I!jId8b;Qh~WrX8j7|jG*X7u(tqhORAwz}*NMU@b z|G||*Ci!hJ0|eSD+vr7szRdT$9{WMc%D7tCXCpAd@FSG?E%Rqo^JJuU)#sR_b*9j;QD6_9}_Z@8V% z!94a}9SOr`3zjNCoyigi24Ofc={SDXluj#pr>VyPH_Ei;VAy+05;;4D|1RD2q|LDa zai-hHv;`0DXmtnRPr?M${*ZrCGS_at_%N9Q++Pq`eookz@5M3Rm+JrzG%)}V*f9T< z2h;bijXK0QoHZHzLKz05{bEDtZ~G|Bb9Vi~&zYp&C0*`zcxYLfeYmuVQyqp3UlzUZ z>97P!s^C01^^rd7Dr8~=P0nrbxuZz@tn;phg~K)eLZJ`0>5wKNdm;|~^L6n#z}JiKX~)S93G$Nivowr}e%9jM0l;p{7;jWj-Vwk2 zWnYsi5p@2`#st$cl;efqN$;E%*@S-sVPH1fL}9_rE{qzM>gEsGLSu`2Kt=qr4 zFWJ5=u=`2zi%~=>7lX%d9#Wb12A`(_nYv{FpS2SiqtUseviMR8%>@)~?r_$BBkDM% z-<7R<_1SQ(AoZ-kr;kGTpU#1|wSt7c4zL~m&ly%gDNh!IeYO7^8Gl~qV>+djJUYti z7@UZXDq)`g)c$r@&djXh5XJx;p>e4E)>NXl#iNj2i*7`qwr7-{1dMoc|ZaSq^$& z+&e?bs(RK>`w#cb?+Ix>l;%sIw(BzLNelKP{=(<|UT|`L=Yhr;3zO+UvddGrVg2n0 z?L*heiuGtaZ|Aa5`6#Ixz{0Fs`;(+HcaNAmwQW;cbdZ}YhCb5xuKtse#^1;cVvi9p zw?(#i1D-B~8Wo(m^6b>p>KHJq6kElqLkR=?b94rNwP;@r%2;vy@LDJFnl5g3@Ljze z5&D9{2o@xUWM99nk0{>L@zEzMrgjx+fc_zGWJ5C@Fjh8vCB;2qygeDU>eV(6U?~F0 zCBEE;7UnJ!aL38yl6^H^D0t#vd8zTFT(Dcki18;E*#Wl{a)kDZ_F}*MolO9GHv}c1hui$M!r*t3-Y|2B2rh$(OxM}}bUkrG0 zb}i>n?`}msn~k;DNzk_=F$Y$|oBji?lzQE-weMb>B-@Jt_rT991^09cOMe7O9=Kfq z^QL;uk`_Nc6*Qj!>rtE=S2bM7aT&# zJNf!e#e*2nPXORg?4u1!)`t<#sM>;4ed2xs)D60JZ#VzC+6!8V@IiJU>Vuv zSyrKMjFTVccRNidt)l_WcPnNO~iOuuBAhu3g zE?HaPPpc+%ul}?-N&3t381R#4tBQvn*R5oLcwW`rn~JsjzOIMNv>=LU2gEr?Ih7cr zm}~6EpVGKt98h$fozdc_O*wUxGx}8M)k(eI1mpAjO|V11>9>Ua&-nLSB@aBi*S|6B zH|YPilE2-G|F7MO#Xggvd@zMJYHJEbnZKmTn_=!0cZpT?v`O2C!KKHwJ8Jnx)eJn~ z>dX1_B#f*>@t?dbKK-!9q&3Uy*3x-DKyHpI2DvKt(B?;!3rz~SQlrW_Q{HQjeo-Cq z)Es+QIZR}=FVGSG-y)FKb7F7!-@hG3?vaUW7K&bz%W|Y&3c~q5z`4jH z#6}}7FLLLn3H7n5o>RYlqFsTlj?OTH9=-Q^cfh(PwzMHdlMUd4XilMU?`FHsLwc;} z7E+_ebO1SE1rLX7&-E*M$!)czi%RB)Aan7Z;=<8o2WeqA4lr6o!Y3I(%)u0E8q6G z)-q8xJ{jv2UvUr%2oYle(4_X-jsbjyztDH1vc0B@uEjCH_I(SdxT3+;XaQiJQaml*orJ0qSCu~=_5&IZM$F9|=={Fod z5b^#w(e_)LRMSWBAQ7`E@aqUc$rFXnV^1KXoK z$z$IOyQ&Dp31-z@XNCM!ffOgJ6_!M*BQ4?kM86$U@@ zoFh9|U*wEl-XEZ7ENs_qT;QzIni_U6v|5ekFd68uHo@N#p(4*%nLZ&j>bluBa8Gzm z!+iwF?q1mGoT~)|S`Q?} zpw+WpBv>j_koPm1v)okqC3ue~j%tr^Q>8RH-UWTY$y}3JPKao>Ty|&{LR33r7&6SS zA}n9K*10QYsh%xvs?2rypqmR0OnwzZjlH^U|*DRQIQnPi)a{CHv(GHz) z)<7z?GIp@_IMNarz-HL0&Vr3SAEVS&fy%T$g>afqK`id1wHcL2@sL}K+x1}uXw4;M zT$T!Jn9Ez|mJln;S$^NBur|Hao9|rvOpv|PUE5uq_P&YH23D#+kF^*x(*K6Xt4? zJL$OdKHrrkG2*>@HunN=R8g7vYUtCbL=3xKWka}|BZEq-QJYI?Y?&HCSul(G9A?|+ z3^K6)^5|AFL}g);gWRR-YH7euF4w5W>b%m@O>c1N1pCX70h#AymzOM6yl;DjO0AYM zU@ZI=>qio9)>D7oe?h}gor|sWnZ0;!a-YGfrEQ^aW;yA?XR7V7NaS#3%yBc}2RQt-odQ#)}v0oirNd7!QYQGtkPz71vTYmoOXpvWl* zn@4_(+tk&oKd*k>5J|o5$~UcJTTQ81QF4+X}33n00jR# zDb~EV^&krYgI`7NKvkMzxO>G5?RP>=8MyT2>5;s1O-V?e0E$aKaWN`zu`?YsxHkcD zQI+^9#mLt;0^evpfbyJ%Yb3(Nx3o)A`ox{l?xL_38@HRr2G@h#}vPIywq49XSbUa-=()H<6`ycl@X$s&44v{e`FI&T)*{!M5M zZgdk9Zx$zp3_9z?x+UIl)tn`EtYWPQ4U4L|Tfe**uO5O}Cn{@;KLm&74E`3W?-fh34HD#|b z#E>k2fxgR?MMUKY-yKkuHyUxrJ2Cku_8Bx4Z&-B7iw}9)cUgf5X4OAV&r@u7_#$g& z=v%qtR4&W8-&eT@ZffXoz9J@_g${=I$cHasI$qCbCtgfJJX*D0SW-VdV!4uQebh4) zRyx-Y8p-=suXe|w$ql(>kmc4iJ-uXwB6vW`J)_*hH8wpZV5+O>!8jh63{O_!GZyaj zxB8nqZbfXFLkbAlN=mF!y!d%S@)$+4e)4K2^$gvc2sO`aeOQ*#ew_$0kCjz0u6A)= z2FOL*zu0%bNaF|Oe9nIlZAh7n6(mrw!2{_9WMFKKn+qKq)Vx8Xetpgn(KV%NJ9oYb zyCL^=olcY`)+W)t-`Jl}!$aJbZ9UY7M0ljf1In7b@}gS;XFgAeTFMlvL5(tXY=<@w z5>8cC;M|HZ!4^r#c^YQC3&GLG8TIJSb+vfuEo+yr>*?ntE$d5IczMgV@GL-d|*3u=OnYqlQWctJ6 z?oA^$sQI{Ltxp||3R(wdLVCs8phfAVLv{+DJ#)=E@uIl&BCLe+B{LM#quPFJ&Ow_o zb^}l|&GcgP&{Qa+p?JA4EEl>-G#!q!y-*-ahJ}~JUgd^u-?RBamjKE#Z6ADshk8uq z-V3Xbi?b96+M?dH z!t$vMo4t1=#D>hbylX77u&m5tw0Tt4+d zYSq_vVf@RL$dq=Nb7fxNK0Il^&TB z-=ce0o?{6KX+PySORJJ9bOW|FxOBHQgAFy&P#MuKGS5=tw*%3_oTgK&tgxzpD$^X~ zEl(5kmW*pT#UH24I~B*^hUdILaQ^V9Q9WuQ&ent7>hR!CeG(W(h&8hez9w_8$sv|u zZzi%bpbL5*b?u~y<@_yhhSy?9By>`}+Y>`HS#W!jt3sb7ROR5~H$;w0Y`uDZT&UOjr z30jLdhVs1f9xU$6v|RNudQXj#?~bWdb5KsKsp_rkcJhjUyW!`=18VHBfN9Q~zLx+_K`be2eA{OCjfN zXWt8p)_Z=2UD3sXHbZHgnQa^CGO5S$$$VFnPW(#&wJMm5Fywb}ie*4GK3~E3K31pn z1@vugr+h)Xq_a>^V@gn?p<1TUqNV(p+^CHN=i_I!in~7LixrBfyzbhEuRb6+Uh*c@ z(Li{PGaUDNH~PX!TsGM-dafjQ5VfWbr5t{m0MS{v9}^IJOQkMnqewc8Ora}jAOKmS z5bUuQPz5FdNpd#w1`j(;Fk7R0os4?+TL$hd;#8)(VFo23W;BcmE_k7t;W}+QCUvQN zvIsmZ7I?H_*(%!2(Nk{a`okmUcPMUD_BejoISLbT ziF@OggIO!k?il*ypUt==r1dm^D7YHC?dz!#_DmmzOaHnya@02ve7;@#)~ZpD`^t+w zzSRm#%$ovPo8{6vP^bj{C5bPrnOt0__3~MC(w^xTOD>1OLK>%4EpNSZ1s*3K6yKAC zEVK~3`rbcH$^UPFuu!?T_kN^ltIUI^3^}{!PS^hUp-D2KY$jo4#giW|0iA53!a2?k zF1?Ea_g4zDoGQAhR#=WS%Y2-?L<%ZuI!oXl+V4uN$%{i+zSKM|{5_uG*v&2MHzf6F z!2W$-Bue^CJ!`A5aLihNvFB?N@{u{#{F&A5eK->50jnGjU@EqSbX9)c%BeH5YCr3c zSWN?Q`{CnJPA}I$VmEVhPPc-IH>Opj(#PMdg4#1LD)DVOV_<#4B6ubrEYB*Q3b8M> zRO4T>@AFaoK1si*s%6*Ip;VY`V83*WfOOh-e683ffST3n>>*u}NOi1{DN+$N+Qf(; zw@!bq)!tl_b7J4Z_eFxTry#pf()XdXi$lZ3?qn*m_OBtvOD43D3nn!nTzBiKWP+U;Ge{lQb3j}yTx*C zi6YBIdB8t8SS1wq@{BO)J9Q8H=@}#y_Nd!4yLT<0b(e?kMfcl3xZ_vknWNdqw!-LH z{hon7JbbJ3^=OcjCs&G)v$tFg#HSO9+Lw#~wJ)F9@X;X$Gf|lMzj6z_-mU5^>@yHPBkH{ZO^BNSE+JV`XS$({piX;R)l94_#1-+R(wxbDEvw|MQ3asjyc zZUAhbcPjKgVCW;}_Z(Qe0{}zM^9-|(G+4jyS*0tpM5TA7Mx)s*(apzAO{N&

aX zYP>~W?g>WB{Eg4kDY`tDWUY5q;xb)7jD2LnzPo3v0?a5bhk@L2TbBL#JiIO|lYP(G+(*q{TWB$dYl z>&LtT3K79dyUo55EXOZAaQYCv+$7Heb9xY_IYoN9TQafNU@K0mld-4mv%i3vX^rLW zovg>#e4Q3*KHnY6*w&+PKT3Y40-CmSaYUu{Z25X)oT~T3WI4x}%WCWS=c@-Wv?m^h z)|;Q!>Pm8a9)XHmJ(Z48E1hL1tkGW?k$L8~s6^hiZ@74j4nxR3Z-RFP0lmV;o-lP0xc_zSO8hyz^0}}T0k17X79D@bz%=v1W@-d^oj_=9 zxw%Ig{rW2vUrZ3y?qiQsmmeXW0q2Xg8je$Z0$(XN&-UNiNZ_0?Xv}~To$&Bkt0zVK z0#~1)QlbF&AmmK|L$_cAIxzHpdoh|bf8XoTxv>=XXXGtU^)`|8=bA&jChPL3?BVnv zTBE2b?3>3&+HQ8f?P#$T^j9}je;T3Z23r(QMU#pgMF+Ydg`cNHCf2(>u(pWam)t}L zEL{WfgQ99@bH%@tdPOK_eN*~?S?$I(W0#3=puLze<3v4H72SGPBzACQ01EDv@QH3S zQt>V0PW^29x^m?+z837vUGyYAF{)iijEeG&=(e+cWzC>`SKFF#!Q3#lq}B0`>jf8H zD0n|j9B2IWghNw#DIe=h0A&OfUC`}|3A1S4ffR6;J*QphD= zFMcpGXw-BHuPFvpq)~2uKb?5c8v{Rx{KF}7yOaAT7D9Fv20bz+BaNG$As(3#n&RU{;L@NL( zG@ibng47ZoVmg6S{m<8q%-id_= zvjM;t1@88d0~Jv3iaW!B6b)oVf^3yNip{Cgv1WGR+uNp_q}bl%qBDw~u#q&z*0Vb~ zv$4chl8L9pzS0WJnERL|(UUw=qUc)rUEXDzUhRimhO*aCr2vCj$K#AO$dzZiR>#-v zMP~UwOp%UXKfWZN*KZPVzD=eMdDXgjM=s*pv3k{6tH&_K=u=wJg?gs~{HX{ux6}i; z^MymQM;bA$y}BnnPf#P%W^wq4!dP=L@v`K(W%w zoDnclHyt87}bjHRHFFV+y5;ZI$V26;CH8 znZuyPPIny*wQM8qUb}rOyo_N-LzTX9*{@v9cHS~^)HdnOf*fobtaLIsX< zBJqB!{w3&Y7jRGIU1;%j>_%jN(aEnV&(5A9!yVD&8ln9<@@mF z?$M^|Z&qpi+o>eHI(`xPw>n@aEALHyYK37R0F`T)Tv$FsJ8$~5)~+tv`Ly*jB;8T# zCDa-7FdwwBr+7au07i8NO~2Sy}x^V6@L^Ua%lD{?R7U(q~fBwTcrW zrgs#Q$Jo-vciY7~eIZ7>LDNyly$6OMUwli4IC}qTJfPp|4@S^Jq7?U~dA64L`AzWe z*ByFCqc(TX7Mj0>94NMJbw7m&XZ~l|eDgDWEt1SF$B*WH2nbwIjy#UWGB{fLNL0o> zYtn7pE_2NAy`=T>n8#njOY zOG<^xm&1ZZB6YuQJ}~~OwQ&7$Rq*c1GcVQWqm7UY5TONZsS=b?`5BLaK{K)BUuY@b zOT`5Xy>GdWAMCsfe||z)H=y=q_1fhl&(2!i@I|qjzhX+I^d&&)MKhi?4_@|ulstUB z(*MYV%R0tL5#J2dzR?HcZSOfap507J5-}UL8aKtaqXT5H-pNjH2z?^7cmcMW;zjkN zK*=Ij0V{jUS3`lfcf*nvc8B9{Fm|2?k8ZpmxkGkFxjeG!(Se<|b=RC}-iXOcH+q(L zE`jOpXH_cwdtGO21E2ZGwzg!^V;t|F`@gGj(8IRiTfJq)NLKIE_XdR&LJHkuPJ?yV zTK2^R;!J5+-(y1b(!!oCq zwHg_#@c@;I-05tsM-Ta45#%_w+~wobQseed(oH$ZsR<`Ak}kU(*Uzty{Nh|3>UTdz z{i#cFWUxjdI3n%sx2>KJ^Bq(oi6k`k){JeeyMgE5QF#%@Rc@=2%lTrp)hOlA;mi_3 zM3v2ddvo-DE#i|~H=N>=VP!f=)A2^O{k?_9!|=+B&kD88{IQq-wZ$nN8NV9=2pG8X zvt0a1q)s-&_ZPp;0&r{nj$@txotDgAa1^~LqGHk{iE|DK{VaSW`OUFuw~hZp)qlsc z{lDJ>u)SwX(Z;4mYIfMWwwe`tuPU``g@h_v)E-r8)T$X<)rcLV#4JTc?X9*X!kxZ9 z-`{;d?mzv(U@yv}v5b6rPUS0MKKL-Vc%LOVcO9KzOGJ#jD27kg72li+FnBUR-s z5x6s+_3?2lZ!@Xz8W`VR$t>`3Se<{;{1L_~y>88PUKWes+j@3|us z_{y9};0CyMM!NvGlDheiH>YV7E`!G$8gptiB?;q;H3X=&Eh#`@w!j(NS4g`-A!jJ} zyLfvw$^BwF93IF?ve(h)SJGxDK=2z!dv&^(LqN$dyW6_qgG%1+c~|didTsZdYb**| zvD`9|z9eP;!FZjNXh2IBdfrPPAE62=NU6`Wf41SimIODHJD83;NFhQWqaRB;bS?`F z+k_~a(@=by2;sV$=ey%ad+?qJ_1p2Z;($6NE#jS|j6jz6ao13;m1&5Gmnp+0bYm83 z_|g$W&nwQxsugLlvv6Y|DG{w~&^-RDV)i3j;bj!W^!3vq8diHA_ zsj$KJr$f7&4G`&5tcU=YS?Vy_CEWYkG+X>!M?)ZILY7n*Vv8Zq7B-JgdhllN>7TaV zOJy=A`ITt|p>e{*xkP{D^<#(s9ArSeC29oP4$G6PIbp((@RgpYc?n@c#uAHRsf9`E z9L^+KErM~!B;{DjI5iJfI{x4keD=Z?wM*Lj6Awky*+L<-2m{?iSaa;QsO5V;a}FEc zr#ZsRp={G~BNaB7NAu^7C>Hx{FWwRTrCiJvZ55ke@Otj`esgob+jchoca(4TJiF}B zhgA^+%?>VC^PX$;H|d@n^bA_63W>Lolp{VIq2%l%a_QOPxw4q8AnXKE4o59 z5z#pDswx8a%VVc^c=t0nyo=I^TPh-@lf*aku4x>s2b6O^djcc4KW2C)y$~}xPZ2x6 zm2qoYce`{jWhf%t^OxFzH+|l8o`1F0Yo<=vh%c=H;AU?qrv@;>+{0!(V7N6`bOz;+ zwP3Ax7f3m9PLz5ApJjN>bl~ z{I(9qgSNqb03WF4bPMZWO+Tim{n(~vhiK-vRj$rW3oI^w-4?5E`_n9!J3hm>f#fFz zSf~)=O*}3uFF#+Hr}TK2CXi5(Onxj{yRbR~TjuWBDcr7A{q?ZEGY{YJ!K4w@^;a^F z?QWy)D)fHzFrRt*=frCNhs=b-SH!v|u-13Z*xceNi0!-#CD6JNs6cHJd#$3NbTc=m zbrS11)~4v#Wl5rZcOn0Ux)Pxk_%Z#7^-}4RBl%}7!o!EWSu878Z@N!&Hz@P4s0pg$ z+2=|Fwt!-Tp~9rqpyKMlgP)6k^UXva&}7$rXutKgpp;o7#5Vr)nt3gzG4WoA-RsUD zu!#>=rZPaXHD8p#^vx+tBK5{IXFW+Emw#Qph)}$fvXGT*x=X~@P^DIYfikC%LrT&4 z@K=K_Vk=pOF0}ek*1c&u8g!Md!Mf_V)BnopA0Ih!Rv8v_Zi0hzlRZGO&e^sOqxf3h zms50<_g@pN*zmt5ICQ#q0)C?sxk1v*NTZ>VHO&##Sti7{mq|C8DKIB8DBfU>Tn^Bc z-^dP796FoX8W$*Hq-LWq^X z47{qi0wO7Dy8Sv0_(lE>EC+rbntv2_3U-#q?2?jp6>jn-^_)Y7I3GN3HK{rw=z$E0 z)Uh8mnbDc<%i<9zHu)z0QPi5=Ho3F1bNmQTRPo-&goQ9&3c1r&E!I7=5wU-(5p93) z%Ocl4XDx@#$bZce+l#E73SpYr^-GOHCo!zjAlE3>L*wEihSV>d#Nxn?SuI9A>coM= z;ez)>!Re-dBqgt#n500$`_oHhf&tTNJ3pRFIKMWxC>P$!`aI$m@>V8pr^s*7X>h<0 zt{J1&Ox6_{4olYqb(H`=$MfG{z8VLG3;9;bU0qTU@Cq7&~TLaQ6lbOvky+>wd_z3oU%?n>d3^;X01rP;oo zu($Bv4Pp8woO~uW$?Lfh@_8+Gdd>Aihec^>N_UTgUO;dx3OCuY zlq62;N|6>Ml1um5-kg^pZ&(LobKd}Bb5rg(lY%BvC<(}S&8x=^VA-88S-aBBzn%A5 zD*Ql;e!{~lgZBni_SO?~^#%YSgJDk0n zU)n!d9n=(^D}9puGK;J8x#(MUB>)DOd3$LpXnmHd&-tn$qlfP|3!w<>q>{-XY)UhL zz}HUncK*NdDt7RnMHPYFnW&i>HA3F;k#1^}xvtA8KG8lHPZGLF$+_MhEO#XAQ6E*A zi{h9D1h+^MGm;Gb6yzm+v3_C$$FB%Y-n0!dw|3@VKw%ebNF}k5sOooHVv=%oq>Qqd z(9vvYO@yK(ckQ@SrgKx{sjipy?pinuP%-}+KNT8)WVlIAFMB6LZkUjJ zkhaA2Ck*dFbRMj)r=Z15&r>{Y$vOa7zna;J;onXnU)%2MV-;@Dxc zU&PJqrW_aA`;R!9huM0mh^7C0a|M!uWO^vc1M2P?Jz#4|;-^cg;}Af9+gkH^A3iZ! zU_d~-y}x{_RIf*G&~ov^rowk|$R`=tCta*qT|kBTV{+H+WG1LwZS&p0Y^Y3)d3Sfv zN^!1oahJUKS|rI|CJNJQ$$i;3{s<4zhdw)gzDUIdPx3m^tvj}(3#Un{ zTIN~3m2LAii|H?kWhwGzYupZ~NH4_TSH1D3(dP<3_ofPY^FV}zF8#9-=tQ{%eGD23 zapB*5VkrYGt#PqiFncV(i@MVH=Q8G9B_=<0#Q|>Z-n4X_(G}K{XfT)$X6m_9E&hi= z-Fd%#>kkV`U%gS8)?h73hE-ZkCEUnImE33S;mWs*`^I8+RfkqF(UcoS{HAwMb9_afrY}l4V3}RcR-ze=&0OL~^u!@Z`M8 z8E~pcG$C8Ef!wEQRm-o&1wT2Q=1Jrn8B^6Dcr*HqqkMUG`*1GVE7hy>`7Gb4qNIAK zsScijj^Mjl9DJeum$q0EZ~ZJbArhn6pWgq7Waq{*TmzoW{cca$E)|mK#JLed_V!=J z!~LsxpZ8>27e9`WI*Aw{ltFuZ$QM_g{H;`ITJT_6{q1^t(wtw6zZ&+&lRA{-VZHa^ zGzLhFeryWf)80F`+qv1vSgX?REeJH7uI!1JI1mw{5)2gE+I+_p+yAQf*lvhxDDII? z0)s~3WhV(LT&E+kmp*z$NXF7T^*H9Z3kChk>>=GEwQQe}a$rO`(4zAP|EEZCSBu|x zM3Tf&e8h(>&J1q!58eiVzNHayBE-PMleRR~(!wK51y{01s8*sV$>_qcC=%$G-`ZjZ zvY|X>K|8Su2ftz>c1z0F|D?%2&5C;rEHtZB1XZNXXQfvh#&d=n?L`3tSra+UA6bFR zT0o*H^N$}d);AI89DCyWI@FyOLORsLkG=xF@9pcBM(D?&&xSq!7B?GmbpR7xKmt8p zEtos%pJ&iHGfVO%4`Oy&;(Y@oegwXZcX5mxzo*Z>yye z1}|og2+OA!R{8S`;a!uA2)t}kB8hocm|L;!dVXvF*M~z>;=jYBqY8*7sGkRT=UREc z@Xm;sLf0qd*bU@5?WftfyYJqF^_)v@U-822>&k+xBS(FcJUU%EZnN13dUy$BJUH>$ z8FzD;zex>?yz1L|u~r~SE__zPkCJpv+A!Pj&MUhS8VXB8X|{wc+mNCZ1kefCx!;+T z2NMmSLzSvG<;c!X(VXm5sy=6e6E9Yi8OCkU&Qz+D;)Meq@E@9{P9rf><0k%r+&mcd#RZ@bCX;3 z1CpzhmATpCYd1ZAhKZ#V%&Ff0Y~UArEiy?zahrbkYL<_faZYPli7;+X)Mx3WYM?W6 zPgR4h*+-53*}fP$Y*s850zbR)I&rNz+~kGfforwhQ62Gqldq+|U?akrws{%bz9zF4 zPrC!$6WFOScwo4lwU~<1gaKrKz4mLB$f6#w%uBR-HaaDZEAJ$?=0v}mXdF_Hq3on5 zVa=OS6_7hUF%rML;dp{}`3s&7lK$~;Uuk3Mq5~yPc(2c!)xm)Iv_gE1OapXRB;6;z z05=O{iIF?ikts(S(V3&wslb5Y;&=TM7w>4+iv#-^hp24NolycVb$TiZWRt-rc)$jS zg}FqBSiF@1$6cQtllF7$OQ-m1+$eyD%5ADv2E;$~a`tk^x$2#9?Czg_0*aF;c&2kR zMo9see3zZl-Ad%g!Y4=5#;6sO;*hV5orOXaSSSCyRh6*W?=M8Y>>8O&waA9%02udD z28OQ|iYX8sKTaTE1FEw}?g3SxPrQdDF&B-C_xrIDY|!(ihEo3y)mcRh<7@~8WP z;u*2oD>RCw*zRoPaF}uX7yEFHeeV9<`$Dp;l|tl9eSg6+isB5l>qHl-UBubn%KgRwiSM=pE}BO|^O>eJUH&2@ zlFz}9rkz;rk#++%-zSO430IRgyPt%+Q2@|zl#&t2fzPNlchde}4=4pqj_P|}=$It# z!7H3~TSe=pf&xpTyH%2Z^^6ye+o~E5;u+~@jYY?se%9}!+A#U2);Q(*;U}YI@)wg4 z4~Y(Jr~Pk7Nzz$7d#oOlFRiqJlrvtqp>Nhmk-)|!+L%lfLM=~rAfwqovXH`Om!hOJ zhR2u_?A8$Mi% zg08PngEsozp(A|n>=3&)=*DD^`1yR0eNe^u&3AG(t;3{BC1}x)C$JY+d*cATWC1$p zhRqub$}h#wJM!B(o#Mp-5WA|cm3leCm}?)7?m~}+vQ92*T6#`Xcb}fx|2tZlhE8-`{L&FhL$Hc z31jVJE~UM*#E!Ra=XZWdn*T}?jqBxn?*!YZns%7~A9`)q>%!L7PeNO;p1vIw=quGr zv!VB3eu2~wzoLx}G2QDf(Wg)DHLL6B!q+n!y_PP93)5@_vhMk}+E=Y_2Zs~!+r3`6)L{!08CC=XBRs8O^_+6qW z?~-EXw~*VW4ui~3jI&|s5~&$ePhYM!mL%)rUb9tOpwI5HiTGC3Z}Nki(-L!-ysEY! zCcr+QEvI)*{cLnL#>Z(D%?}>CnyM%XqynoXdcrE_Pu|f)Nl$5)*$Ku;hONvQ(AKfi znQeazc^Qj6Esfw6c#1fp%jV6Pn!m19%5+-oIrB`ue9J*x=P4Op?!y@v_D1nyJSz zH8D~FfEZUIlF9TdUyoS*Rl?AQX}nVHySggRtsI37kLV)&Al7`BtCsJXuiVd7zC)xJ zT~<2^|4a+l@fb#j-q_NS1RH&yu>us^Rd-}Fn-`wkaP>2-LYZR~QL)GoWo2k7P(*7# z)`Yx2a~jL)cHIUr7i>=d%77~C4BS?VbTAPtd*PPJcA{qfo%`(q$4{_Gg(4;s6yGW0 zHg~`vn#+nh#b(-Jlu8jr16Qp8vlVe@^4*TPUF#DTp@}IGhkrutj~3@+hg~kotM)H$ zLj2IBj>1hYD~6+Xg%UU36QgY6$MDT>Ekb0k&z+%~WPj`bA2vCKhLnoHeL~O{$rXEL z(VPn>F3x=eYOBVoiQGm@`Q)req*;0S{@T0WQZ9aPj8xF!CWbA1qXEE{YnAf%l;-LO zU3Fubw9=C)8TvY2OXaPf&=}^SPLk)TDw)6ay z_Wm;pt`3VcCDGo8WIg~Pl6MKwYH&M0)$5KzoWWPz7JJn=%Sx>-N(RJtt$`9*!!|M) z9z}dR#@UFmWpkNky~leSe@JRH2MK+$dW;BbN<-jIc)$*?TPg%#dhvxm&xjLV%3#*$$Q;rc$ zOGQ|a{}!V9HU*!vhsxT-vs8tHpLY`SQ{)kM!bx5^M(OD8CE+=+`|tL{3^jo^eNm|m zzizw@_z083=bubVft{guhA;Xb`8lA}IHcd!04U5O=bSNEQf%EOa9vxyJ3em*sF+zFDLMtcRnF(Cls=6?&`3~PpR_B17|mPS&4w6LloAobI66>nZb%$ie}k-W+T?G<#3-G4C-S-e=bm1V(Lq$VNp#%|0se+Gr&!LPF%3X+ACTZ1TnZ!%|eV$;)j2dl% ze$p8Oo8o!9Y9z;PPNs&}-)%dMJ*p7eOZmS6q9O2gI$Tau2@!ZevgT&9Fs)xlZtm4{ z_;m@BE`S8s*Ju)F0otrHPM(g14b^pVym7BkKQGqBD%6Z80WaBjs~YV~xPe4VM&SvB zlAK|$(Cs}1e=qzz`;?=;GAtAth@HXl#Z*Q+iE8oyHupv1BE3Gn>Zku|GN^Sw&B0n9 zH+2Y1-VE?)x;90dTXZXB_T!g2CfVD*pKSVt={tPHibfiLDx3T6QnWEw_V`9`!sH-U z(|{%NUE|BN4PHeh0pyC7RYH47NK(cEE77m1+^&JxM@lQoZr*45$D|hKfxFg{b;*-h?^wFo?2CPuG-q1w8pT-fb%R)UJ%1l)Zfw<%SQBUL55x*lan(c=`) zx0drX7J^gHR-kF*O&ut1Ktq^)~ot~N7)p{tx45c!w2z! zZF&H_3FrY_FLy#(USvD-W!|O;olf(E#lp=8Y%UaERpEREajkiiyI1$vhY&77-P&+n zYWcY*bYO-jR3J+GsBR<7w^NI?2cDYJHHe=7*Ve+)u@4f3h)y0|D8MrQJkXK}R6hrf zg6Cl7T@f8cCR|qd%K9QG4pm?)^StR~{XaAtdG*sp|1PO@(xwkS`%&VERUsmW7whRd zNAmTq$@!fw;KwT6%W%4!E6GcR9XWt2k!_*UkB&ny+yP%#QQ|#$!h^Jtzj-34H-Ec> zPG252obMEk`U$W6g_sL{n;q;Ld3i8fav6^OTh8+(=Y!g)N;bVS<*peunUQo?Rksh0Ae1W;D$G! z3nCDiKc?a?e|CPsd)p=yORxeAQ6l6V>4e$of$k1T%WN-sfru^Pgrf59CwI?J&h5?l z2_k&^;=&F*B5k+*L)*V<4mbh2kb_CgW0Rmg)SlEgr$%Q>299*36~lF1?o-%>3KxX`Hel{wqHF*{fegyyc{evSS~ zy|6kMFGF4KEGu7nTT&vasD}5y!vq$D+-KjBh8uzcS-^rL<6=&%#>vQ={1x>;?dnq7{C1`hbS;^mg4ENAUP76P$7eUokC_qIKo70jCQc&$3Jq5F% zZMFfOrZQ%qLF;o|ksEc{-YdUJNzOV_N~^7ewQ(f-R>R|oB^dhM=Kk*!Hx?{7>o;%A z#{NrU1nz9XR)qXBx$kD`6UZ`?|cZrJO5(K6D9 zzdV6|;riG<^iJI(YMm$v8qq&~_)9za;m5X2ZK!J{>)sRiR8Rb8zwewGijerz5pS&n zdkz6XMc1X>*ViA)m&A16ucO0NtxA#=l4%Ix6Y^(y&Z6`EZ^%!zC8fSJG>}f3<7c$h zH$gd4*!6KGewoOK=2<>oMX}b^0*qIUvy|~ZRK-sylEE`Jt86}r7f;P}8|C^j;B~{9 zv8?~pDa}@98QQiHu;1K`hE<$n5H|xUdYeN6QD{&)tai67a0`KNa$&uKa2f~Lz zPobJW={)RYz$;W@tUxp4nLvs! zGdVgXkZ;P)JtsibB$kAgnd#!tjVIQ7RiBAnwoKYfFpzl0Nd7O49p_J0@}L$!e74~G~D&fSusw4?9fGW zG3trN7$Splf)ROZwX z>g(4t(AZA058wTIvO<8W=C92tfGth7ZenzQ;vB8UhTlUyvj7;dnsmtwW;1Em{}uBn z(gkO2k~`G&b9`?^0l(@)J>U!D`v*@ZI%L2u7ZD?+bn$ZSy3D6i9jHOfbP51-xc%$W zBQNdK&C+=L+q%$(4Z*xA`WG(rgqA-FjXK`}H+}Rw1*2qRdh7Lsi;YXxE*RINrQre9gVF1?4SD&1wr z23@nH3Z$#ku=%m$_2NEL#MqPIPnftvSIpvlFiK`LWRO>B`i=ywl2(Aftjab1qhlfB zhfmJN;FN5PvMv#QAXDJiB=Rq?K}jJFkTRbi03%_gEQ#{U`0*BSEF;sNt;BcD6iMQ2 z$m61b`TJq#X3(Qt{^N)fe)yAA?sYB$DsG6_h+w@kB4n_p56=N}PjUQWIb@N#Ghdv0 zvU_Uf#SM0}{60i46I;Ighn$mNg8(k7 zzH4R-Sjzl!i(DRC0%0!k z;5%IoPWCMHl6`S<@1Ku{g!hm|3W6FGTUBcU{h(~3dr3$;e^S>dCSji6heisz&X3&q zA8BZ>U8gWrx95FjnB)NQPmJa3*15Pmit8m}i&O%V<}ycB)jBcDmE?S3XKz4<_d$bJ zI4kEMVl%j9A1z&ARIh1Cf+mb9>~^nrx1J=bLlfW_NgUDl$-!a;fX=c|_&Ivt4eq*SG!|9~85d|GjS54aDl zz!pJ3An}^P2P#~{QI@}aM)s0`r*8cU&oXW@JCXgia7m(q;M9~i&pXnEV=%XO?#)MZ zuelj}qRj45d@N*O^{3qZptF|+k@Jfi##O@bB)-$Y*fDl!eTe-&)Bd)<>c@uXljF}lQKBL{p;v7zrP+c*X6kVd}7SrTgk}5Ba~QK=Xw_9;|Ic0_y5Cl zw*Dl+BS2%A2KO!n-NYk;IHWPLSCxR`pp5tzB^XUwKLa_O*o6(Fv**h4u)3CL@`~5jrs~Um=4?c zyD?h?NZIc0hs*x+3a6TAo*B)cC1LHiTWr^>Ywra=TqPmfLe>^fQ z*LW9gG$M8zoKEMptC$Sv3hT{`6BCR`{vjHqBk(kh*YI8rz0hA@rYsSk!Th#XH6_;o zr$~x=(i11)!#;w6&!6)1X5@8Ti`7DZI>ADoV`M|OLx*(c5+It(xz1T;kIH;)znfaq z_#`M~$U#~U%1r4`X3L%*N{%9(3U3-p?1kw!j?JYIUQ$>9{HAZ3#Z`S3!cdJWYO;)DYvD#Q+#sAHSHvNEXUuzv4q%e^Kz;k{-BMl)~ zt@q0+nD>I|m-cj&_wevxj&4Ezv97<+xAptttpBv;1jwn)^wTb~qa)IMAkD*dO6Pus znBP|XZ&MRFia5#aL7Q;?DE`EQZ^!R4e%T3u-XaA7KLlCIRR?{gw=9Ih%Xavqtt<`q z%!QhDU&9NO3C}Tn1O(AXyE+QPw%g^sZ50Kau*j8I#*_{2?ont5S7i z&!ny)Wstese;YqCxGQ9?7SH-lj4IzEmN4|ROY>b#TET#S^OlE7v=?YWT)V4$;y>2m zsm+@ht)rk$bTt}y!0ZMY?^kfZ;+cygqprekf8RmGJ~ zlC*Y5T%;p0>(m;4s9#`5waz4R`B_5jZT87o5!3meiBbx+3c1NCg;j9;o7dNRFc0Q; z1yF=N324QxWBQ-n^QIW0&zT%`^EDUe*5qiX4+WjtbrfV^F@G*6nDCC`LePy2+FwSo zOuUODBjq0Xeh{;GlO|Gc?KVwd)eHEulU=i>Ch8N4ZvTbH$YQcZbAx~|M3)v##30MG zri}-_sVmq#)o@v|5N4wk07j#6ADv;JC7DWvUU<(OC8fRrA>Bcf2ocj`?7Qs9kDEhz zPa}S_&AiWTvO;dzycCN}FK)M~KIaVzx43Ai>u@osQ`h&q@rC@Zfb-J#|BzRAI-{0! z9(7O;{dG05e`QV&B;M5LRI$~Lw#3_^kRz?@F&AEbKP^8HJF`iuglk$o z`}3|#1$l=y>8o9e_m;_A@%mRxCYQz9Jt6{%>ENfpv7z85qwtYst@Y+bm0O#5ud$tI z6y#T2?L`L9Iej$mQ^}%v{N)FTrPoA=OX8td*jJ|q>m;30&XzoB<0k@R%qk4L;oiwm ziar%Ol}?odp*Gm*60Q?M)GG>qka|`B^H>*0?5kzPK-@9G3l&uOJiT8ft?_}kATK1O zScIZ&6kgd#@RlaA_=hH$ijZR#;4eh@YbgXY}4f?Bhe6Eg_Y6xgyc8+I!++xBuevK)M5+ z90VM%wb#3?IpVFD8!hL@d_KQw+EP~`TwUIWW6In8 zz2Dt@Cwks-SiT>1JZN(A8Dx-s(R4oyZ?4)rzS1NV+}TVbGJMoIPKysoNkijG5twj) z%2Qp4yC`AHA_-Y6zT`(uWp$?ArI_PV>>T%3u~)k!AsiM zXjy9ld4w9*kuZjrGA{lbU(x0!uoC%l-8`k@0TyZ%8sH2DeZt$)d+U!A$I#k!0+;#5 zDN5YLvWd=2I_cvnIxs4|vhUw#q6%{w?Vi55G29Q|yvq2(X}HCarHU9<0?f8-u?=IK6Yh{~Knr%NVLWM0cmuy+m>Qyi5ov>(5eoDmeR?2Ik?{ zR>6#|8Wi}g=zd|*3e|VNL=M$KH((%x1=sbdSE^%ts-NaUZ7Zf4m+v8YE=h zNZS6y`qRP}fs`D2Nb+rB>D4DIGcA3mF)_(CuBj&CyDRN*%YDATC--M^?j-c?&BQJG zH(117oz_+L@+pWAze0<-NRm$B#G2F&81*$!Pv^$&$e`tGqM!v?k49U+R zw$CCmo09zbPqC-WdZf{M!9M)_C|A;SS&(fTtVTdqzkVyEUtA_p>8;9)ggWg#8l>Aw^;{2x?QvQ010PIX;@MLAt;j|q2{4N^<rZADYj-(9~uPh{p!O}x4@ z)9@6;)a&Bw%guYRmwTA{8mZBqU!<@o8)DkZg(Z!d|2i!<+z9E>2Cat z*dBipx{iGJXi+*VK~s%2O9mPaN8PH%zW+?&rd8&%1o#!%fq7-6vfcKf=l zo7ci0YtwPDyOnl=;jmlQTv1ntd@qU!s1)27k(KDqLn7f79m(?$ zBr`ik&W2mNfOP-zF)ulpZ&>EZqh{^1kvmLQ z{TYY;i!XU^I$!&3AvzPYi6{!t)dTZ{Y9@Ln5Ptj&-hYctYA zF}E;uus`2(QVOX`_8mh$ylwv_H;A`ETXP2)mLc|&}aI)t%rtoD`5%!~K zkbdw<)g|0ca^k*BLCUM|LslmNqQI36ZHN~2E542^vLVyrdi86d|5lK!(oa+82jl4J z2E(&;;I9>5^nZ_l*#A8O(tzYhl*wOGIx-qG>nNKNDmAw!^QpU1|1Rg$5s3^>{f^CjEuu=kPP4p?)R&bNBA8F{JU_ZvHrJ7 zkiK2>Nnu|#ab12bzh0hY#{yo5#|#>w%#>H#s6?G!Ujs%n=5iPW#to3#O8g&FplRzd z)hC?=31PeZ`N#6 ziN!@?d~%bbKU3hOUTC$Z0b1F$nrC& zk64wT9yQsEH{vT)L$j~x8*?(ei)EMiV&yLk`fjx7aA7X2z+V^T3ew+wozr&iIH6U- zpWr(WE9gd?trW3#z?ReBfsej#GANAiqh{#4AA7QQhu_ig{tZmQ^xq|AkM9}1@J!3S zFYVX-6=rmRR&3Z5^>{QnTFYPn&5voKvdZN zmYA+M+dg)hE;p;a#ycPnp+0{gF9oh&n;wHc6nQ?$UD|GAdOYsJb- zO%uFwA7hEzTtlv3=B}9buHTM~`;oK#Q_SNry^-up+C75!&EmX~{$>HHJ6uInWA>Ac zS^M4ZopNb}S-~wY6gNWmusU~x@B$xA*^i7LzXt;V+HJmjw`;Y3X3nG{zfyjZ@m=J@&g<9*>$=?lAX1!_L0_pNRjKEkSKDIn1NoX&+4(jR7iKI6BJ**R3Uv;D+7S?S2M&8Lb5XW8GuWK?@1-mY<07fg;xXsi1{~CKj4CJP zYD!rMOOXekEUt$4MLmh7Q`MmGwP<+wjdP)Q14B$ml1D)OZXZxx<^z{BXp7VIe00-= zg{QWW(~mm4a*|-+Kag+n1!=Uli-vrkI&&(Gka}E`Z&biXZI-+)Rw(~}(X z|Eu+%qZYi0qLlHqUS&%irS|H-T7Oz)4@IAb2&rZDLMTu6V4v5rmxT*>k~m%ai5)9%ZQouDp5WYLR{{>?ZI72GC%3_$I94x;1jfap{VMJN(!S?qVEj)X+LZ4J_+|3}@6cVLD=pbb zgGB%sOs<2Sk9*lKL^kx#+g!BYQoiXY+{0L_#IRg+*4PWtsUr7B+vk6r&@MblsbjB3 z@<-7A!ue9-(%#dZ+khZBC(`^|JdqSc?zOK6U)c;sI}k zRtZ{}u{LMunC2rg&~-~!n4^=+VR?PBQvTG9kwY5UeYY;hdqhJ|QK8u#LmD>EdjS0aSjXAcxZAyLp3l>zi@tfMu_<0IT+mJf& z;s0FxA`G8oHbk#drL|egLMW|K4Myrjh)cfNh!kH*qYBULsT@GZ4QXMiFuQSyz&*kHT2) zkN=BEkZQYG#NY$1@NGK#_^n(kDhBtVFWfTCL_r;=l|6*-d=mn&4Qp4&5uPU*YLYwO>O&(>s%clR=hGEcq9`y)?;1ReJd2hc(-D@nXp z6_y4G3^7=3o16rq(|hGJ=0lnb+70(^#;_?=;U8}OKsSE=j$e5?2>C6suDHorhMJ!h zGloxie<$M=vg`i&%DN5{*KA~$anuzou)|I#tzV#gs;ENv=>3t>_OA#$g))vmMc;U%4io_Jz;Ji%Q|c3pPn6;=cM z-G(#q^n@`yf7PGs=TOsdXW7vqq(kB#LDn`)QVXkqlnQF>+y?&kuF4Kq0^cR1&{Z?< zl)>A|sPI8Sp_z@2qdl>4=$=TssveXWVyZKfI6nX@ePQvz022=J>dQtxvw&=aXA^a! z@tPb|fBl5aKzl><9A@zFQ$VG_l3wntxv8mqSqO+lsT!&&LsN#+Q>e_|{Khf6V)4N# z@5$PU3j`P3aamCjXL?2wpdOOAshAxfe7Z?Ugg#LkwyNvA%3#KBUv-5U-8VpDS|~-(2Mmm;GLg)yEBo zo7L^E`Kjn;y2_uL<&B-a8GP8uGm`e0DEqDN=w(?zs(=;r>`;Y!Tnz1E7<{RyS1^PZ z$-Vs!Tlbr8<7d_9qCUj9s`J}J&1AfIMTummOxt~U$p8Sq_ z(#~YYi5ewB>Jx11?!>4~;_dKq${uvi0bC84&4x4Cc${4`T*Q`JPFVs$_8~O4| zQSO9IHPeo4E&ln`tG~Ifq&e4ORF%2GpiSWB0`xo{pAUim?1!wq_FvbtH0O^;Eb+^j z5}(3ZOey;^&H;AKx8N@oj=n)=S$5An8Ys*~Ux1*GJDPfh;g32{i;(}n7>zkDTgG(y z{|gd|d>cg)&p<#xu&n)5<$2e*P`Y2Be)ip-`e)ppp^|U;TUr2T#dG)q78XQ-2#og8 z2Oc~4cXS-fsQjG8r$J!irZWDQ5kW&4>5ySJDB379ES75Y08TR$aq;SZK@l`0k%eh? zWM5{(Zr?Z5$8#1x(d%gv;eRVMx9)TPFUHEAh?7^`lWINPom($|!EXzO56pXbdi^)B zZ{l9<{B1Trc z@t>$L$V9K;kmMSQ{}*p=K+jbb{@>k$y^i-V{vn&qoTft?53{qSmbO>zjVzw+ubpyI z>w=dZ>>p5b_epX%bN7{o$a7zlJ-X|9>iRbLjw20H{%!JlokN7%QPuIA-}Zmwj^6yX z=?F3n+PA@;L-G^icD}B*U#wNDBP|J1BYDh%nb7o zQZ6z@&g@pa#AQ}|Up+HDo4+uf=`=$_TDod&aZ154o0*DgXoFxa)mDB*fLHW3A^DM>(q^;)Apt|H zfue}-<-f)>39{RWaMOQs=n&#|(!i`3SU)?I(&~RSe%R#(o0H8sfB6>>VUKwpLssLL zu=uhOVBo;e#FZ@;7e&MetISr`z&d?7^>Wyk(!4Vv-jSkT@1L)XT||&%%J21OXp!Fw zn`2kcnpvcQG8;9jkid`BM*CD?^*DnqXbcQDA%V14 za*_cy!?aE%Xu=26=R{5bUUIu|?RH3GCq%bKLI zQ1CTCaGDSOwFrgCyp_giQX=o2DM%2A9C$Ik=!Y$!8C|a5{lA*Y0L^}{EB>D|zEqBM ztN^_s)YiG6D1aIN5!4?+ z{n4mDruWB3{c+~Oe|Ek9Kfr?$Unl*q6A$(M%*^yfj%N$U1I)W5egB4rEF#EE8XoF1 z^gbG*Ko%8l`j-{L!=EOD) zXZQJPpR&9{WAw}3d%0JG^*?`pnep>i?dN9(0LXk%>HqHU5XUa9>D3K=&jl;}EV-oC z6L*9?V>_Mj)c<$DBN%c)@|DA|!BB*geuY;H_w>@+zu_^at1f%b^i3~6INdocyaT%) z7+LLXWCozk*Qr(Z$>%_C_6+FZGZ;QSjoknXC!) zILG`28RZGrNlofmtpf9r#fmdkf=~dMb@fh)|0SEPbJ=kG@3PUM(diR;dSG~t^3G8D zc|>qFl>hageY5~TlwBsx;QC+x{68Dn)tTj9|Cf*Nh5=An`%bgnzn|tmd%h3M_v&R1^}k^7e;kspBA{w~Y?k8x@yWkE za21$uQbhbY4E`_rkmUk`hGcjN(_b3(AEg?c1kCqd3*x_hJeyjlZ}Xn&=%xRl2Y(lk z{5@d4`Lovl?c>@&RDjLrag6xyR#z>TH#gUOV5{`MeVh+a2fqQZgwlVh&wtea->=*s z_5Zgf_DB8yy`THz{r|n=`s4lo?Y{lT`~UlW`!Dc=kuj^UEF6CRrBGA4@Y&P{AIhDl z=U#m&rD^E1zUQ8szNw6{zJ0vG(0li+M^IK2fYGgOPN<&$r)6qH5SM1tzU%o6@pjJ;a zS{92lT8+Nwoo~H&$xLce98cQB^*y=p7%Z=P&GW|V!!D1$zL_8UVtoF{&0x2wz=2A+ z9QW6-8%tG70byBARr*6&&R+_RpF3|~q)I$fj^L}b7H>BP0KB~YP+u0;R(NM`GU>CU z{*T2e{R-?W{R&)>%ltSKMbK=Y?OzQ#X(0jmg_JPyHGFaxs^b*{z7m_|CS*n-LCJ zUSUu8?+spcvnqQw`}K=N;@3LIzn|f@6p)tO@xXW4O5{x=04)j~JfGIrujTkyxOt2m z*ra#fratf1kUMz%owQ7i(fsl7sjZ(?*}v5bpRZ+%IDabqYKg|yUmputYg9Rx<2CJ>=9t+pEMD0@a$X$A z(CFsJJ=vQvKDZmr!%#zgf0sM=txkA!>r1zH~qC0_-QXUOJjS$$83$qRA^$xp5N=7iXRk_7;z z)p=bGt!Uh{kI{ukM~K3eLuaAV=GV?hS7&|MEhT-=2?y@`r#&fwjSI%I8s~$4Ek85* zqwr`0sAJR=hgV}fVhOkoG|Fa5!m}FJD=e`{BYkP~H6Mx9b7UW^kqe)^B;6}>TANc? z2;2;0@E!N-aZiH*pEL}CRb6aznJYU-^iA_*bqD^CWxV>E5h&)>{1@rY@!aQeS!#OtKmC!%r#qqgt^90tpLHku|*~2N{>e}^F zcnS6_sBOj@pKp~GJ0P|C*esaD;vw?T#%7NEe7z{=de)0$gy>*RLUepTM(>=d)h)31 z_LjO8-p%oC+Q@qP$VdqO=)`kzfaG(5?VtS@Yr0q;{g zU)Wr0K%d!&D(;*KO_mQ(01RcB&)ID6oE0y?YM;+6u1^nmN_L$w$u?&yU*5!B<=6##iFDVWz(yh{z(W zCKtp-gnriV-pBB)n3hWoA^pR8@6X%;a?Gpz{+z+vpv7{G)B1CGT3s{5RTz(OCPy{x zZF8zVxNBo82giP$PJk<=7*$-W@6Dm9CIhgPB*Cw$OY!pb-_n!KB`LX1*k^U;y8zx$(y`5zICAZNg?9o}m zICtbMJ3A#SagG7)&*Y>uR>%#AE{$&2SbirVkk;bI*1)I2@7^+aj(ahYnnHJ@Ec;#7 z;f3SgswdOK9)s$=4}j08KX`HST<_xxjEVl&7$iMeERYn3UQ7^C&*c@($#QmhQQj10 z3O#jX7XO&zYSFUZ`Ftc#S0<7_IUW_T>QyizMemB_4$c5l%dq4b|A`XPnM;0Lizo)l z$P%2!8&zwA@)oD{sPOeEWBrUp#VtyfE_bDPn_lc?zQ%|8s&eT^j@a75mDiHM48Lf* z@O{ z_RCz0hVFh23eZ+YdP*bt$*3sWc!$IfhHBl3nNf$TP(6Kn4(B50-7!t;~wlqbah9vn+crF|bL z2MP$=_#B?wcztnlbT#fpTinv;&(FGumbb@^k6>RZDz^L~qICx%T#G=oi^j{a`Alf* z4{hL<@3>l&bYBAAn!fx@`1rP2CBR^F2NpiHs_V+G@z$fi%LE9!Jx?}K zr+@~vLf^Nww9~8p#OZ&w28)51=)qNO+;i|qYp(6%qz>)x{`cYY`3-!hoXS}uAJ7O~ z-9ntcCzd3yN3u9p?Van$>Z&>roiy_LH~0*bcH5)D6E7O2$EL<#t3E@0aKqQ|_F0d_ z<4oNMQPsahJ{KCpZcU%Q_Ga~b;BrYCPo=+ig3Na^aA9mJ+Vpl{mL7P3Pie|Kvq$E1 zXGX$c8RvHJ$|$;JIIyNr3nsmYvV5^Bq@~zhT-17U^!R7M;@oGeT3Xh^O1ZQ(k;Ju( zO_t`N9^_qO+~HXF%nd*aqswMx=R7{TUHih?dP9IG7IALwXduA(o4HT22D6sxHDc>4^GEYykNMm_o~gy1$r3KUl~|t;uQY_F6U5BEu!MDE*mdn`y@3u*r zKctb`^leQh5%@mO<31!urA*noHDABcJ-;_TnRu;O+8i|_1Ff)+u>H&19| zGmg3POHezoqA>f=(>p+2hg?KSi%%_#BYn;rP3iG6|G+a}< z)ca}kx3r#V)OOU4VAxkUkUrBaYQsXFJnTdJfZq5*kEusT+G|AW)3~ux$QWd-$4{_l z{KwwxpMsES{O=KXBebHNZg)&I9kU+G|EqSIqxlqzem!;IH^j`a^qhCjoZ@VJ7vLg$ z@$ztwZa-dYbm>$oPX<||t5x&9-Eoh#7wFwK1H%nEX-uUhg%*r~4!8 zYQ`xgSM=9SSVH75r$e(A>NSiZ__5IbI*QKH*EKMy(_YjicUN7EyP4x1gt@c!q(i&S)^yen0yt<`Et-#=tEpzqc-ybG3ry?TZ@=765DT)d(VC``+_*Q>dc2AbRl>=5 z#!mJ&rQW>0RNi)hvSk8EXVf&qKH{3_%k8c(CGzcRk5XgG`llNhL?7TcsOm|a8qXFu z{LG1TBMi<~{tF$Tp+mJr+d4?p0qDZoagKd)?1g;~&1q-chRLwPV69S0^v zz=m*#ZhW(7)!-B?FAn(>K<3P?|3$7oV;n`A* znB}N$m;?wXj?}~iu$nvjIcIL}xo)JL76Z7*jHM!iU;D>fhPSNX>{6l(({bPuGF8sv zERs29$cEJOSKJ# ze#2-sVKORDY5HMkk$j9K`O&eG=+XNpFv0Ieyd}zjQmTWNCQ?)7;U9TvFdf0X{?ugm zL40%Fg?sm^6+o6yKM>n~m8xNSSQekI8^Rb>$3>HMpc^Z~^=Q_mtp>c@*{#S${z`)N zvfpXxn1q*$Z_?;PD*`;iIMDSm#d?Y0rGOELg7oM|$h9j*m6#l|VB zA^HpHYpJe@<izK!Q25n=UuvFV)Mi%2(qC9?@-tIW7$9);`1u0F)Yx9^1$5uqCvF1tdv zx|SscJDT=RPVY;c3@k&gW&U!3gc6M?+sG#FOmHVt7(d}!bJw(yu*jUblXIHk!M<_3 z`h`MNqYM9r(HvBxld#4kw#e zdt)8nYBh>1AeN=?j#TltF$!0a4$lb{yqN=fBIn-}ymto{x}(5FGNe z-A^%%G??gZ*a)@XHA^(uzN8Zp3l7jeLI*&Zdh5j<4_5Vfucho9gn7euLA{;yo~ zoLdP}7b1fO2x76Zd`h$Fl%}|HLjU2Df9+rZq?W@oDu%D^{zOF zJwA?8+v7wL^K6ZKq1#%F0dF(fQ&QT!32wyVYHyuhyFu>$6a-uBVpaWt=UK9jk(@Gy zRxEM&O!poUr=7|+zhqjH;MFJg(R8u*Gae%-Ohd4_I(iPXNiM^{42=0u&5F zjB@GD#1F%0v24yU*g*4JBJCpWkRnXqsRq2E4mOoNe<8?FTUt1B%UoJ8Cw5zSywxYV z331N{W8z!H3ay6Dnkdf!cVVt7^dS^-z5L;>v|2W#6Wcl_dJyILs8f|t%;kWylzhQ- z$$2&6`x(lR)pIyIJQw((%eF}VA(T+Em{zv;kYBko8Q;(| z{wn+XET8+K8`YuACc1%JwH%e8z|;z+h_cPkasxKNKCk_{xZ(`` zE>WD@HrYRI+m6s^Nuk7NC-h#fYPY>uV!gQ0+UIxvkJ^rlTTh3!MU)xo$@v6!`yzC6 znKWz8)=-suH(Jd43;S57*4fOPC`0iOX(#z@88w=+Ro^BiGMXyzy=?MCe@182R7ajhLWU3UQ# zJEqgy9_a*RI-r^!EdYT zeqdx}>v*D5wQM@lJ_LDh+XpNv$y9)Z+1a8R0w!t=_5rMwGx&-EFngcc2meSb<)_P= zL3onbEKjjQs6Jco1*ZR;H2>OdT4fQ(aD>}tD;nZ=J+%ncr=n`)?YJYl}*^Eznkw& zd#~EWD=&cCfPE;d@$e<2gd)fG zNjgK36>r%d)uBs`j`H?Jey8L00)f(?xanhuF$K37&T7)Sv51?>@ z!O#{aN;_LWw}|v33)rPs0KU@k>+w1>xFM8ec;T9`z=F@DWDn$@L{kk2l($n!sxu@`feBg!Yuz`e!?WcUekF zED3}2G(O4r?ktpHlqY#&IqA3v!Eb#xi+O?DRoe5%VJlS0P34<>#bnpH80!x%6XW{j zFq_e|Kz0%`HwGc-xxY8815Q63-Ta`zt2|HJ7T7CNPy=Z`C_F$uT+BEXe$>Yi*Js}_ z#k-Vqa^G=m2a(vh`uWpssTPK2-}R`i_9xilyT{1atY8Sc0<>Tfx&Ph6G-^%xH|mSd z>+aAF-%JvhQN<>{n;cpIc;YCq0KI50)s^SCvip=)Q7%Ld-JM?LR$H^_vj$Z)HIE?>fE~N?i%%mUX>p-8L8Vo0p^bX1ixd|FPK49a3mmxC zH8ULm+i0_-m2@*S1Rw3vp_tHkbrJ2}NROz$JtTRQJcvbSM`4t+3i8-I+h2cV(1P4H zsjs|xDRIk&#v+q{)8H;~04u1YIvmA}z36iz$aT_{|Hfv_ljE<@H!E+~GdAuEaeR&3 zf4!c*iGRBL9asCWY1gKtvoMkA?UgfV$at3cTUfvLuiJ;xm0kC+%X5RMQz1}};fZnW z?L3$1Y;I}iXtlR@b_BRxR_b#aQ}&5%AH;{>56dh2@;Es!d^E=gw?QZSuwlcR>=pp; zoP*}{a6El46gc`x2}r@=q&jQI(;b{aE^c@oQt9~cCMCKBld?@?GxiWIh=9fEmGeM- z_AF4JWvequmApd%fnnYT!WBi?FH~JGrn^1q3dm0XbusVWwi|yGzI#h2XmSLLfxtAg23D0sm4ljen8ZFJEj#iUU z8|j-LKAt`siX%MA2L%k^B&8s+`#B-_I9!b)9vB1{+p_JMfYKrX+n2HMe|+Yn$$@lz zV;x%|K!d3|r$`KF1gSC5_rO`BQjTHFF}4WZ-AfcNmGe0>Z@tQ7#w-;miJofT(=mrR z2o}|s`*68UHuL!Lh+GRvq06+V`rrlKdrLQ@bnDWm+nXohbVJ&$d?Wf4M=>V+PT%cv zlAB0@#Q}5T$FvFU6{(>+oetn5%W?zN@2G3RR^RrzY|$_3uBrrQT}=8Ni%ZPF*oatZ!jAL9qKmOOlZ5oM}7|bYyWClI2i#4NwMA!G7`|Z(fHm2g`AOEwHoaSbu%9dQR0A(F9&(K zjdQi56aa>`19&X>SlER+e;~%S18h&k+2Ehq(p$8gZC4*$6(lWF75gL()>Nw)qwCCe z=%jEFtdre0SL{2rx{vDSw$QWbuMTf)BR^XE8ZAx`BCB|X;)ODDb2>-Y@%eXXwq9TF z)WT*bpuRbm3>?R8&e4y6=7igbgl{XcTPEbCL1h3FT^!EHN+eX?Vk}Xy8ox zVEH(yoC$8{P=(|l{k@^Q|dn^V~VRDCQglWB1 zAh`QNsiQJW*?gFS-c5q-JcD!l;{hvPP(5Q(Aci9iHw;`;!;yM`V z@rG(jtb7Y%6i+`JcEofra`S$sHXQ9rQc>Nvyr#lWVZ~=9;z!R{)z*y6f^tx}I!9Ss zZ5K~tZql!41q$d2z#w@++^??raXZI{f4_Mf{V}(9ZfY9KOF;5A^40;Ju8SGv+Xgd# znXLi2jMe?tkf~a|TT~C^Z76j5iW3j5-OgeF#PAzwUint(N@l)~N*wCVq!NLM3B|1_ z(hg1cPBZuWYcl)n1SHbkur)02nvAuYKuVt$=xP}E)0t?*-|qrwzm|84NhGT03y+Y} zLBSNMIt7`qiR$4UCeu-5IxZgUJ<3z_$>r-BNBuQjEo0ulqp96rUg+Xb& zwlgfXQ56FlWu{lW=--|!Ltpyrjm)u%A?zIr>rW+qtvb|JU(JrM)cJ@VG)PkVvHg)@ zBA^S4gcLj9Sc<>Am z$lu{a?Yii_q1eP-NmNS_Zn?pF9nONF>M`BX?ThT_ru(1OoBcS@PW@e+;}>lCqI2T^ z<`6jxPL$<+R+R3!s49=Ss0s$v-=J?X%hC;^`Qa|G|7+fx{>ZS8J8!`;|0iGB>8SmH z6^ozgF!*+H95KR>x6W6Jz6+jDtVlm&cH$)nYUgI9aHJqRs*5L<=m`!b_;gDkRa<4Z zPy8uYvdVp30Uq*M`KAySD=Lb3!Jz9)e0KOuM!;!Czpa=PCYz5^Uc1!{4k7cRIsT4W zzbEsX-SEI(+nZ`p1rYqlbXgiMrn&Z4FD!B<+!X)lSj=Ryt7F7H;Wq1(d5> zQZV&?rWv`gYBV~=JQAa${CM?4nHVdH3n}W2&_DKuqeW8AGLE}xr@Obs0<1?DT)s6c zc>OkcgGvg2StH&!5=pcQeBz^k@8&}`sf6jS=8288DxNRpSQ*OuqnO{wmm|(S9jKV?pwqXt)gTs{h0VVeA`UwSo!X*vUv{H8!4R5r~a<#pnpWM zZCNf5O9COr-060yd#u)H|Cy;>teq-vx2lf2g5#nik-TgHUMiacP0AYeVqF`{_IY>; z66mtw=KtXTzQAR|P34;DHP2aL`%Y^7>4K+;4h3G*BZO2-Cr8yscb$!faUWZ}B zfy!dmFL&3CD2NpSn2+#&4}a4t>Avy+b+FRBZMv@et+~jhhsw77cbzGX+4pIs&Bq-;lxx?$*a_Ga{As)bANCbwK~mXd z6M9}8)h$;70o=|lJ}cCn3=5LrX@AK{H%BV&&)9><3gqL@ae9$F|U^Cy1N^t&rr zk3gUMxo*!x8xcl6eUM$4N2`ycxRxq}(dmtPB@wI5sihw{XaINrl@LPYukt8ey9 z?KgZ}o%KNes7TrC+rMV^4$JiIvxlg~+P6#D0OPW^HsM0hc!Y@t(jH(Id{GrGuFVd( z?)d5IXx!H4a?yw>^UFCiwqc9SIZ11xzaBG$i$#>$Ih|BqI?FvuZ#qT%v{IYp4uHLrwc(?92$uKsQS=faHUp6`K61?JF;e)1d z&rG+NVy}Afdf^cRh~3;7gX z-x*NNx$<;hcF%&J{<B!w&e$a@qq`e^6%pI*sbK-=?$|X3eKo{~2{qw0;8Y zU4-195}r={hOt`pH?}t$PHJrUs_IYS)cdom^2KJF`@@MDbY6#C8Z)8(uyQIE;gY2p zzFLnq(7VXg&YWh*D+-(@gn^=9!f{Ot)Kjh=DmHfP0XKEGcp5M=bMj7+R$NZ&LR=b1 zHuq^4izzHE*F%IgJ!5)Rd4Wu2mdf=tSu5y19NGN?yq%6NuOI$6^(!RME3M6orawnz zGlna_yu_$2K)y2CD6Ek=##aUrGTIZ_tOetY-ip29NYr<~uVCXq`(iQr^m-32-9sod zHC#Ciym9*(BBa``Cr(o}w`KbhyJmmcFV7mc6c| z#tfO6B=M(VnNn$)3BB#k={ zdCahwmSeU$s;i$!!0*ae!>34$OhIu~2QIm~DoRS6z;^1~-+$50|82B#L$-|A?Dws-4^xh0K|Db!0I42`Hb#_ynL{GLZP(oLvN+7>l81Y`vtiq=l@i&Ub%0hh~P(m+Ap_=HH zsm_Gaa>a*U!5FS#KGY5nPVD&{8#E8grqi>7JHWfHQC3Bvt*o>} zg`od%^0Q8RoJjh}Qb79>A-{Ej?G~^`e$(MnoCL<D$(p+ajjB#Z|2bf zRJo~*zXrgD{$fXk{k;}SZ_2bMvgtnTJc#8 z)EPLQi`NXP?W_W*UKa|^sa}E_oK#_+D<#w;BJ7$J=97%5ZLQgM1}fQs8FwWOWy*ex z_Zrt=xnYE;dBI#X!K(~K`lFRWYga+Z@JrtWeYcq(J%r=`I*->##P-Kj>oR_rlZMd7av zKXzlh7H3Wn(?k0*j*{u0cl@)^)adAMBI{Q%D`oSs1cOQ$2d2HfsD>b~Aukp&!EwdM zu=dM%aioa2?SfiCmso>NT9wGH(;f=fy$VDCA(qwz8WK1)?O;uVz4R?ch%~Cqp5?-L z^WlAk9OB0}nPx5!h_^$-0zkMoma)>e^}O3*8v^{O_S9~NT{-JpDxVY$=cKlB&JGsH zvqFtB5oA%0SJ64?1wSmYX}?SDf*_^qL!F5vk-0m^(B4{uwcqbfNutb0IIhY)>}Yx2Eh`tu*oq+Be#yn0Vi1KFj9* z;$)9Mu7XHG@S2+la3McIF5Rxwer3)}ejE-Q%`wM@ap_&>qfdLJH~X z)Y~|!aqAu0!7ux8x@xc&^`u_>sOpS{aG+-n+>+L-H*B;R4sk8LO4mOxW5z)F(bh7C zFEazqUVY6FVZ9jF%~2W|dpV{P;#rPdEz>lWc|HMm4UyBT=aEwp_z>NY9(Zuj1kbJq3s{4(6(@x>WracsS+0krt z4#(P}XyOx^x%$=kw+l~(d9Nh@yRT+zmpzJ{8!&Ohn@AeNvy7oi(@q~AEI-<0H_})< zwiX^1<4sTw745*B^DSi0{;F?H~z>-b@Dylu*5?`V)o> z^w*edSbQnEX8PzQ^V)*H!5c)h?++1WvBiCx=SKCy)1^_qi=Sb;hi57KU^FE={uonxy^cd7ngF%+#2>3XK^IElPbRU zv+vzQ&udHZUpPwXzoiWf?FlVrE|fa;#&?9jGn~-4wIJm^xi5pMkKayrE6+jfZls@n z%57{0z7|u&CPm%6Zm)OoLw#Z?S#Y)>^Si`RDYpz16->r9c>HkR<0Sh9A75#kf3Net z-;}51Hzx!{frA9HViiMwGe9iWj>Y8W^fxk)(VAz4(~F|iwI;qsnsi_Gd&XjADr2&Z z%vQLG=dAYSp>DDqQLDK#x&Bj~*vA*Fd>8y+!j+~-VRZ zwbO}}I*(cMh$c3`HRc5*@M`- z`#5Nf8|O%R5A^guoI_Q9zqgrv2y75H)_4+#r{Coi6#}Uy&rIuxNvmVsWBlaJ=wBpUHicQ=zBT-tX4zZ`;bQnO;2mlh~i^HQw^SDk+ zD!*TcKV3szBYH{0*JMVq1Fn#XiOC}08|93>OOX<&$BXwSR)Er9w1ZYGa@(Z-m65+$ zp-acycgEgp%bX9)KSFN;s-$7%6L^-1{SOEIrX2x#O3tsqoGtwuOWKTEx}SR$yRoE4 zT$2EoRB&zcCGi3Eog%ZS3XEX_-`Abazi4GiqyxOW^Q7O)NtAdG{PUUNSeGJSS2^am z5_#siuT&offtq?0ni({dIhQJ%r5zwdc9oPP<*^C;j@q8OFWRY?%+=D|SBQcS!VYlX zc=ugc6)Bw*U8JxcG>_;D`GzD;dk+uRryBtiOpNq$ixat=G|Be#rlkUjP$725_d$2~#=q5D@dsZG5(1X42g@AiD!1gir&#>qqU+fyvECv%ALUaS>}gXE|laacx$Fe*zQn$nMV?F|HW)d9X7ts>sd;70B>P+X3d z_OifqC836$pjW(Rey|ifn7j;hBmh3hyLot6@*@8o6!+Cbc#VtFx?+pX3Lv%9d4+DG zdLr|Ls7%TqSi(M5d5Ka?ly7e=ZmvGTjf$+X?#*5ZdQwj({f2mnhlHMfw>+z@!TF8) zex_@Tt{F7(gyPkM+LU#h#inYb9i>>lq-yBng@?sFV^HgZl)Ue>3btxbL| zUKU@35tkkLnxVGG8NWSK;B;ftJ2jRk_7fLx(uxX;#WPA4#lew~A~C#J_FzsTA_>=S zc|i0l=lEw5HW^#JKRo!Q#JmuYaEB{6CI>N@jqf+!Qa}*;gB1 zr6VV0%Ad61O8v+l^Y_CQuaCFd$OlL-*i%Mi9e4nz$@#b%!V>D1vw3$X4=GBQPLDi@ zSrJxpPyM5JY{1!V{vhVO9-}_hTbu*#02wU+$w9GStdV#9w#{2+f_>oQT<>O;sUKwQ z4PPAu!wI5>NUzjRr-PV5zIiRB_6WcaWS?eGJLJh4=LLqbhEam9D6@kIxVuNz?{kx^ zk79~=bzB7V5Q?^?CwEtWf@4Y9A!>smh2X;d;qXbzzR)4Q`Gm`Dp&0kxh9V^*D^#W)0a8| zR~3On(xX@jWhgZfjnugP7f(-SCu8qz@r!Si*fs$XxFun2=OLqkz*U5Q&Xz0+Ht1Kc zjbdJL6*{7gB*>(#{K4J++Sg4eea{#myyv$d^!;CU?F79i++RwnRV~HngNe7piN*g< zRR;+kB)_Sm_f_Pyb(6?iN-9GNV#rpfyjWLxY`T0h7#yX^ zqQ>LD1iOYHYQ4yPNL>kzq$YeuaCFGa>EdLpIB8k@q=?8h_+c}}tF|4MuQya5&8daw zD)~wLy6PI@pr}^d7W0Hh@p9s)7a+bo?izLo=f}j`vtv*S`uGoMR}aoBVga6VqYHPa zU8!YU1{Kw2kE)+TWcO=-tgUj5Phx10cW!%TlGv?=$~uC*ToSxjls^61RjRnYN43Y5 z-^5Yy64Sp?PJ-IwH(@XqpP>$dC@Jvgvj^(Ae5_u@Ela}CBbMbUm;_yBf|tEK)~xVL z%DHo~A=QyeEKD>{ZYYjM%zX2M`BSo!4?`1LMaHe40>}y(Wm~Pin6f_Ko@D=1g!TC&*`w{=v-qtzL{K(tV9-ZwD8+(VHxU7#2k<=e9*(kPcCz3s!>P*Xy5}SjW!oxt)AC9jy#LB+8*+baC?Ew`lPWBO-;9K)$#J z#Og~~Yd?Dhl?Ge@65R262;gyCc>nM3005vt2e_0J>hmB$rP6m>kbagLE=Xg@dlqOA z8gH~XKT5FMKAnM4iE)NmWClk~)}2#OXGx$y@*}6sUzf1lS&ZU{ zbX@rbNA#U%Z}~w4AQ#Zv#~-Vu$HAVQxj9m4AaEWEY`5RQw=$Mf6#cM;;#J0i>bwj$ z3EE#G>AuoWHh@_hP zE;`@?SVjxRLx$RA;)a4>?d{dyCS!k3t+Hav$oca}4BbH{3FTmoK(&}K_4ZLKi5^G4 zf#_YbM8?u5(Lp;L7^wj~oA z)upt(M5UMr={~i%QqTGxvd&a1v=g!3?`hj(#eepYE2H7OH0vGY(r*7E)0)I7YZfTA zD@sEPq7d{LQ{|NUW1`ML3l@HVbMLH828l~to&cjx+-ZJudkM>Gti`*ZsC^Yd&}uXt zZ%yUJoY8U+5~Z=9u;-T?fY^SSw7PII$`uya&edbSu{=Gj^}<#ayv7t$QVuW$_Qlui zp4zLLtMUUOnJLf6Pw%@mZr8|=-JPy^Jjdj4*U{6PJX&6tq#3LQP|Q|xYy5+e zFVRH{`_*k|PuFoxd{gjC-;8$C_KX0__}MmBNYX2mB&AWTSgcB|6?@5Y!$T@`o^!gY zi;NMeOTsAIo=e=Qt2!D$ThglEEyBrVFDUHYy)e}!)dc(CEaJd{GM;v7CN&TA zI2)JFi6bj*%vJ_kG1KQJC>yHDeK7>e-42w`xqCAP@?@>HhNxq~VL5mg!6ifpH17gh zVe78c@z&`N6D8ccKaN~zpLk3nkfm&)!Vp|DgH%Aq(8eZb+DThLTXocS`&Xa zfQ!!6sTKuo1h1&SCFA-)k4BvUEa*K4HX|SP!78qyIKwc~hve&|I8tS0SMG7nrDUr& zazU%xu-M*C%ma>)3FoV&#Fiv)qkZp&%eBh<{Af`%irGqk*K{}o@OG{drlbsgt6F3Q zzq991rvyVZx6&1ZzOO=+r9;I>%7u%NH$?y+;nRQVn7Gxrfi|i>zm?BZdY|k#_*cPJ zfRr02tCbwP`C$oEGvpbbJwC9vGeLn@Vw10hbWVJ4?jSU4`WAzy+5JlPL>88Z6Fh6` zDJGl!M3#u08awkq>55G|5Uyx#>7xZys}Iqhpc(F9V`43*Wy#v=B`znn=%(rNLwHv7 z0+f7HB?{3XUevE0OF*|`of#n0?hxj@#?j;Xq8_SzWDs!y#tn%Gu&Y~uw6ok48x+7H znLSlxW636|1}UwOTByL}+}=?*h9HM&P0pQYr8zq}&f>SpEq?f6wky_fQ z&utdf#9!Of>;3GNgcCt3Fh&xn=i3pTh7a87rvm?7{2Hzz!YSb zRe-|Tzde==oE7poyB(D$A-QnTj(GVnjE=m!$d-IEQ#Z#K&K`9e zA2-6is|X0hbbvCsnTad(8%Fu`mB67mRoIpihh_*wz#DvWk4zI}o>e9p#vewa%g95+ z6~-BbCe`L<2fukrgbpQ1jH<;V)0pLQoOiLmcK#pI-a8l$Kl&P07AwTA5+PVELPDbV zvU&+2kwl43wCF@<^_Hk1Ix8ZPgy@7Qi|8e*M(;J+>XvQK=KDOqcb<9wd*_|;j~Qdu zxc77KJ@=e*2RrFiS(E7})K0nc!Cf`rw3sj%?nbVF$zF8~SZX@z)(X}e=qzJvaWf{; zwBVggMP8rLV#;!ab=~82f!_CbYw5S?54FVWJk9TKhshlcl7VfnrQv|e8FJ(uaK4UY zymI_qUMGXl`cFiy@0N08h#^b?Z$H+NmF9uEHW_(XPDTq~1YNeif^URtrh zaAvYegg~I2IT29!z6AX#@BsK_YTpSs$P^9RW9by1QLGx5VYkuJNz?B?w|zF!TfH;X z5rA@h7pm62D3KK%*w#EP?Mf{b|8rMybPbc}sn>sY+-mWr$m&O}n*2`_70$e8%r$w| zgZ;`r+jjjc340VWl*IkfcV#JL%=g8##TecPp}&Gz#973>qODzvMDDDCW2Y_WF2JWZ z2cpHSK9jSrJ^WDca}XFhn}jQYvJm=J$JdwBIKbQfJ@m8BaSH zG2`jIz-jdA)8hqw^5^plRni)an~FhC?^hRb(x6*7KSam{3I_`?25P^Q3;xBseRAs` zfi#jSZIP+4Gzn;wg4+Vqw?Qfw(xan-K-*w?=6lkEDNX}olM-j(z=tnoMsr4Y z_JB3|48sM~>g?}p_xh%K=Ky`wnJVDhx`r#f%Zxe56Pag+oYgcD0Q{$A(Oa?uq7w$+ z0s*KN<1?q&%y|NE*AD?NRcr>El6Z_;f9ac zwn%CaTvwUrd{b{jP2UwE1ZM_`?~ow{POf&KW0~Rny|(`zRKe54Zjgz6-t%=h1w}RI z*4J99UmZKz)9Y1)Y5N|MYiY;Reua4@MFv)5H?>xwh)VgLeS=7B&s^D0uG6Vi05 z*P5uvc?_$(^5Ypy=yI$s_~I4)+OTm7vlvGb4OE}tB9rA&9sk22YnY;~wxe7%jo$>v zLEgsiEK#fQ{*LnOzzjYc%+X@bk5||_I+xa5U#X0(I#p)Kt|>|hS60UmNlOYqwG2UZ zvDoD#uJ%hXb3oGC08lg8?cUpstX-C**F@QXMx;-{Cvav8GXLB2zl{YRx@I(0F{CrJ zM1~PV8T~uYjOCy4wyi;u8afP3@yk7sU&KJUw+&CW10N++EnF;`+zx&2-jRf;P4!zv zt(kM*T3`s3zHsUPP^{pc0fNWmC5%|@1bR*o=MJ4|$GONy+|wvbe`A}mrc2cLj23V4 zDdUuSCF4;DFzEvC{j<?XI@8RM?i$#cx-5Yp+NHJP_fk4&RQNuBm8C8o}9J7k81UQa26~5J=uMRfo8k9K@cn#4i^75$8m=<=xu$r8m&z9yz0oTqwNnwY%wasI|BW9k2oxDmhCm~0k~BnwKO zRwuD)mIOrpC}r9xd7yOMe|yJ=R#js3z51pk^`xvJXoWJ(SGUAO^QFeJiZ#D&`m@0n z__je;e9cR$s4IUU8Bpm*47l~d?32&)sOBPrS?ex$!|T$Y0*K)tnZTbR;g#$RN8hUFZt!X?3)})g$X{)4k>a!> zj`(35-2)-v1zL^05pT5vqQ?wJwt|!UTEqd4g<2uPx(zo^sISzMp9)JQ+IcD23Njgf zw^E-6HOE&i3oOK%#Bv6m5-&tMwjQ5sgJC#|W(I1-j<^~uhp z)&4i)wwAn@^73|g6p_*z($GlxCK9$UUruwc|9l%P96Ix@^y{$y zbG5X~@o#cd@tfO|nDw%E4!w`~bYy0OG9oS-;*zgbtPyn_IB;kRyEZ*iU6MzW3(U37 zH8d3oYtb)(wQ3V$U5a7qlNJVi&`yd7GOzbedSkz9f5-(F5`P3J?VG#hd22cV*~g}Z zS4Z;?q-+nf@nC;n2yMXBhj=YJj&h*aKTxCz5&!;F`V6ff%zar2RWT&jdB+~qZk?uw zfmP?Wo<0*){apc`BpNg%tF%{+i-yTk8_B#B64=%X%dB+#a2F0GK=?;C)-<73(t_Ja zen5jFXqC%NGxc`lY3GM!#^MoTSDx&q7fD~tGoSE4Q)d1=SLRN@iyk7x&H2`8nra<* z^^t1q0?04OUMheLm6(BaZCYTS-pTmBV)8t*;hd*5xtEkWPx~`7du>3X4@;-i8ZF(+ z68o*jQoKk{F1hG~n0=o61RUbMVQtpkzx@iwo&2kf27INE9&Uhl&wchjmfkzgT3X^C zhHmO4pOCqAuqG4nmCI@F3V#n7mKAnm<_1$)Y?gC*_g}IjNa-6_!juM7lz9~$uBPeN zQUt%lvkhp$shsb~aO$&K&Z1ft@stj$!#Go^L$&`ah>4#-4 zS>>0~5uVqaOzSlSe5Jr?E2ezcV1bu#xVbnQ5L~=lI*))6h2JUMgRL0;I&V3}=I-V& zD#3{)(0chWCE5DPCsD-9DbjXcxG7R(BR<+f+5r>nIOX`@p&Kx39@_r`}CJd z(g~9p>HOK-(H?mZcEqL7-R-QUFJFn094kxg|Byc=PSp;sFp#IO^+kK_oLI&Boh@9`$)=Dy$SMdVX*1K6;&eo}KF|8*DqWWjKOr#CX3nKR4Ns zw)YjnR`qqGq?VAb!k3S}@~W;nBwdKhZlRnrTO6mmc~|OjA=4dBb7yQE%s*Gd^4?W4 z%4`0Vye1t!JLB_>XeuTvVi-Q?;uJJL**Ny>LKV58ePeTSW3hMoz+1)pgWT2OY#lnuviUqD zqwg8PXas%_f5gvu!@OEvos8#?Iow!Fa!9cHNr|eag6s`javspv{lV7IoG@OoDu3Jx zBEHENH~5S(wuaU^b=X~H%Y@aRT~4+8T)JD- zZ0zVP!q_(7Ht}cfCSF}YEKW;AnAj8I`54k$w=S?IZ(f&huv3u(Qnr-d;qMg4M`}G) ze;N9FRm%)EwdDOMweq3?^cncv0GD!G;QLY@&#w!-I>WkY=}#EKcVCXkw$-I_ATC$$ z1RSw_MR>m8^3D_pd&qHg_5SuY zB@@lXek{8xIrzUtiPtgVG~D^TZ|tVvxRVT4!03(U;PZZyw{2zL+zEfdV`?M4+;6@c z4~W5w5c|G=l|C+BD7rHGRkL&chUU56CX=RnnaK8Ojl=~Kr&QQB4#d9{XxhcuY;h*{ z+8B8C*@P7B3uApl2%OYrm&5B-Pid|^K+S<+3_tmuwR)+qTRyPK4r*gc`rD)YKpk~Ep7WSuBYGl0qV zd;sw!JJNktx@`zaPW*9}Yf^Pv&7kT1b+@}QeoCO18f8+SGyv;G0 zM~H&32lcyuq53KP>Q%vE-qv4!3|(ePOmS@}eNq6?8E1qbg4UI9QEd%|Jm%a>sUT9w z{55JU6p85~LnL#B6VrtyoTW&6CJLzj^=Wf=QA6MZ22s75?**fLHZ}=l=1-A-ei|0Q@ZOq$~TUn&Jn8xb8$R$g> zxa+Z>c2$@n$-mxi^7@DrI$`XCU%Wq-%_ckyJ|Y%*;h#2mO%Zo%J65%9@mhkr>2W|i z*#vn1W#IjdwThniHNT}&P6{Uu9o$nCJ7K^4weNR6oXLG|*UNI%k?oeH3=}78m0Z7qhHqg9_N$^o?feL76h0A{7g|hzitz z$DU#;PG#bduC2F75Zu3U3>A{B8Pk8NAC^p%4RMmzjNyk)p0Sq~Vy&kO!t|sX{IpxM z0qaB|>`dlo-{MO^Xq#Q4;Ixvf=-K-<^55pu(@$*>;Dj=CVP>RKFYhU*Y^vo}bBp}m z*8L^UwJyb@0y_HKV3q?^|AY4I_CTqnte*R>Hg~rT29lE*Pg!|c!>SBkGKBYTHc8M` z#n7GuOQkDP$n5@{?cAqQ%82$N0-l=eB;fOE$1H{I>^Py5ytcZrxC8FJZE>CRaI$sbHg>%Bnu z23ud7)m8{2k9$e8fGPOps%G7Fa@VKPBRNr`*_^=7R>Us@#w2+{=sbHX-$FxCsBllAQq0{O<@!Q;djt*b;z2{$6%M1)daKqEeaI;PRDnliz z6Dyla%O@kK=l)$G=&^c%XOGjza@ql2QO$> zyo}cdY%9-J`z8J`*aWVoS>w%zTJ7i_Me}{Vm!P+==8*3-CFu+c%pYX!oMy{y(kU>i zbILz&2^Yo1UAr1bMNycET($BQ7d9z%dfzPvKZyf(?y5gs6?xbtf3hJK6ACfVCQVUp zmm|fzpOU)Id}11Wm-l+5d0Cdv;(6O`>yhQ%r8Oi!H(7IEW0ccyDA11X({FZgu?EKq zLV9m=@Z%#7{jjV^61x`#`ijwMk9_`$@?49gce5|I+m8z7n~m@xbsI<~_xjc|mT3hO zXU~S$-jtX(aa2j=A$u<2>SnnQ=gz)7QTlk{qMpql8zfdy5E(#(b5BAl zOa1vqK^;5)b}}>DYv{yz!kiA5>ckq1`8V*%{7KraTXEM|viL4QaC({5_)gK8>SBdI z|1&VxaIT{im&Bt30KRRU&F5O*Fjl&2DvS}>wo#PmIR#q;$FCLOhnvlbilysEq?p9( zz5Ue9I_omK66(>sMs`r5-OBCglZ9R>u|J-OPXB^x8cduWXn&IhU|ZZnpM%t7BX-uk zGp58n@?%*BmnSm6stb+hE}OD?h8kygV@^-w_u1nt?1Y2o*A*a`|RUiIol=B_+x}RL90{4XVi01!g zLAYo|dDav`S+-cC{t6%XXB&7fe54CBZ;ZsB9l+1Y|78_%ociC+EyS!6%22!*WeJDq zFp}Njnion7N@7wuE3qcIDZ3Q?MNkF0RJF$17e0JlFTLu$;#NwpvIVM_ksPBXk_Mxzt)*uZPgz*CXuHqEn26M29w%nPpC|%d;Mfb;}*(MW{ zg!sf7F6An&U#jfIB+2=YPP&d2)dGkOTg|k~@Od02otG*SxAx@M4olP`%gK(Fo}K&C zu5nwW_}?}DVc`^?p(lM`er;7Ho4>i#F%?68NZH+xypFE^=YH;S8vNF?aC|rOYYP3f zfWkh-1L{87M#ar+p|)D>6tqEqvUEKc=i95H>YO~D9+5v)ctd%Cn52(f^W_WgcCH9B z!6*EW-kr1wZz-K;&*ZfFSIbX#Uw=Sq2nJH>cm}YjMBH%(B&E8liVbnH&<-1mQ9X+_~V|fpY>g*dxAPM zG=@on0Q6T|UGM{(N!`cZhpk^YBPHxNO_M*p`$`T=Rxe;}iAKL(5BN&cdpt`tBD7Fv zF%q&1XCW{6!ryBr+;tnv@1~%8E)HHBxk%ImtHJ9I-Q>PFO;|0<=8A*L28G}Pu3jAN-rFA6ju z9U*iW!a@!daN80p@aKTS%gI%^bvq#=O)BenjgXh%M?#8ikNRX`Z-k>Ee5k7ptOycotn<?Ay-9?~d}Y7;;IH&$G7lVIODFIt=J;N# zbk@E9=W3IDO4eXeX=TbGrPB1z`y2Xj0Tg(R`_6ihye`FKqjFkyKWwLPv#4}+>>h3eFyIpHRB?T-%*%K|A%Pev=| z&B#P(k5vU!CjRnHdb`yxZI@Zh2t>&#Gn3vrdfUMiFZ&}_rdxT2xAzx(i8G?fOgh(R z25~$@@IA~s6FO_G@XPM0+*NK4!k_c2TEa9hy_@uN0=t`wbNW%CV;DXE%)MsL~B@6#q4Y$S{Yl|HyKLsb}hZrDNhPRNjju?Ho!T zBJbv@GEn`(Rg}^`%r&c12`=BT>sgwqi~(a@>2DvV-L{nmlAyKK-(Y_Y4F}+zlt51) znWy$rcQ@LzxNeR7p|351pc$fLl7Nl%A>-(c&YsGF>OU8sS8eoM)tIO5G^?uUD^tIwBE_TgoJBGKF3&U=1u~ zCOhGK80*$`j5taDR*??*=U0k9+sJB$OL1@Ne7r)r)J3Pm)thpsIvkL{${#yadvhGS z4*IHx?7phwPmdN0IDVy%KHAgG_q1O{{zfao^gM{ND+1;1!Ap%K!lrRR$ft?ww@c=` zNJ=nW&c-o7YY2UXr=OQXF!h^n#Wv^qOs)TV-e!HN; zDLTBYlHef?!%Ep+k2q7Px=V(iyy;yD-c#blj8?2~(A!x7SRDZ8X|I|ZW<5?PD_k>3 zRY%4N3Y7hL=f4Kym7gEk7R?qwSvzXhMyUuDzz~f&3EgiBvVOd<7Rb~j-8kCKXrxjc zcKnc%Ui-H4W&2Xf)(=p5#Z);y zt!Yhxi4LqG-VtAKv|e^WQO)cCwWN-qAjKWO3_TcDEHYL%L7NG?wqy`WDz->Dju*oYb5@XN#%r>4TGA>L*=1KzYU139qL*OxxvhOK!lJ~z0eA2~Xkd9Xq_`jkE^NhzFm7`HwgKK&W789YEDO?j3^MTR|_zKIuE z56RE!g6*lj^6t`lH@go#dtsmDM0KZrt{$qw>Qk^ST<%|3#kNljze;oMu;G7p5k}s% z97H)=MBe{SFdC2sAV(h_pVmv@xfSEIZE`EVYw!83FTsQ|29d#tL8s}wuWe;%nlW7T zylu`*&gj5Ofusk}OVYoekCoHAYn=`QipsA)9j^6@mT8bOxEb=5WWS5*$Wh-VLPn}dsOnbr3wScx(WNs4)X@SC62l~^%FIx*X{-NtV zwAp!tQUqDk63CRN`CzI{BK&NE^^}B54PSZYdybY7drUw%M|6`l45C}e?YX8>b zFK6dIsYFgtu#~RZh*s)WTS?S#YZq|~hz-f{N=0GVI8ZMVe`ef*eWn~f%mMh`p}v}g zzLSQ9x9hh#e&p5f6kaDWQ}9fsYV$HwE}`rSQ0}{oMtxP$h_skTJwgl^CJg)Y9gBGF z-5flQANGERFpsji6vt}8s@@zw=gv=SkEggj3xK2s0B;`^=@cxhz|g+45qrqu7JxvY zmg)Ht`cOa^Q;tSWn?J>nmQu62nB3R$mwWb>leA_Zv;eauCk+C_z zS%&^f|Kh)Z|ALP|990BdGpVd4X0rxFPp(XBD2dEL1?&6 z0;b`wF#mq-qQ4|z9vRfIm%~psvXmVBz3lM|OhU-JZMH`{9zWo5(-#LRnE{~MhFXKb z&uFt`j~`Tfh_q+BKbR4@QxeINoLCW9ELOk_M%n2Y>yeW&+v`=R*4g zoul4u7KGhb!zeT@pbhxw;)_6o!BC^wNgOCcsv0uL?^U%$aLef)EDErV1>p?;G7d-S zp3Ku=i654X%_Fn;2dcILA3WQX>tz?h1#{)aH$kk{$k2lh7?JPuzQJdRP|Rp9l>N?S z`p?ime$n1rTa`}IQk~!3vYam`yepok_`$~a#ugFyr85PKi?YkBrn^KWC$ICzO4N;0 zd9J4;C=%x<{vV#G<;(dL3)!oY>sTjt>Ak)c&f7ak4uMmpbjNCb(|v3uV+tYWV<3rr z*R8=Q?riSFbOSp5*s88}5(qPY>(jai&%3G>zEqdIVE&_4wYl=nn(5L<;bR1KxfC0= z9uOLtK9LsU94%YXl%fAK?N!tOtkfC`a|->`}n zws9aQ-O-7Jouy2ImOW%0nGJZ>8I0_C<&5=B^v} zBkpdmm73&c^A(wzhV!-M!dPucIe*o74o8@sMZ_HQf>V#>RhBG6XbDfHrk^OyrsN?$ z0w6vGZ~-3jh;HIT!S308Veq1+mXgOflPmc3@loLA&t>1`M`bwZMoz{zk=IvUn*iK! zrU$W~aN}%B_VK*oB@gLwCf^auF~5lpt+1MvZ8*;&XkKu(3xi*O$B&RxV6p&nNf?CF4;YR5j7I76WO!sZ!C=+J)US zhaILeDZ3f0{=foX_|LjKQEuRprYzQI4XZR$j}YWpk^xZ|GVdgBh~mfI+v)&$A|C(# zyu`H}57cWtL^j=*d2#ULiwXogr`4=#X<>LKA3rfP;`m%>`WG8QS-w}iW97oY@@`q- z$KxIY!%~jms9&`)6mVM2kAy@kn^ZD#Ofp1x-|OVE*gL;o(TZ3>a#v>}_sz*OqIM%| z^cnKOPXd(etW5TjVb_ovLwg#jD>G3M0_Mg;4fbqP5MP%E^cIyKGLmz?$-Lhu7r7_W z!COt1%f7vF4+B@a_iQa?HN#>Me_bRng42mGy3k7AUR^ksqs(tBBB9jQKSW-9VMkDE zw!84`#tTV!ZrFb1F>WBdP!sdy!fLmc@}O)v0^xC09b}FPj7Q}~DqY--^7$g%i%LBI zFq1qP@)P#w>U($Prg)Ow@Yd}6u~^ZO^oF;`M>?VWpU%DV9b-@>l%lQ&-Pj zd`jztT^CJ!6c}##w%)5p)y@*EP}vXW4tHmI;|Oc%)J$wg$TKvSb)R+PRRd^zz~YB7 zVeSX#_e)1+My%_Cl8e0I+|+P`%lPrv4f8UHo0&naZZP30Z0h9oLgnV})@S09!{gdg zEQrVu2-hPT_LgWYe8RaujOKJ$A6E%DpDM1!>jBIlZ<4Pr9>IP{fXD#yfyUqtfm^p1 zFFKFdbxG*BBig|UIl?U-_{V_dT+$*r?=)ponAYzt-P9y5u!aHD+ z3@#>Bwn$w47jn`ki)%VniE;S~?aK$i`~~x5{wm+4rMD>FcFLQXz>UyhpaTf+zd^L9 zs*2_~Q^!(r$7H%K3SgC=E5( z#NI*uRC8L#evbJk!mRC|B^AQb=<<{TNmk^+G^``)Tv5TvBr_AHrG z*w~VGVeN4mQ8rnQMd9Dqb9OnNu*kU#DSnuIk$L?vu z3;=gkR6ldwOzP+?fUw{sA^l?Ku^B^|`RWdf5Co3Y}rA zc8Aca-{g0TTT2wXd+I}H$_bm;VZR*Mr5ZtG;PlU-uayw)%2M_9Hofm|%N(`MTezn+YigK~})QT~0 z!VBjMO;Pv0+cp--e$zohELVTkL`2BA2@@U~1;wqAf}!u<+pS&|?-%Qxk{{szdYkuq zW&9gboX__wUlSiE21$xgqfF8J&l+3LsSb0MFjt_WWAM*J8+0h2t z>*{V@-e*Z7TcR?d=lVY#U*;PkY2;7tQ)1#1Zz$2FGW%woH{Yd(ED)BMkJfXn?vwiN z7{z#znGe?;>;Huck#a8Xt0^Z&kPZo{OsiNcJ8@;lEdSceR$F@RFafS|{jL3-^ck7q z?Q_0q9j*-P$!?&ZgXo$yeFeA@?Dx(3QPJ*Kt7I=-*JN=k$V4@)vL!ju#H{Q}nt;MtXleZb&psw2vdTVCW=u_&#OJ~w47k_!M32(hTx#SH41QmU6{eFJ-6b-6|= zDPRPa94e_Oe=Fh=#xy%_u-Wof@BxdJ_@9>)3N1SXM!PpmrcW zD{bi()tFW9z?sQ`+-UYlqF~03+W{X^1Rj2_=O4Z913s z9(Lupx_DLV{eUubwe*PXREL%6O<1F6O@=5)07!UPROlBGdT13|C%DfV<|1m&>KRg8 z7e=dK_lb2DKq+m*d39AMNju+7B?0H3ZE;?6P7i<-;p_66me^sp5zSYx>z*QSa6XhD zsHG4TojA9z#RuP4;gpAi4q9d-{4yJ29eHPGHZLy=Cs{48etBCQOT*H3-#L<8QJpff z2n?O6E1ncZ2h-RT+Vy3jOuvUmbfnMvHcHGj8{0vTH*YyvUH}#9l1_6DF6W-aR)PF% zQpJ#x?7v9dDS5^pqzM>95?R8Y`cXPfw5c!JgdmtsdmT?yCK_=H>K_j9PR>i)x_7rh zoq9gP0OWJpp*sy{^Oadrv5qE`nd3%`NS>8lQ$+z#bZ}dpY@I=%$W`#eycg%E7XV@y zGZvxW|1%fh`)5EL0~~<1^cgmkv~%>v zK;?I{r1!`o0R!O9?D{8;jekCI&{tDui}J>_3v@i}+PuZ`os79tu7mfthIN`c$s3Zz zuS>}Yt7Pk~<~1~^g1P+)BR>Z2Ea(nWeL?v75 z`UhrGMVbwNK-8Am<hQbk+GUaKiRYj5mMc^P$Bqwa+x11-m&IYo zYjRE3aS>ySt8+SMSo0LxkdFlK^?>QW78_SK^837QLOj=}7zIz-qZ{Z)p#%=vfQqQz z5kho&yUiOAtzC3Eck{p$&Sr0xn5MC_YZAzYYs<(rRR8xjh(7UpRiaAE>kA zL^NJNn- zTGuc~z4&>k%+A}4yG{Dy$w?}aOqRquXAu6ANENR4i2DM4< zU|vsLs6Sl7B6vm4&FO(aRmmkcq#?3gPr2yrG7G5zrP%x}kWf5t2=d+BP=+eg$3kE* zKtl)MPW`)P{1OCmSspgipu}DiH@d z?ka9qdQUqjkwsK8IfDbd^RGX4+1y`>`gcXT)_IX@8-5*r-r8q+ zlT;5$!625tUu80ffq(&4beD^aJ|4A^trMxFO4MYEqjMhVaREd*3Aq#)Vdf49ywC88 zmdlzCguQkl8aXB%)rLbPj7;Co0H}A!XjGsdD-zt1Xx*h3U z>WpK(AK1l{mm28NVbisqj%>A1rNnT)^(jYBei5Wy&UCcYS{9+H=M?IIMTZ7(0k||p z`lz9rig!^FeIVlXHn6#gECtkHCTGTfFHFE~I7rp~y84q8Q{p4^^2}kdS7_R_f@;=c zC$2rF2qXy5jC2A%1*%FTn(K7TtpI-xF|;8jz8pi-N*-sv;RWDmzXkJIje}8zc)kjc zXyF~?yN&6XhMOz*mu{l;E-jaXUM!aoH+^fPj-jhz_mvWq#BsH8u&=;g*#HilP+KWN zi%o@XOb^i6pye`Mr@p=iz=2kPa!{^Jh=(t!D3BZIxkl(CenN@9#BAn?AbAJC8hnUo&!N1ZmGD z(#D0%Ou>tj?^vr@;hnIDfrlmoU-vIqGqR|q$PL9+|Zp>zDDKYA3w*=U8VBV>i{oHO%3ByRr+Vy?H>f zf^mD2BRl0wg~O9>P4hT&7h3D(+)v5qQDP_{t%lFDYaRtrkRXFS?R+^4q1(`vA1;>x zerkvPf=3DrQ-slJg5f;*Y0T5tDJrqotKKfh44*p(pCKE2Y7S$luTkJ0ECM5Vv-}wB0q%A1GhSfKw*Jq<*YqdUry^w8zr%b!J6g=T@mVuR{8CVrA%n(qokbM? zp5|tW@k}mSq7C@Yx6P=>7q1bT8yRAe7M|9lFv3;75k7m9r;o{3_7pFsmS9>utSq!) z;t}hZH~4V*fJ%44w39(4vYqWLdXl&lIZJ4t4+3eHgWGBU$brL2W;i(SNvW0aK7QDq zY_CE6E-lO1hB!)Gxc+P?N-^Nd*0&~m!mg65qBTD^pW>=UVtuaENqWJFSzy4g9>5=v zf&sQphEpg|3;^@!6zGWEqNvmzo3RX4qoqUF12BR*+wDJW_G;pFYK>+{EyV*;pNzQigKX8A79bK8{u_ zo~1FwyZ;GOA##>E<5vz5OK7cbVXq~z@t;j&hIEneYf${;nET;F0mJL5Cl8rsK>!BI zfwA8bm^5_td`c3Gvr}|c2E)54Jzl=Aea2P}9AjFQ)f!Kk!uvPF^%ZjkzmB*|p5)Y9 zeBE;SQ++8NYJ1?5Ndh~DGap^HS9ivdQpN*^P8|}iJa7-4%~+gJJumOP+=9o z=kepF+jKUFoE;LnJu+U06lu8dRfAgSgD-xnCT5aFXpT?$3R$D@t){IRD?zvA%k)pU zK|uj^XpG;xH6M&8PqxC-5WABg`*4Y`3!v5J%`;Eg*)50QxW?XxG8ySoSuTEt#Yv@v zuJ#{K;X+eEmdTgRzwR8-_K06Q3;#}@WDLBwU0HOP^8NN>oGV9aPt(yy+QZxlIyd_n zg{eP`UvSK`R96zTBgMmI$g-~Vc;QjlIAnaY{AT6s)h*hq#A{bkQYkV2i67aBjosp7 zzX5GxA4@!zK%7$Dvd;BV>Y9P)fL`jKTnZv~r}e;aQSq~#86G*_{PLle#D_m9$jKDb z$)Jg3TI$?jq+!!gkqc$&tME!XiAGN`$Q^yH$EXd4>9r4$UF4Lj*E-tRfl10n;>!)0 zMMzp0#3+DD6xFWTs3oGN%v%*#guDcTm2e4Pf&stK$vKiOYxFB6tnmeRfb~ri?Hrff zKyS$6hk-%Iuv%k0Nb@-&jSq1c7w_;f(6~8sbcxw8s zAUHFGI=A|zC@hv0Lp){N%} zLpxdhFy2TdaTxeH5pz`?>5$GM`xEh^)Bw;+xWB{PY5a7=k=hA+kTXpu47d*kVP`KmMkxJbI zpMnwhXC0-A;V#7FIgodV^)BGr6OT%d;Fy5#MzCV$`KKdDz2>Fki>A|Zidt$MsNWGa zD~Ww?^w5G8S|PrLA2{(%iz)F~v~wP&CPyq8tY|tyi*hkxD4JUKF?M8Wp{V_C8EbaAXuAn*!U;FX<;m88CEl~a-9sjF^Tg#cf9L7>$|4{8-8Y0 z-A*FE*B-v;fEoEg>?l?08K(?R-doX5yU|UnzLdQT%ODL$2T`^x5RS$;SHcQ)ufV5A zWd>oZbNgPF)LWbxyCiv)lAZ_0f};)HxYVqjl&$jGCzg~qYlGT5?Cg~XBhc%td_Pb} zt#|Kk2h+RKMEvZe#Q9wPxeYKgI?bD)d#{LXEeDt+yzLh^{^+(5RFpKlq|po?|LtwD z_hw7nGA|P@jeY>rdrEwZ4K&v2{)2TQ|ATdi4mSS>)^WS(pu!-)`?!Dase?P1ooBm~ zoP$(D?gWtiEug(&j;B#8oH`<(s*OGT>KQdwB#+W7$KRZRo2y1*R*t<8$w z5lwkX2OgbHEcj8yHPb#~SVwfmiIp%JXvGGics7`CN6Zu$qs@D7o;nDSuB4%e4*9tVg!&9t36^laV_l#^ zW=08GUShG9+Ao0bpBS7hHH{=$)mzB=@BrR|<)vP<{d{qq{H8+_q>(u<5-%9Wpcf_S zc3w?rt8!=Z-#}_K(*YUAq(6XCsTg{2k>R~%cWR0_=je&2fDLf^O)19%Dx;v~=TKT# zv&er<9Ix0M3+a0dgrJAG7$*4qy*cK%0qNo#q2nSGe@wmfJN%GYYxWW1ZLFlZ(eO}Wk_3bIeIO-ZBrV&D1x}&i_KHwLpJSz9Y+)5i89CAJ&&&FI+rUptMU!8jt4p!X| zWPJiVAU$BKuu(h~$wfMM@5bJ6g!q zaR@{c-#7lSNTmwheI^poHV%J7Vi}JDFrJL%(0lUAsDC5< z?x$Sgj?{fcFa-?35JJAg?J+IUTY0Pb^Z~02=YVk&x_7J!fEU#HNBZLq)}Fas6U>IM zmg(r^Ro{Bp+)r_x@gnYhB0$`O*;WsA)yq$KJKsIjf2r(ae;Eqe_k) zf4APkOY`?9+~ZEElwOilZe2A`>~>STxcmR7=%8?4bYE0?T%6ew$u|Nmm| zJ)@f1wuWI!L=coFO+~thfHW195}JUZqM&pLh=7WU3ergkRRuu=DI!e?$f1KYDG36i zbWy5EDAId})OYQ8&pG$K-#O27#`oj<@r^r1ek414uQk`~bFR79WcnO)7wNUKoB78( z{Q7t@dGlsXZkih$UNqLT6Wg@3c{7|~bBV0lw=b7dek9TStGq?XRaoWX@X3B-$=aD(#yvhZRTdYzqT3enIQTw$=;Cm=((qo75g)&u$XZgf~e)4OfIM7AeY4TN# zAsQ#pO2F@@*Niu+opMVUcsl1i{ZRyJXF71~m&z4D-)HDx&SHMNeU|C8Q;Y)ap#PrN zB1p7=`y*DFFRJId@B*(H>t5gpYgo%wr!WNycLrLGQ^?XRO~P}9@w7KJ^jPLRwkxAqvTw0kUF`VQ7Utz^2ZJju zM0S}ya=wZap*mmNh-p(AWQGY1K0TKcK)Fe&PHC2QXl;*+rhoA0K%S>k>&B&pk0a~_ zon?a+o>hDcXWmL1e=|Q+^z$@>#-_`&16w^k?h`LZ4mf;V6BH}}E-ac_2vx%tKjAJ^ zbC~fWSZv>k|2TzgTx8X#w0jxp-J5nwquxRT z4!GXUnvc&hmgNc#rm=Np9bT`PO&_0h-$OpmdT1Mdt+hCnXC0job;0;h!c*NH9FQeqaf<%!joGk!!$2EW7Z1svczs;20*dL z`A4I2XEk7L9Rcc%#J8$bBYw`8OW4Q0lZE6;bvfe-^hb&W_+K4i&X>rPe-lPZyuZ;1 zmr8qdDIDuX#>ARiV~XBBbeOrxBqppU8dCI4qJGM9^-fwdwF2}2M0Rdxvk-=;GWU%& zQvy{Uqji{{3?XEAA`td0_d;O9Jb({f?60_gwkNgzfZH-LjJUFNFGZ@YI5O*Pzj3BJ zjDP7e8;}L-<@978Dl9vpG;uqtjMVvM{QfX`MM)l<&=iJd%{GpsHbSf)3aTu3At9btCeCu z7pmXn#TBbA9<-}1ow_$udM_oftu1;f;iqlt9)jzSN7f|u;CuY{&Z3g)B$4D#E8WzI zwUOeUT`wnHPJsii5>oM?+yXM$F3Vpt(dK-q$L|OQE$856$!%}o{%`7pBB_o}aT)}o zD+}QvZp8>?Vbg=g(DtM1ga=wAO{or;(n;VnC3E~1!Y-~gQ=l$y9 z1wkIIdE+eh^*(V2oG>}cR$)83?}jbb#tikAk=19y!meJ%Upz-nD{b?)f#Le%D#fz_ zFRV__9YpoL@R+SmV21%-9rUdu6#_Z{{!XE}5dY^v0~dP4Vg$|#VXK5ZbUjjf8TBM~^PA6R&N z#o$uZ2>VjrPus-sVe<~O_tP6Gq1%WkT$i~K-6Nc!QIeTf7Sp?|cKbY~*yw%YdH_=P}rURc+@ zA(1OZnMXEi4{eBA=xcoAqi3xCmd>BXY%U{Y7{sVn;>|hrGd5s5iXbl^WK7ZL@S9F( zxwU)Z21Pz1YSbflHx?Ydb4yZ{$`- zUr6;(Irh6@_h1%KVMnIL8ij>ZPIErow^%bK&N{5nH3Jd>8bsCK4%F@Y@@9t5Yax26 z_naHYpoY2km{6Z}eda!XT$+*+*JYRZ)s?8Jb{ls6=ma!>>+xwv*V_@#nsVhT7kzBm zg6%8As@RLqrIz`k{6E&7^Q)2`v>j~Q+@mvR!=*C}d7*R52gI04x|bYxNrhZ2xJ+9A zA|$&&C#tHH zgZJl^-L^MwC@&I|9&y+D=XPt5&D5bPam)^r?fNYLQ+e5An;u(##O zQfnd7d+I4q?Q8Ows=rW9!EFZYnq=$7V~*&Q(zHnAZbm}lOsC0qjbRXdWacDYv~)2! z6d$!9IcBw8x?H+4uaKkOL3(rBhPocRj8(%b$ct5c{gQTVL{vH~rbEB>Qv&@#pQ4Yo zXS=4X)VyUn+!Jf{685j_yF8Evo!}!QfwKL?cK4x#F!~EHrFg5aAZe&Mbv|K8m}D}`SK{%k;i?&!w`^6CGdJ zgs*C);SMU9vLBLO`k>pdFD_~DLxnl+BXMMKglsJ7n)UQe#3@kiPjC6o ztEeSd2Au0=@?$skNo819)5GE~t6z{O7mB~kjZUw;MZM!RlADHKS`0pOFAD9K_XPdc z`1`Bf)uFN$zG?J#`%^Zb&!wI!wG)fIU8RW!8?%iTHE#s5M%ZYyxlj7frb^!N_iv0J zQ*p{7x(^l<!+0&gHmkl7v6QwNh#I37ruGS1Ylt9tUpc?xQIQ4^_|`Q zy2DEsySCGVyxmKoSb?=S;byoos7QQPZIgcKt-Zm6W*GiYdiX=wPm!zUl$!ns-wAVz zTI{1Ad88gt;SE^-=ho2EP5IPti~A91UbbT+XBw6Sqg?i`hGSXH253_F z5o#apDq8^R&q1Tk-3B`N(4_s3G?<5L4k}=iPceAH>Q>pJmXZ%~O!4bF!2;snPI)sJ zgyYSfyZSIklg=sLJmPCUI=vdfX+cotJReEdJ0RC2waKgSGRfyO|NaSmf5$8_dNIn8 zK$J?5L1|l8AzG(}G`19{mZaeK)!Mw%;MQ&q&JpHA4~FK9&(cYUyghWIrZ6s+p3X8d zLfA2N=HNzBW2hEuo@k&Tvz?+Qx}J_G)*q%>uhMXp# z&X~`StEZD1JTjQ`G(Ot1D0lOstcl=*;@6yeJ$`OaQ2;fzfv2w#@z%il_+Z@eU5^^PP>3-rGB)RVSA0=trFfDhsVn2cOx&QDf`w<1k^p)J@=Dsnns^QORE=j@7Ep&GVFAw(0%r=+l zqC(f63q3}11PGedjU0G`Xu^fuEbNuimsiO@Ry?gkD*ep%`7 zZzApu;M3&@(%twG$NaTz+kWzt(mLn!BfCD#wXG8_HYdvS}0QIO3Vtd@ z6+o%szIMz7xmPuuIK0qxFW!2I-Iy`dYz+ROP@2z?i*doLvF|&&5SD{$TVnQ`)cwV5 z^p9~)?s_If+Zb3?$E%|cCHTWQjO9;6T)2F(rG5p*$?{wN$u9bsXMwz)_BSs$BROlF z2B>PR3N1mvOGy+S%2;}gW`Yag&x#h23||pHCuAF_9fT=jwA5n99^3w`PB?jsX zcl{{?*p8Y)qa)*TE6OtWI|JI)L#IETf431Lpv%?BNg6A~+W8->UpPY-PJEfJ@|vZ! z38BWLF(M}_PwBd)C3HWQP6fFaj<%@xV%8s+Y7FQkcaYm#L}4G3+CNnhIP7r1gkkq6 z-_ZiwVdjf@HbHpZK&}0i{x7+7KVU=CTBZCC-QoMdP^dQgLgAhM9Af*UHrd@?RU0vfE>CbpYd!Qi z0*?Juf7tre&#~kj->CI2GTQR7M9WO9XkMxYXPp$&q5CO{{U^7u%NI!Uj%p+YN`i60 zg^(NmipZZhowFJ#0sd^odF_ic0gMAhc^}t53me7LX z9loIc-`=ghT1o0&97PF{LT`9sQDox#x|ej{yQV(M$Stkk!JL zo9D6B=T=RJFV|=u3{j5)XI;n*`Y~;uG_AcjgOy*7oI3&UJ4^g1*(cg8q+(W}!$iHv z1uW2v%RN#=TGZw?XdBjMLRwe-B#qcoC^`mLKC9ZcCQLru(O$~QC~P0!*;&h(qmkCN zG-{P?TRX0Uv3W!}jW3x?1f3?LKM(d^1wOtQl>M;!aRxMHccc{>nf>A8XDpfQ3SfB7 zd>*#V%*4s;A~l^nT!P%&k;W!6t_wRsJ?4DZ{|)S}cf}EB?-K9 zV3;Pi&pkDzwsUPNQzTH|ap`E#>4b}XXkwEwFXv;75lbYO=9roc)zW2}2j-T{jGHw7_NVNoD>KvXTGKwdrYopM;&v1fsWhgN33!%JIBCr zHBNN7-sv3njqrwn{&2bpUmK1I1mz9rzwkVPg3@5P!)GLw(C)Rr&U6hOe=5k7&g)=!&WvTFN zrIs2(%%#!vU=?59s zEgK`zr`v9i+ub1;w`glH|HwAh)fhGA-ws~7MPPA$&OOp>EZiu(t)v;aS~!IoR8v8G zm0P2N?Yq=i(&oE`hm)#fy}mWawq_A4p=_SlK8ieS59WU@`<7AkmS3PcN0xn+-K{jC zDBEtda^a+wM!(0K3CD>S8U?$}oM*c0dfwk}m+gLtqtkTOaYraj2i42?ToU1VLUeW- zv9iJ2P=}0dgnY7G-c`TD53$P7*W%7;q3l2AQpP`geqh_JQS?nE;1ZL+LS@8U9Hw)o z;tO8;-Su=HvNgf+om`o%MedKc$Raars#ocPfuFN_a3VilLx46seu{oklJLx5AN#Vt zWJhw4nKsVmygZ}Ahp&?K=Vq13lnofcKpk(SBZ_0>DfpLKm zLE+uf?l-n#DbZ62=N5e1D;_i*Cdnrzlh#HWD|-bvKDd9)Sc(!Lzy(o4ws&?HNEeLG zHu!Wn;&{lZ=laa0dZrKC zVm8UFxxKz1!OoB|AFL9kqBru4%zWswUdV#Xd40Oax&i%ZUNc;AVg|aR{c)kCr4v-! z%Y~Dht9Rliv~%=%)>#r29e*0YTEJ-oFoa_kqLj2zc|hLLS?v(B_sQ1JTv^}J0;RNt zm#plfjhq?2Zr7h;wt858rNHZ3N}wf$7-jM#OHd_@JW9_P^l6?wSB-JI_pKi6%OE57 ztpL)Y6HpA?{&iDlNSQcRH%1w&6VbddMmgf$P7W;!6sm{cLrnNR78n;9xsD=O%yXmH zoi(=g%lic4+pdg16xmKW^eu|`z$pxLO2#Cu_}z8|eE)o{o*Kp<2(Q;4^LvS8$atis z8^r^Rt&w@C3lVNUQ8$J_RFA^n*k&BwHt|M0I$Tpljmtt{?xSsvG{61Xqt5T;8Ikm|0T?6YYP#Nc3GYh|kdb&A! zAmGz7{d)5%bzQ#!KT$63FCTy-1y>3oLpCwxtgdATd)Rwt`_}H33pxmEyF04~AnnSU zYC}=^WCIv^erc?3d?&CxM93^gWMjhkR`sO^a78tW021v97- zcb~n6U-!K$tqBT3GCfut|GE`od&CyuRQQDqN*z)@=NAEwUlDlxVoH@gF9KnD;H@i^ z>xIJKJ^qLiqm&gLAC^zIuT0y{ZI{)#xzD^u?9Jm$ZxOu;n{bwLBa$|#2UOGoUe~;N zF3h9k8+9mJfGl)ei8xm`7jRM9rH{}{{uG@qezVrMVhh2!E`}?I8SUevqk!Y#zMj0a zBvgXvEIj|+ZE(|ybPUGXrT>eMouWOd_-t8jA|h2#2LrJzM$ zWbxsOUKd$ON{P{!|D3H*X?$VfT!xWge|B#&wlBBy9r;vs%3AE;j=cB2`bf(iV^XZe zuF)@w_Rq}wfAWM3n=K|En-<$v@Js6>sjcOn)x9Qvc)`n^>DZJNE;Ah~te3l`P$$JM z|5W24AN5-urFyerZs*is)HjJ40o(S3UbS)V{WuhI)YWk$VdpBZ>T#x@n_b#MWF?A}1ca;jBn^JAdl! z&50Xnj#jmEl#_@}7x_JX(o(Dh{Z^Nqk9aLjEqp5|G_1XXA?gh#8t3BaprI6pb^w&`_SH5YCKs4f){Z6lI)d<4EGaz3WJj8CinvRdAW zpt?x842MRhcRr78JFopopCfVoZVPxX%H8U8b#KL@*McDnn|tcUEIEuE%&q01vSJw; z?dd~CWq=BY&d)X*l+Oc&kVTTyOK zSdQ|3@N?!v#&p*Fq#PLxw(6K0m8;Gq6l|6a3WhwlZ0YOU=)&xqH=Rq_N*mwMnDbdZ zDOjp#lA<8^QOkPpS}}w-9vy|v=~v?3~>13m&}EA z7IF_Eg3VT)dXG{sKo&eCS>gM}PZ3V%3eNguy~sZP5H-o(B5hVsbHy&Li7rwVRZ{aL zeX?xW1<@(Fap>!8ZHds9HmOUaVCdiqFd-XEaXsVuo8Z(7#b9e&u+8ZK0?Oa>N6=ZG zy$9@mPu34P^}}Y*^TY>>AALA_y^V>~P;y;vN!zK)V3__hF_Han#Dp+Ad$mUV-TlxY z4$jh2==`jK(WRl&iK!Z1ek)xrsC_Pw0-El4I17j}b_7w}fG;1p|G0_=G;TQY8- zFh1#cA^D+CwpeR~=+9V=efw(VFxAN4n3-dTFIaA_SHQ|49K?@pPA}vx|9%$0bugSK z$6F;c&gnFAS;@M_f>P;eH(zVEqUwN$IUIGz)vh}8tnHh!*4f62Ij@4!D*R3@;ce&+ zNW?gtFXP{JwM(cx#e$dG!=Ko=#v&cPd?qF$`vbNqm&0OXG~O%Qqf@{2Igz-yPl&Iu zS5bB^SO!J`!vnSMlaB%L=*l3ap!Ll_ug&G31ND30Sk@)F&$eL~0+Wu{xNc3rrI)_v zFL{Woc+PpfaN0~tL5*B!CrKWOeQcKo){hf~`wsc0` zBXeq!Pe<(FEt!7GUIqogsD>frx=woCwb_ZzcpllrQ8GB+fk#&35_rX(!1r~N^X;aO zM5-~z1_D)(C;o9HGnk!+>e(FLGIAs#53ixxKGU{`;LhM$4wZ#4K@F{i!;1K!ACuq^ zSyV@$^#CY84m?jV*FtgeV6BfbOdU3wX)^vepP=s;hOBW*pgVF6hu+q0g4uCTE;#okJ&voHyI+wc=Hie#isa$OU01qzw4(D zPT8oXEuSd@xnY<6cDp>M-9n=MT=T?vgQM5l{1?gs!d`YJ@vKiZ^w>y%{VldoK%Pr7 z4yU05a@!KX5C5HVck?bc?xbaG5CwJ=S!Mw{u7Mb?@KcB{ui-5vCQg({!bkxk%!kx? zQ9^ck1dGAT&@qCcN%@&D@h)+kU#4Z)OQ)@yrG>9#t-Xdg*1dk1z(}=TABmZ7_TmMp za+giI5TcIzGL9u6F3ylXKaf+<45KXOi|lt_N-R&){I0Pf5IV(eCsrcIWS!XX@&(c)COs#(9orQr^;GICRf=M>;x#i>nmCUjzt zTxFWKZBtYw!Xo1drSK+^<=uj5+T~?8nQE-G#swqDg5SL0+5~}g9ZX`!4DrZEU`tT0 z3!uB;uT9-^Gy(mAyXTqdft2JP>$00Nuo|0zGl?(-`fE$KVNL;dVVpgJ^<;3i~&C zUL){BvlF1skLiqr>MpWWU_Stj{&_Vt;qfiQAx*rVuUx zJTRV&Y%Xk^U)K7s*7`S%@?Wj>UxWBxgZN)x z3-X-)2VYC5QWhiTc5y}){=ECk_6f*)e1;>j*wjAno?w-#-7bqeO!P2hR=1A1TdRyw zCb!4kMOhnG3?+3BCzp@qrRsqD17UVPM=)-bo7?2VY+a7!?y4QddaTmDt<8!sJPKtP z{sA0Z4E!krdkSoMsW3DxJB`{7H=5~Lrmh(@K8UK5z^H!c#P3AK-F+W$*hRq@@^ZLt zaI2t>V7^9YdmJN?SfjH%C?0;%NNVSOqR{dCuhe&&43<^nN$ntr0stK0RskIXnkiVW z3hZ+C&j3<69Jo({7K_8wJC~{UY8)?r2cW>#Xl8ICuA?bf=o)2Cu9N!el>=T@N@P|@ z&8m^~!>?kpCKYmHas4bc6st#P@W+rs#o?#yQ*TbCj!n7kL80aXq!!qzZjj?DjNaj{ ze2~L7s6C?9hUKTA4nMZ&s29T^q6@v!7uAj?gWY{$v@)0d1E$v+NXk%zQ_07JTeC~a z-Sf`fV;ymrPC;RH#!kDu~D|=ZUhJJpz6IL&Xhn~k!j)8`Qg4% zI+`yljK08KiB!kTiOeEuErR^br)1{8F)T&K4EeE#;n`80Zx;otXL;dJXlG!b3iIyS zQx8APV2P19t1933?#;FjEKM!8ncl{|>CPmCzGLV)b5N_ChMMtb8~HHRDZ<$*5*Ca4 z(D$`d`FEBW=_?G0&PEyIDOTHLpDbIEG42Y@g|^;O$`W7}Cf zPz~&-gbq4(l>fhQKgawn5I+=}VukTu{w^SjL^6MyGCzg{Nq??HIGSUbWpg4BU(-ag zNb?fI!GJ85HRVIV6IRLIXs5eRfpMU)I?GX7v zYk0V0=imJQ|Cpp-9$2NlVrz6X7XN!dejg8rB1i(J6}$ZJxBgnx?L-icgvI#$?Gk?< z`Cm5xkpx}YJNuO&k^3i;oPshR4WHEh;qBt`D1g#6mw5qkd42{}&n(c7*ih9o(_`YA zVP64T)(u`=+PuD?U+oh8YfA_W3pxe@ahLq_vvB~(A{GY1&IR1sB?x1znwoGdonFZ! zE`EN2lX(lP3;%lO!}H|xJs;rFa(~vNZ0z7+*vH8hD@lscghsH_WdjX+4iS; zA*2cNMbgvo2ReX)X$6rYDk<-#t+;$BE`w<5JHjZn@HB~#EyR8qbSBv_#(3_)rpMY3 ziIJ@j38uaTIaA+-f#SfTR0!_)xIiTG++No*Nci3{PcK<(lFR||anCusI^~kI@}r>1 zcor7H{?mYxpXBIrwC^t}N{`^4I0*6p$!cGHAcV z&>&U>ZzMIahon&p1O%Z%N}c*5uE6&4H=)!CnHUJjy8!DyRH_o7 z?`I{`u?P=VDmS}o1?=E4>ur1{zGN_!-C&8Rk0>-A*K|xEvQ9tI8A>A9;{f|Io{Bn4 z%YZNfM6q+UYeN_e1dH2!`%uIY^q-jxa5&JBd%)XpMB1z4)FQ+ffJVEM^65?Z1IM^`6%!@hXhTs6$o zQ&VSq{UwS>U{wkDkG=OAVV?8nZ}ttxE&A&lE|EG--xF0$2Kla076W#*4Dem6BFph; zQLRxi-MO7n{P9;4AXy9qIApzXA_C%YdnN{364@PberCYJrxi(N3si^13wH{j)2)h2Emn7M6O zsE%MjA&A!v?hDHzI=)!MleNlthpDh9KMrC<1$jJx72$1dze(hrblH1@g zn>#31ENDQ3m68Y{`Qc-`W5P|(cwc`QT8Lm>q1-a8aanr|KdQ8W01BxOVOyI<{bix$ z0S|=VK?yLqfrlt7klgBcqkjFi3eB@;64dwiHrs+V*B>pr}~=qQ;Tt% zhaO36;^Vumsikp?zUG#`C}Wv=xKxVzi|#eVJ}v6nN=G}^aHhxQ+60+{cmV=iT>d;j zq{Cp72`%TC97w5idp$`@WLbb&8cM=3w8Wb-U{t(ryD`H!l*O|9ZoMXPq-bk0Zr_iD z?jQJ=ZXKb>fgN_tJ6B_17(RAJDVvd2L*qKqKu#Z<*xg5QG?YWiVx;+aHt+^5ZZzx# z1MKm87VU??m;~w&?)G64$h(2PAEfOQur3gX$pa+p_;8;>u|$T0wR+#abD>#NxRH=j zZSsC%5z-~jlUP;ca8h@jj+AYr)B1p?aQqK^tz09);uf_911KAfEo}DkGxgm`mQttq ziOi#4Rnx#?Dq zV@M4LS}hi28;xI&d2v}>Vvrdv-Lf~VXf`3TQ!n{z*f>ujyia3H73L0$P$^p&k2OIj zCd$rNsoM1`NmZ!t!A^RV8MnA3f-%*{uu!cwU09|Z8eKvH=x37#D_MxkWC7)QIjA7V zEnajTkip3LQZAH^Ob2!d&oG7spk)EY>UJPQz`$Mw23pz8WPR891epk5)DH23kw_an zwtJ@8g2>rCiNE{n*&r~i_t8k3?s=L7j7p3Xk)midpf@t3hne<-vAZf zB3#hVG@;jzWn*FGP%<@MY%7E-Q3y@0$nw}mi~d|grsfuBQU*+=Y7Z-|%Z1-SHCvnd z{=~A?!!)&Ww_V$Y}@fRRB`G=Sqx7 zi#9J2p8Z!aq5_*7TUzL2Z^p#~Wy7zIuG;n?ZoI)3z#st)$Jvl-$UCF_{m%{T4?;Uu zj2^~PXdh0^Mz{yU&~INHsLNlG$wf_Elc~ijV<1WOqh4etE*%Y`C;^?o3d|GV;in75 zHq^Sr43w!%GuG9Zb5$<&sj(3Al+*Dw8i^bjYKc2HBv)~=z%Ts%fe~Wwwph(22w`gN zPQwu7r_QH)!QHUPhEvNlzi~~xt2(U$4mWmBs(|rAZ!2pje%3ZjKZpjbJt%lhV`;hy zS=Ls0yiq;dT&7rrV{u-OBis~z`hDddBXR)-ir)t(3`Ht?*JRi>9Z6 z!&gcn6obP62}GKw2u+^^Mt4d7m2S5O;=!o&s3``=0QO!G!j6_#SlYn(93S$2s?jz= zu0b*!mLh|Jqzu>t{J-o07yl`+q@o4K8EAdCa4^Q8I4@;U_HL9hTkos1VvF#d240A;~u%_9++FxccdKpeH1mle>A z3lLI}2Q3(&N3FGhPs|gi9!~oQtP`)XnF`50&J0q$yPP^j&=XLtM&`8 zAvD<>6lnbP7_9%IeY+V177n=myPY{AXcox@FhVN&UpsP;lhP*4tIxm+MMgA}Di1L) zje!8D5l{%k?^zN@(}+fZ+?GLyX*Rzf;IyGU8UuO7MZlfgp}BLB&@Q(39k+IUT>#9Z z$U9eovAGOjDzMN!mkbco%_ru2~41 zKf3brcZUj)g|3WH5Hsg6je_(ezX;69CQ<$|D$djEfg( z0w`pj70dhv!7Kon-z9KTga-37KwtW&&M`DT3>= zZAEhRKLDnv{7{u7EjGXaQdmFv9t)wj{}hnvJ332`Y1moVLWLV2Ly-l~Dsi z6snxT^9)zP0!75Bj8ZE}1=$`8w+s^yz?p3F@-unAr7l&wwUp`6tU5+?mr`HbiK|;K zH1!>omSB(iURyX{?b!Tuyt8@gg+pT;CF$}}%Mp{L{P%vx{xy0SCdZ(xOhD9|8Y)Y2J_bC)8K zayw^0q#;X-G=K%lCDoD%16sZ`BZLOD`J>vv7-qKPGp?+D`Q$;C$Yg-y$_2NDgh(94 z+Vl)b6J(KJbfd+behLLqLX|`ajpvDP*t6F^>5r}DpC$P$8c79ARG9eg4xhz)Jyu#D4vP^DGhBC|q!cORk>`-8F#fy%h zc_1`h(7OvbN4)40AZf&H_15DMTF;<);BF-6GcsG?zcWz%MabtA@gjbh#c?4A$|lFW z&^CZlFMM=0RM7}zO^O%p8M={U5}T-5pbIrc&(N51pCot<-aN!bEWHj7rb{Kn#@Ug4 z4W}ICaWDG#-caIcMO-D^HwZhs8#>1;REawkFX$139Ve_|9*EA0*kCUqS5ML6+Zljl zE{h9*&tHbqLOP363>VRaP=x;C1@zJWSAom`8`~*$kGtt)IhfZqPoeUI%!B7hV^c&eRP|A ziYG}FL8pr(bmy7fYJ*M740-r5uPbn`r7juFP_|xBZ0Or>V3m{mv@cT zYVg?^&J2Lm$`D^GyV0WbWc1Fhl0`5F2hddEH+Wj1ltelT$QmW-B~7i0K$hY2fp}VN z2^El!taMVJB@z^@P!jXqek8=>7Rd^%)EkCgSk?G`>X*31M@a;kHPE=~+{_@9Ofj5U zAD2=`zJ9b**rb{y*VKmmm9A!kRFn6qVqHjxCxHV2Xa8-IR-mQ=W@#u6m7u9$cAzqR zE}b-HiG*~P%EI0JgjaNLsEMJSiUhZv43165By#eSjM1#|Z-j|8sT2@Bt$~UrSE@q% zAVBr7R}T5yTN2Hm6b9OMw*wQa2*iAlQFhm#|D`gefj%i8+xkfZ1$bC+uf|kR3=aqn zKVWM8MS-hwBGuvn)h*Lw$Y$)WAr>le?vj6PCyz z#|263zh^V*fEFZ*-qe652s;NRIA3jX47w`-pm*zRx4*c&9nhuYWwj@T=+U*x*&eQs z2wOFh@xvgq5~a4^3i)U!lC4b&MxI+)>y8D%)>3FXHI$YX?L!FCAIhcqqb!g!g?nTM zSx2Cp3B-RD1>qPVf23$d<~^k5_W}I=s--D94m9S&mCDNu5f^@?A&!Emb7y>IXK|Zq z4??EwZ*zka7FN`2GG!nMcbL2omrA^5n9z~sakB1TTNBF==eUU(K>(x!Y7sS#zNGnq zqmVA@I~9K(0!9eT8F1+Y`@boo6wS7e_$p}(q;=wjo+J_QNnEYO`D)A+Qg>BcM;(XK z8i?TvQO7V}5$p<7!`s!x`>7@4#WdHK7xDvWeiu;29-<)JfpiD>E+YeI`Qc?p5G~x% z1>w%GvkP6pRskaUskL~~37CZ^i0@xuww8R&4VCehzj=bhH8``b8;Toq;1_0k++;7j z_~m&1yKkZZ4C|w(Zat980*qJmhc+{yU--dg676!)IY7c>*z_ou1&o2np6&6|f^%Jn z08$L<7KY^2SSj_DB%ru<^bGmz>_P@D^dJTR8*N>k6_2;E5M_Xmj740mwvvYac>!Xj z=gX#$26zBaDVJ@V{{mJHIU+Rc-tYXc>>J0r(772wlr-i$-fGYKHJM^+4_?-2F%3w< zi@u7g^i)WRfXpHkl(8i5rumgBfKH)GNxNt|S_FvL-q8mx(4^paA2{~jky0EKmFyzd zCw>Srut)xxV01BDqCAdU^wvM!ms6OvK7mT3iQ6nDt^kpY5NP`q>`#JG;PCQEnl&^4 zlBQD<=0wXz!9g~PdcOTKw5n%-Tf?vKZ52f~f%+k=#b(JnuKxXKh#xkEj2?Fis+#z& z6f1(}hvUgDy}^*#HLw>3jKruA(gf*BdEmklPcxqKgCPHY9zfpCw7gCeV(jI0W^FBwD-YmRuW0WB?x}fQd0B{?en0N;Vqvzbwh^T(H~2Q#Y2<}E01YgA$Hy; zFl*WRj~fgjR!M09aD=O> z5o9;gSpnkt3>Sgte*(B8$~U(wVDYbT1?o}@dlgVybtYe=`E=TV2J2}xvbOCySd z=fxD>^#CaKjDdaz@h2T1+jB^l{BcMNVdo2if#Df049aTpC+#+{N3%34601F%ZxO2v z1q%>+f#{L~0?>2UztB)T@*NC|cOFGRh>TAHh)BwJ{Q-&Y&pz=3*cEMP9Yu|>`A)gw z`}r?nlMo&>Y}z4;fPWw+fQ;wgeg$3Xpk_f+_5Z^?|2_zM2L!lj?B?!Jg!u1&{M|&i zp}qrI%Rb;ml7Di$fxQOck>5*w|7X~Ljh75bgOvZ_Uz<<=`!W*Qfu!!-J8AyEVCHvd zalaa;r0s_I^6wG(`|ZFw|9?V)n7ODFsfEGh?hace=Mjg01Q~B6t}Iu|wl=PN!tZvT-@4KhFfD`jy(oR$6}m^BxZ_p5{;9)GQf+;sSx(aqtw zZdtGFcUzMJO3U*92)(R&Bi#dWc7De(MpL|bBdk_7yT}52LG7@Z3Of`c|B2PV{>lrL zTKB#LrMP;a0jyHc+CQ1r9y-j{Eq`YJy30>+@b>>drXo@2YgzmbHNPkju)J!)!Ufw3 z;FN8R&nv%Ov0&%#mpmhM8avTZ>y?!q@cqYdnZ@&*P~yegjdn&ITYADc?PASzE@9Xw z#56tY(JRh+b(}ADP55=NaI3t7$*gaAP0OrHuazgd)!lgIVi+ zl(^;T|G@{XV}#R2p~5?~v7&BZ){&Z|f3(6xT!w^w1{SO1l7Wb+0JGj3Zu>_oJU`OK zRn&#YXOwmoe)jy)}OE35S8zH^zaHFi3)pSd6c&HK5Yc<=5j93pNDEM9>ccEI{7rJbubZT!LopwM~&5v_zU-~-= z6Ab`ygdFBB{S_nGN3)#JQ+2rho4~CN(G!Rh+OAdQQaZ?J_N|u1TP-bx?BVjK1U@uY zC9($$!WxXL&;=NgmrU_J^k%k0g9@NReN*Y z45rg+4u9yu?sXttXEyr!e8)@Mn{;1bQm8mSVBzVOnAddTZ!nwf`^@^CN8dV}T=8B_ z>Q{3(O}DZfWG(PQ=)&K~^f8X2u0pyye90Em<{>~n{)1DG+?&_tBEy!m@EK=*(O1<{Hz#BgxU)U5|B2Eu*9H$tND>F42ugY4O zJZlJshso#4kb6s7mMT`I$v6I1h=RF`x`bX0+${%C8Fc=q#QB}4%RucE>Zsg*s<U26P@8w=MVPL9X_>q2bLUg9$cN)gI_B3Gv>>FuMBb|VCIPMY>_2sFzazsBZ4^oO zm=fYv{&tmQ`|?QXguGuhv$XNY4}<=W>N-d_)uI<&Q%4ud-obEf3*wsywPTKxl$VBh zua|B|l#n{p(m3?e!6#1BCtf5K3WxJWi*HD*yDXM{_)w~t*U`gcf63IpJ~_v(s3Nmg z`R@=JtO==g{rbOB>+=d|y9~HE+V-~=q&tQ!LMcD-1%>3S6Yb_KOUw_^IJdnXe4cz# zma%C4Rmm1caU()Ldv&Yz7^-t=>%CJ3YF1U`{3{lR8)|~PbMRK@B(^NW>?b`X)kvA# zN?Wci-zvW0=VF#Af2~(XzMq32Ywr0Q$jn3PV5jour;U;b5uMuJSYNc~H zDcZp0f?vtFb;~P!M7e2(kkZeG=NZ_t+@x?LW>Ps92mCOru^m|bJE|GR={ox8eQ{J2 zDO*vEYKVu+d>9-X-f)P%*i^0X`~3f5s#OAMe#fQ1@;iWP>o=J$6x!|c^nA1^+12S2 z-bZRoSeoaTt={FNh8)C?=5J=NgW(hy2O$xa`^Gk!Yz;8JJs43%!TgpLx0bjka3O7lpivhMRxoRct2W zJhkCM`-+YADBWvug7AP`Zs)~^*G^Nq3d}p630;|Z-z$PChh@WJI}5GmB7qBCT@;@u@1Zn4Rs8E zgF4S+>nkO3v64yw8=;w$A7+)rQIo{#LX)+K@r>IoYGyaCh9XTB-F>W*R(|8x{OosZ zx6DN^ZNIZggtHsB$ujH4!_4TVa7jYil>(bCn=8+P3B}7-0=z$&OSWm;(9=76BgE?R z5;oXvrZlBQz^ds<2vW&I+bBBa?P{<%vi4AXx6rCv}zh=|XQwnkfrqN8c%rzsu&>di6CiW&dZCc|g8^$@2k?#|8-QTm+>( zo)+F#VLdzV8H7I>QCmM5NJhQZ=-s0eE1kh_`=Y@uR6F~ywLSy(7E4-Mw> z5`6a6F5VyWV0OW?jy9HY9O<~xF-P9eMvHw#WbfgLcuBqcm#$WL1Y~xi68_421{yth z5BceNa*v`+NBM^u#q~6N-du6~_4kU~yKe$wbv(S+& zjf*ctC=VWOpZ^PYWaPURP1%0=8)b_Qr4Uw1KONDb>?GWj^=ny@>ogeJ_bdl-b$aWX zh3O2j$))8PHn(j84h%r9oJH$l+4pOw-gYhhP@-pA-<5&ds*v5qFI;jkMHZvfi`M*NMy5OOu-_HU7 zUj7wc!W{V%!@-!2_rppR2t#tp^;s`gaxN-Xv3jndWznv4{_whodzOYGbyVMp?Zos^ zBkcVA`L`~RslM61I{H}4@0lG|F@b?~%xda4{-W&@pZiPl%CW(9J=2*V7auTZi>G$_ zBUUXRU=e<+O;+1Kv2tNk2-w7w5X(k-z2}^f@SXstX@sZ_!lv2SVtQ-Wgydgw+cQom zvZMdY$jv1<4`J^e&i3E_f1_Hn zwO4D^sM#v2C{1j&XKPk$Rh!x*_NcwLphi{Is#POeY7<3`Qla(?5-Sn;<#XM~eSfd( z`u2C+|8X4t@_N0`@q9kc^JSSf0sZo&-56U^RUmn?d7!n0Pa}%td2p@%CGgmb7Q+JJ z3nz46X_p0GIW2bO9hsI76)JwL{IH}ft=+?U*zpEmqWv=hD)+ADj;@?8&J(f@e?XvA zQy^xWw!+K_u5G*$@D7!hwpEA_dVm&V4DlQah2Q#10VOm2uUL2fe~EQ6-d`BTJ7<&G zdIrR%d$$_QKj_1nqE3^R`m{r{u0isu>(b9!7i^r>#GGeaM^nXKJG@G)+&j<$3Vq9* zvD)IR)OhK;KS=0?gnKWeW<=F1(*t$6 zZtRl{*!(_cNqO#g)s#bgOGIp#&_J#tyKn9eNUAm8&58lCd~4^5E3m(6W4-09z3>EL zu56bl@<6_$>Gsv-NOW!~S9Y@Su+1(t&M@O`C`}S`sd(l+`x+te2*-7`H%ex2lwQJ9 z5~a}Clfar)-_LuB9nWK$*IW8KsiVK$XtA$Nr*Kf5k4nCBae*je?*z5O-k9ntQ_WT3 z9-EwK7Fo>C^a*x=QmtCHLHkf#c1Z1+ziWuh0r0>OYyau<%o7{8W+s@}aL{`eVT01& zdVUUD9?X~u64hTe2B?r%<(wDOx3{16f-S~$O<%e@!z%oy$W@2{;J8ma#s5-8N^7Eb?^6*JRu*yP5f^; zg;2%B+{?E?FB-~Rydcj(aM28m-+IologSK*;ZYk?k5S4y^0!zSkJS>F>}rzX9IkY= zLsTa16nUcG(CNaAq8LmPdN8~12K;iW5ti9~P+bZ9d`?hjUE|r5!>jF!+4wQKtU0IO z-i@R1&Wz!M0e+&E9WcU&o=U?Z;L!=tkyJ_)V8E4W_7wk~TJo#_!_}&EfR-u?Ou%t| zuCkF77K2!Zv$B8%{pcs;V-FqY!}Bf0A|L_sCm}0l0`i#8OvNtR8wZcR?Q~jARNY0% zl~#nZr2nFPB%5#5dR)fg$kp;~GU4aMgBRcC8iHg(k6y`mhsbb$c4iNDelgGn`}l#5 zl2vFxsW!D5IW@6wxcw`qn*>`EJ(V+uc(%xes}d(1^W@P=JCzh;T zS-j*wI{{wm!Ca$*xjscPbOZosyVkI}?1t{U{At}hw8mI%Suz)RpRK#+mTRGUk zBi8oK{Y0TuhvUWiHc55fh&Dvlx?^r|i(vQdqUaUE(QZK>soSBe!Qdi!k017Y-9hh&CDJSbdJG7A4JbDsG0^! zFz1WKV4Ma=($nugl`zmOpXG~65S$kHd_&G9hDOK&3N{`~1*K0zyLbadOM&=$YOQq4 zScMW#b5C&hV5hINd z<%ICu(`%Kv>(Ioayzja8hHboSid(8 zZzn*7ZZX?rA~r{P-m9R*=+njxXt5J`GIvIt0$k1!ld}EvVavUv2c?igqVRY=G#m4N z#;kf>SHCg7to}bW|k8j?9UO~cRa-5H%}T~ywd&{;~l{N=I!(N;~<$} zRV9e`qg%3C4{6a2G6?D>Jr)vTni{yel9L%RrK++LjQPP`v*&Id(q_3RTb~Ms?2TUyi#Sxl$>AUHea8h|Nj}Qn z7;T`c`0(^|?-rvLss1F49NkJtIggM`Dgldin)5(d6SF(cocP*~2XA8R>SpY>3a&b* zEtgy0a%8(IeJoy11AIRv*>t@C)nzV%UibT01x-#Ynt+hyfpGEQJ9d&@YCc!#xuCnL zW#qlKbcB^KED5SrMcFsFQ?1<$EV*XBHNg`GuiYT6Z~JR{=qx4_{JR%>Cmw{zK;Qh| z^5GpG+Uw`3wVs$6BG1h;|G{g1C%zWf(8``X!ry{XF*E(#v(Wx#l;`zs%ZXEx7hbEC+eolv1 ztB~kg6yQ(x4zaym6X%7^YwwPT(!jF4%LHJMP$8SiOqlW<5F$c=JX%V`4k2F&#X&1Caw~!|m(wRn$qu8RZM0niuE&(GPN?=5_C##vNBU4o zD0Htj)8DF*jCaT6T)NfaEP`dnTV{X87P<9LqQgi~UwQM@ZWN_aI& z+CeRrHN`!>SI@|fl8 zGp^Lc!Qa0>i-LVl&(X6L!`(UJ1)WprQ{;+}`=9F{b5M+$^6qirg^J*8S$}~1cXfVVR}r^MvF^sf zj7rGT<)}8a!@}MQqMLuOL!B+I2iTXwU1-NH;bJ@-?c?%xw^Ck*n37npy~=Ouj1eXx z&6~8AXfKMinUlZ(~22k9rVg31MqO$enJ=vDE5r_u82;L2-#ugqPxNpKl`K^3- zNuZPl=EWgL6-q+$9@9gABf-2S3}tWiy{Awzp77okV<&}U+nL-^cZD0{VnlhB_`n53 zk`E9rVr({O*OyK~s&aVcfauH?w_&ed{(5!6@%2gKCQ1yAao#r$jf^h%n~bjjAv8g_ z_ReZ2!bh0?zdyozVYGZ50XzAm2pU7jM}I!Oc6_3gsl8_D$Z-*SsJg?e8m#IL=vv{e z&OYghpNAXyk<5tez##w1_fsj&y^B4Z|YKg zO7j?T0q2XTE>3M)LA87Q)21r$@S^P;zsP(|+4F3>|Lhg-Y8pBbUfpVR5pb%=t)4x8 z27zuY#XOs}R4;=D11UVM1}3TZ7#dQ@oxt>MsnK$qLlQ)L!KyyFJ`USqvO`GICotET z7Y$_!nLM2+y>~0?@&N6m=N8JbEsCm!jjCqz==x6$UL}`Z z=+c)qo$;BS0KEn$dM8YMihTF=-bavxa~Sz^eVsfqGoZ^k*4k0pRvnqT-nX8=s#%pd4&ow`(v8l~X z;Xj5PLUd!MJZ=sske;dot2}a6zH$xA;GgXC^x+xHCCWqt8}SykF9N_G{5M1&eS0MX zE~WdEew{R9D|Zx<%R>k#+$hd!J;chBN}tNI>o`6Iq8pD$alI*1Z4b|e_NolhBF|*_ zan*6eG~)?91;H?#MQcvQIFAQodPF^H->)C#|B&y-?tbR@$`tFs+1oSrmm27_AWRZ{ zmeclE-AjPPZXk;T@{TvWc0feRVSm=PQeroBU~)4(IPwu6pDojhTqxC%NzArfZEiK! zubU1^9!oaWk8El^>preY6b4ptR(?vY5Fi&FTUUaHssWQ6!*ry&1F-;MP5szUdE`)>F+ zb#JTZ*=5_x;96#4bI8?pcw?M>Xr5mzpE=lfoE!h7dBbU!@xjyce!S=0_r240zH+w6 zO;AMz%@??@NU7e;Wn<%3o-z$kJDknyX5g6EZSSi~tcL!;B$iTM;Z5(w4}o900||;m z^uJ0RIO*!Yp1J&4dhMsAezK?a`qyr|c^7%_yQGKmFA@x)3dF+n+V4L;9Z(H1at>23 z53Mq#TAxc`?KvD5yZ5cSCjaB>!H@`A83fq}YDLJGoFHq?yC-t_sG5X*8J69*YPs^D z1SPI{-F0dbM)FtIVa0uM5Oa1zl~DQi0vq1g`SgGQ&0*s(56OAI5|fiT=hj_ zx&ESZJPg+%e7Vvq0jofQPpRVi-#%q1kNzX`(#I{glV9rhqdIugF-fMjawBlX>yI5e z>1y0G^cy3qkdJ$Rr)1m4csY5=XX@k)7spK3-Sq!9_#=5ZLW>gg`k5!Qfs$i`Ubb(R zTFhe-jv;B-wx`kDO#nqU*Or|-&`=0J6xr0va z9Jwg>AUDKI^mYY$=#?HZsq+w7u@h&N1N$^c)!>b0XmsW0AW1Oru{dMpxONopK+LKP zZx1{~Y{)m8A~);2J81%zr4nrSk>TEWYuNY>j}896!g-oAdZqlWiQhNOb+*JZY9F|S0Q+o0Ve&DOt92A z*ZlZJFL9~fPvXXNu`3x0jLmi61UbNTVon<|x$Afmg~`5y=Te^!UAkthyzGS0{2m<) z)530A`e@)gi~er#SP?u-$dE0tn;_Lan6c*9xAt!zs5!4d8HvKf9|R=Sx;S~mmJl6QJT3TD!cI8yAm)U?^ns0TvYR& z`&!@Y@$|a%yD3DfgFGj6eJJ4}z*+q}N|NDJ&ba7NK6lym$JG?H?~x22ls(R5;@o?1;%xw8sb;Ze;GKGOP*fhsd4 zI&BhjIq`lN%4Gn#XLNyWLRO})kg`zc-_--sre$h4;rt}h(%v=sH#7Fks%GZ>&O@-t zZECducB{FJ;=N4^FS_}f8?!~47iq4u=zk0KEoF>No-y}`CF5%C{lnmmrqQhg%YBS~ z(I7}?8H2=+sL1#zms!X+9=`mUPE4i8R9=cw*w~e(qq6I3yW2?twmzWMZ{Vi8JY*2h zDFjWlH-hE%o0agDeV_Zp%`kC02Q$@Tv8Qv|AK4ENfZA^M?>+2KT{;9HlMR`wy6_mj zjq1jpMSx%RJECj}#w1(P`6xEmO|DvlJ{z_0FXo>mm&0l2Z1987IWfMOxx>W+$j^pO@EQA*F5_iEo;?Tzu zX?Bw-R9xls2XRCFufVF!>ZE+$tVbB*iXWSNg6{&d?2)sB#%fDEZg7IkH6X@hfbMVH z#yy@}@!`;LpNUuBLR{%cQm&p(>I+&tIf|UPJ*J;5+DCope|7Xs##e{?75&(tqrE3P zaOhwTZRzTby~7jmlqhU74c0(?$)e%6nTXww0*vr+{|Jr-@NX9IP8qz1VN;A&&{8^< z%{NQy+E3RQo-0py4L#$^83lShC*}!xscSd9OR-`7GN?S%T|AeUDzr!sMEJ{R8eoZ? zj%dJh?gPy&$=M+=++}Gq^elMMr?hT4(y}y+JG)K(tAA#xSg9UA(d|`rYYI?7zX9wW z@KJwzN5f6U8>E_}zgc!!oEHve&x8M(MRN#Q6x2AEf+Uo|hX1P!mR8{&Hf>)J&!Po_ zg@%S1(zD48R6swqak*I!h@ZtEjj~A8-#4k9zhLW@QEiM8YSp`k9@^c3XXWs;2p5Z% zHBGBC1G_XRz~g$SEmcRxwFWWoK2~4hcpP^!8uU_o-EZ^%OdAr*T)$dtG9+o6_9hGr zE&CRHDJVjIZBbq=dIJ%Px2aIz zjvG~k@fd+n3H0{_`ZJv70qn6dXyL1vkL5KK&Wi7hB0IiLAy$DaC2&j zCClv7;5)MUji{CIo+6npInKbdxF3NIatx-2Dooc*bJ)WdMT}6pXS?Dr`1Ogdvu$}* z6JLL!5-;Y=(KLT@{sq7jO)+R)c+T0-)mdBtLHY~TUkkS~VY*>2V=MpZ=t2DW?$8#* zPm%%$=06kK*6SJ4a3~(&XB+2}a2>!Lz|pGNPu8!}<_gup95`FHxV(;9(_!%Xh5oAt zGi8O9Fk=Uq4{iZZ9Wb&80_5|tB#dEp@~GNgk$^w+rK9N&NqM+OS}9GpIn~FdDMKmz zF>GfQXO#57>_xn?Zz=3c{kRD>8rk669TzwrFBajF8+%6w7q7F)a0gfeG%vEV+kQKq zKkkp41^*Yc^8ft$|9cE;+My}{!h=U6UK%g2x@)Fqu0LGwvTHdFsstz<`MP*DjR%H) zuJGQD){L1wF->r3GExhbzVEaH2MYG~rljkGf~;JHtaooxTzbWa#M`Y&+Cs!6PJGrQ zWW55;ZRkJN3^T`0NzZ%d`AKx>ZHYLQOo>S>1TGqG?XXB4*#?rX$tZFU1~-1Eu;h*` z@^a$p{8Tp~FjuBEx=D}A8=5QcUk)0U2jX-;aDPRnZhhoP85J;seOi@-ut}E7h7*ya zH_u-5j)BHv=jAAA&17%Coq5K1JQ|!GWF1m>Deu;F7mb%Zfugu3v0;m-{}PeTwFLe_Xe0@r!~Z6F z*=Y$)omTvmZtjWw#12G=6?-_>i;wYow_6Fm#}T{GgI3cKyQ8;fGv$m8zhouBXC8=#*2)!)^w0KnsO}V@+{`-`!so9Hsm??D%b#%PZNz%!#*;`@*uAGOaEIU&m{`T=n za>Aho2%w4@%AQ^X@rPUj9shLSsU6baeYPcXvxqzDX!XQ%r`CM&;5~^sh60Cq#@4Uf zYRmNi|DGjvc<8ku&1CJ&@mI!jvgl`j%qExf_UUe0wv6EZNh+}@h@rYJu*ryHOjS?dAx?k@7ZB5iX8s|kTSQ6qL1_+gLJyocV^Y) zY}QOIDh}wjlO6_kGdHc%?L7-^Os@GnqnTbWdnf3?s#@y%Nd1?%<;=xBx^4R%>3%++ z6qc2(d4iegd*2~$fan?yT+%YmD@LpkBO96*I&o|2EI)j5t2yl1K6C9g4x2+3m#}2v zS1ZGM37$K0s@!&+_j7c6u7REfVotuixZ@V{tKoa@ckp%>@5WK`&(CpFkXSWu*JX{$R!j?wX+f4;r?c-sqI*fJYpe_8FPfH(C4zX&JF zOydIP%_OH(?45*udU&*V`uZFt1`@nxH5ufkJ;FY%RArj58gyEFAV^G}=I&2BszqQqx1?^zB*~yY-ICoW*pT&mY$ZrCgZe2qbDeLMc|e*JAvN!PNGM5g^6#- zDuY$?Mh#36V$HXL1nyre@%_2r@?C8v@a*DeEdPl>R@Wy$eX`tcJJL~u1QXuJFUe7z zG64ZwuYl-Reg?KbYl_w9UlBq73JH2%m}sI*33TO(F@9M)~u4M z^lu4HwVwD40lwPk$0G zU3GXlplPCQ#uT$@xr1@vn;Zk)7@;snnxTNC z?p_ZzU$(Wn7IMoqXdL+8V1G|*+&|pgt;)aYe=>rknYB8_4<_HdDZAGIE6#gxbX#(Y z|McQ_4?i-)iP}Y^s9qkCeC&-$P|NKwQ4wDdyT)3r_xNh!RcZH@mUp()S9tdw+m?_D zZ5*MK6S*O15kiT$TB~O<$FAT zSVQ%J@!?cj;eHeuS2N$~&<_YI4^znSmF{CZ_{3xC$x$H@*k)oV;jSRH(_zj<9x)@IijK#GmA6+fIb&eIxzjnZkheBC> zU&Qn#kewWE@m{~P_4JP;i5!2}sK6Y*Chk+x`=f??DNv zsaDa)&7lMhYR)NMhoL4*RgzoNd~Z^pWa8DBGph9}FUvHWc;{QqYwidQiw`Bj4tXgS zYBVa37v0T8zt(l@zW{u47HfdX3sh7~4S*U9fEE4hBl^(+$xfjy7N*8{ZsC9cwTvKsB5|DbN zX9kYKHWg^>F;J&tbFMMI7xvl+zRfkpqIo&+_SOy!w(N=0r8z=i&zUr?53zd@I2I`f z(?v0q>ZKAV2{RcxTf^*kGk(JH*M69cOH*m@^W0$N7BkabzL>4`L?NEvRgf2I2L`s# zJqSq^rbvCN_mJfI4b%MPBY|!9RpME^bM$(>c`!j9&so7i0Bqy0@`!GKWv*ZD>brbt z=n%aa;F7%6H5?q%-S0sDUa3TO6f4y2A4qrKPsXbpo~UYhwVaJXRheKJT91IiYGLAc zW%xkw0Gh_kYd)JL(i5r8_dMo!i4wT<&oI)t3OT~%RwW= z#)or5CsMGgK;HvT(ESwzB+DduD-Jd4f7GZjn(kwnBhxF@Z8^27G41*d=nj~ zaU1m}k%qLhOYWxVHaglSvXmy|I;NEY)sUK#~Qbn%2{($ zDdvikk+X9f5!&Tr@weXlBAS4B)jk62EVl#Tf`z5L^j+hLcB%3t4w&-}U zbqPR-ZfqCqFhw=nJ{Z41;{F6!PI7x2At>l`nL4&Dj0=~4ido$947*-&MlpIfENq=5 zz-ax#VW}*k^nSD{OF3fP*K_M71C4hAvQ;ezd1z(* zfZ?zYn{wu!x$#mU%--qZ)`pQ0$kJ)^5yWh=RL%!TCR5)>`Bie^CCF!a^cM!3w0n`> z9Rws4Vyp>j*#e*zo<+&AfB`8>MqgCcLmZMFl8!q*ZyPDct-p19CFW zLrJweOjI-1yq8(AqWy_LMRH?mZiTSdOzldFJJtB#S|rboszuC?{~(R5Z(=vh1F$(l zJ*zrhXZ}97%+@4R`lTQg-FNHxsX~`rI%9ppv;Z%_AX@ESOa51DKLqk!zHEpO9Aop& z=&?M@>!Wf)+ycmV2)jjO&_l}Og0pMbjr=8!RLM_g+wraRPMCr^f6_krOTn&K%svn} zyC4?95&+wKgYAS~&UFAk349I#6kZEg6%C){I83!V*@4MghZEf=321xZ+ZqX|`-0s; z{s*JHT>md+H0|#}sgUR!d%OnoMmyKD&z2d9nV#!Nn5NDog8`G|WB#PHSUrzQSDi-w zU`Y{__**sgyor55s2INdQ}bT)q&)qL&;=dk67KP|Z5p63lfAx*!P|%Xv)+}DscBW> z1R9|}F%F_a6CS;8#UpGV<8oplnSx(1mMmZtHGZXanT7ON-c8^N9~#zkrK4_v$e`-F`!g96123?70;|At;r+_lulz!qT}m zOE{rcZ0O)L|8aWM%v7j_yx4Vy38>9}RYP~3B~QUoWulZ_@rc>%Z3o6G60RW+;{zlY-1 zTnBFk&J`!1j8Sp?V(ynlUy?m<8P}X2<2Ayc%F(sEHbIjvdQT!;{j*yXl6f<#4Zk8M zLv&G19U#)qhY07ffv`A)5$frJ+nl@U6{gyD01p9>hn|my)Dn@0 zGOxQlM}m*xU7NEOUpGqNH6fhfNB=&X72d1;uW`D6rYAkFVlxXmdoj?!e`k}zJ-EVu z3>iwm7@`3ZxIqmETdhjows-8JH{_i7s6qB5ymuTk@8Z!T*%Bw>g^fV|7_@`3YqqMK zqcAkIwQ}@@!#4DT+hBR;tsur4G3kXPfBQE){b3fRQW4s-|c z%Kw)TR&^sS4ScLdo&&*rn1EcVKG`(A0a6Cpou51(_7uNSKxFzErqo74qevwDi(Z(pujUi@K|-zYi>OjI`qTF(EF{RFG*pecAfm7H*! z%mR{qvqe_z?9`QclOHjGI05zlw)d(7Z-U==>=Iy1B2=_B)rXLfO;GyWrl9fGIpK|W z^YSz0)~zKhCw0q!JmmdDI+09EJvOxc)D-Qw{gThMgi+t}Q2$jN*LR>HNRAhEdEjCA zIMjbxY4j0Lg!7B8B|-PYTml|YloCB80ei1@snZgFB|Xx@R??6v-#wvVR63k872ZJm zPo2baQ<7M23Dk9!&_U^Qao@cH3G&@{BU0NvL zj9@EB1_F=%V1^q>7IWs?9&D?Ud0DNG>|^+4pu^cIk=+D{{(|K{#)=MO2!qnFyY&`( zGfEJuhYL^LN?{y#@k{?hVBC@pFZ*$;xuIFmXB+~N(SG?!U4y=oiKY zL#0J>eTGg4&S$W3P~F)PLG^+qtBd32`#l4bkvz;pidCwHj~7{7+ELY1y#`iOIpM#a z(DdE>dcDE8qkMcX(l56iVTjlosJF~@h2M3RW5PUZeD>vzw5y-jxEbm9m*_ngk( zAJ$|G(}-dm@{OL{zIwl5^gt!p5BVz8>G^5<6biNbm_f!>3A)RHw2?T5=bes-ObxH#L?% z+2eeESD7nV8vc0`l)V0cJgyDLC8ml~WWkWCt#OmfG@-w*0xQC--JQ_n~o-CL8*|9T0V4|pdI}K^4k)8?b}w*_uxP`0w#X{ zway=&kMf)}EnuB9S{3rG6(xzVlym^>~H!ia6Ks|{oMu`ll#OQ;Q9P6Ul$H^BjMt*A;j_yD_qWY4#5Rk;Pk#&XNc>9sBY-g6 zxFoB7!nH6-maTKQm&19qkqxKWw;4uM4%mfH(L@jhSjSBSkBNft5v)*_?sa;}AWro4srhcEDihhl z&de4Y@c%xgd*1)g3IwsBzs+`fnv!pgf2M>D6ioA~zc$Y>0AI$rf&5vvx4seNp}aGYsF&5ciPZ1Aj+4P!(d!tfZK<6r+_Whb zb8qx{N03bt={eegIA6N&V3@WN=G*Pf<{!-uC7PLM;~=|L>ffeS9j17dXQBs_vFX^m zbg^5eFDGuJn~DxE{p92^?-{I~?Q!*Mhw{-)zBT-1G7}BZrhTpl)KY?0u!y4l<&7@C zDQj~ld@age8m0;tGp&~Q&z)KIz3eJxD%?N zSGRbDIH2}l6CYUlrLD>!G33J>m|2tFZK>5^(aE@zpIDOkQwtqz?^qZ2+i(j~k%$Ft z%b0D|flSgq3AH@;& zMf9H6m|U$~vSAsX;I)9dqRgK;0>0#^Eb4A;hS8Crlef?1CHm#{>LHeY zVK3MRkRifYIN+SRpu5&h$|M&$YFF$u8gSq3`On*?K~=%U$5$oCBqL<#;HOZAk*WOI z9kQQy<_31nJz%!9wzJqMpHL;-_; z<9kiLr2k`g%5p00Alt0HdU24PgGPE6v}xwUTUU25?}s+CG4sJiL1Yd^uK3+J+Ve$A z*m-zg{|O&eTSIq2|JYe>Ed*HW%oUyGUM!vCkl_8xc9wW>+Jh;BcUFKMMz$d$r4mv%x_|t2p1` z;|huJ0*a6Y3Tk2W>g{c0)h;_8Iff;jA2}gDil1H&`~J>kn1xSTh!rCg`{h$7R zixuZsodi`${i=?Nc%;4qn8lv1aAkNNupEB=dDRjGcxy7GPXc*xsefz)6(Z$#ZJv2L zOFN#|yo$iYGNOCB(^hGw#b!zb$lt=_B1O@%WB?>>Poe-2)jr_UiZ@F6a*)9RM|Rxt zzBqR%!uvB$oZ|R$p^z*q#nfvfJQFSSm{ErEM-c^B*$^FJ!3elFKzZd(LBHI%K)f@F zGWkPmT`$}LGeWj43_xf3^|5Ge*o3;hWHCB#wc(Nk8ObX6Fo#dd8+@Rc4-3h56koz( zo{v$GP?8I-glrp8-*c`z^T&n55rW_Od-W_sFRV(_X*O@d&)Gy~&{^=1Tpuv$gT5US z&)JJ9R+en_WuLuKIcF<_qKMf^fu%@KD6RT!xdIT&~NTG80JjWMQdo>BuFhelRqeLt|-L+j;?YDINp?SEF53D z&qpJCP_M*wYRk|^8$m7}hzZ_K`xx4P`>v#zP(;Aq{jXgt1E&D(lV;Oa2w$p51n#jy^41v6Zb;EYJ~%VPtyDRn0MJd93E46Vnz!$ z0aGvPoZ{IVS1DL?e{>!YrdvfU4yB0+C<#*^w_KN)=}Z{>1?Jl(18&)rbVPq8yuzBC zb?bqf9^}1tANKGP8&45_bBsS!%nW|HjvL&SF*#LzcmwJd0!LhPdSrF^yARAZyMQm9 zeZ_?{zqJ-Fkca6+R(VwH2Oy8E{Qlh&l93^xDQEtof4?OH1U@^BVsTmRAOpwjk#8c? zDp>E^#F3`^_zB2BNrTmwDrV`W4=*lmrU7`79rdNJufa#PQxWKCtq4yhY|kaO0>!i>9aGX zvaYMaVAMh(g=V#b%1xzfiha2}wu9Dd?7`7Lm)L&?i-$+5TI-iWF4MC=oEPx9r$tz* z(?(@JN}06sre8@|Q_()gZ^o&v0lyXOZcoAafc(T{X1{n~9(1|VS3>dHr+R9~?@}@* z)WMTEOZ_btR>(I=NdNuY##=t_yiZJ`PMcwRli4`VmUL`LvbWpj4Q?ol{Y=Jvy=Bsf zVy|6+kJ7>duR_ReSDy?UUaNQN)^Cq#$(o3jp0(@Gf$JoXu-nSdvy|nl9X=P$sB0f6 zECUU0u}*lQDohWo^20+VXx`Dz%JZK`X@?H4`xL=ZnI)u{a3TT0XhZa6F+4VzGi@eB zcWu@m?gEgMk;W*75rn`WTaab&|o$_Bum8exyR#3y-8Wam2gzW z$+(^=Z9S2sFmYd4v>zoSw@jEFEu#W#jx5yLnH2L~UiI{!vWxu-zIzBM*N;2~RS{7v zo$YyaoV_FnI8y)C(w^V@S0Us3*Ft8HLPfmd&e0%Q6p{4NkTCB^o#B*@yZaj-JLsL| zMQVX-bYomcTluYIIxiA zTlVUSqlD949p79Oa3T<@i;dP`uCTOdyW%8xaAsQ!0gB2WQ5TBuB8GgD7MdP-6@xPx zm#YiI>TlU z9WoKfM0#FWc|sZ5COK~Gxy)X+I3r6hsLs298$Fu=^74O_F28aOy01t*x*$3oyeS3c zGVa+|lly;n0oavriQ>XSs=f#EIqT+yVn#Zbo7lSb0J8vJ@NXmwMkk9*nCXpA7%oE( z4Gx@X2H`?8a-x=-QS>Wi%?gtxD(FqvAk4Ll6%q#hoU!@7NAg{q|1_yq4o*!y*?} ze*8R3(k|QtpZzKJw@Z#s?YgPuZfd)|*_jKuzPsNKr^X(Njneg9N9M{nH;E4v)%uXl-V~N3Ee*gNLT{+I-WjQE!{=ef+~pwPZvb;1Ap(5T%W7#=$aJc@K4?EIQxD^ z5V{Gyu$;C`oe!mcm>tk&?4s%BBpxG{&Z>S7Wv0=?lDC1l$jMP8?jk5=^?be4&BpU# zIPNsMA1QqG7m3Q{z3Kf8gS4BzIag0Yi;n{Pj+RG5tf{yC{g+jkF{ET@izi{rFe?M?7t=U!DWdRBxCk0LE7|SI zs}OPJ15P9h9Y@*_r%_NN02^U@X_7?KsQO{v|Z%- zQ~NLZ0>WosUd8vD46OJ7P?ezHY&d5T{J;B(hdKnt7=a)dL4R5EhhQd2a(u7%vG2|w z55mIe!Zsd@`hrSVW%mg6qjz^!;oYyJ3t*O$Nrb(2!{V|Td|(+Aybz{Mwz9F05KxU%@dw?X!pjZQRoA!PcfPX9HLIgScxd zL~JRs_l+Rxo4YJ|J*~Mt?gY4WXh=dW0D#jD2k2#xkIlXl6s+T?gWVNmB2tLCD#B~C z&B=XD=qZ#b>ihGZYd7cvx6o~+qY->oGF&Qu)cy9XyhQuZ8^@tEFO4DaP2%BZ5vK@V zJ?E~nJK_lesyjTI6t*RG%lbzUzx_=%ClCD{->?v+GeoMAR4D z+Exi45UUDInyXrS>Qm-#HPE~X$t5vQJs|B97u(ug)niyaAUvhkXM-U9)WgA9D7Wf` zY!!p25JiBMpe0l!LXD%FlClS&fia|pjCZ#&pY`H-LD({5vSNvjKR*$M#O2DbcmpFuWT44{+%g?Z0*-QU5LN9J<}*E(OUoipCp7J!QerlaS-Y!0qf>E<_)#dlCy`3H-ysU+){mfQ010F1NRqf!@0qyV zO9{5>vF4{H#IJV!?bVt6`RHY}I@L!fnT?#2)A3eAH0!SJjfehUy;E4HThp}fk zL`D1H8To~KDcsApJv2$uwh%sc@OY~YJx)Wyg}E-zy;d#wwbNmUf@H<&3_PTz^3cQM zGGlHht%pt}^z1_X>`OENd9zI%(2D0hLXKY|0VdtZP|&fY7Rw=^O(6MDahidgLH2dX z`9k=r#N*E7;p-MjddQ>8rhh}td#?U5Qb7N#8U5dgk%f}=n?{CvvxKGMp;g_|{Ac{} zgWao+7pS_C!}F#Wnb%a!AQfwXnRm)Vp&D;}G8bW1eWC^uGVN^1+P`#^%l;2z?;g(t z-~W#za?CN5b4f%Za>!|pNe74IkYh-coMTSga<1f@oQ(=a%pp0OLqZM{F>;vm*@lhT zX4`LF_kCU8`}$nhAHRS0mxsspex07L)AQxaLeDBWzF_qE$=KVFlygN(IwOp%y4N0! z-JoYtPBJdpNJpgMLW^>#{oh|c?ITxQ>pm^A6|)mKbIWzXk4E5v>>rm+%ZxF&h4@!-tu zk2l)!y*7u{;zDpq$*|oQ=h`Y?JaCcBOl@dHTnA#>O;uDR+2$PS9Du? z#ULW{+^^JwSE5kBO!m%2Bi(jJdSS+URpaCt)JC8l>3-hnQY0(r+at@uIGysth%@(J zRiE*D>nxcLI{cv1f6|;DqJ(4hdxy9yP>7n*hx(tbV&EYk^0i)D9GN*Tn(5?Y&kBmV z_d%=O!?%D03x^L&ukdFuS5iO=H`*Ddn-R3A#10}Rw|h);6B&P97Oy8x(dppgb&tu? zD+b?fxIgXlh?H&rLfg%tkO{V^{nL(TF?Zp^qTeSeGI{=g30VDh{vi3xqV0tJlV#Pw z^qKIqiiwaCv}?>;qmTEMeqNjM zmtVdKsEntPrRW~m56y1X_9YdZ#iq6m0A$=Aij26ghw^kA!$rBZQi=tXj_c{Zkvd&d z_4cbz#<ml@fXoc0kq79MnXnoYwGh`)ljVPlALT}*+ze{+?{N_4ae}L zT(?x};7NH+Uz%z}*{5b!4$drSWnPjt$V0BYJU6XSXs89KHLKjF6}4IJRCZJ|wzQSH z_lZNDGlwr~(y=RLhlg^=CXHG+QJQVEuwvnO=5!x1l@~k87^xz50&h@T`Dmh<3pEkaD<|w zlWr>1SHw{KFT1tz?zHDkp~-2a+L_>%%G@dI9t3F-=YLz}oSXY&*5Y#Yzs1@3JUX>^jeSWCCFhhk zr_dUFt?5d0fdrQaI%=PD<%$hWA(90#7RC!7&F|~+_%`u*wx7VwTu~C}`$ikGkhO%p zI$gLu=GqT7oG<*Osh%RLuICrM)OYP6m^OJe)qMnr8?C=nB3b;b8F_X#m>C>N>(@t~ z@w1HqbDHbFYTM^BF=R*`nZwi0cqmJaHJXm>h4NgFeLjHHW?XK>r!TDv^`+614Siu% zgQtgIOL?5YUEKJ1qf>`Fm+_=T3#-pUY8&_Sv6G_on9UkrmES^6 z%x4{g-jH}XrOaDCI$V8wZ3p!NXiL%t;jdTwy^V0z?f)_p zh}>bgtPD@BL8Y*Hbj1 zHC`o15bD(ziI}RMMsMZk+_=ca7NT~grGh3PC*C>UJIH-L7m~*AcM|kVzULPIYfXOW zBaA&Y7l!R!uBTjfv}b*-Uj?AtO(P^cO)ME8(Tbu}o^NDBj5Yb%T>x2&DcE0RN&766 zyk+Q@w27q)W%Ctdtr2>_8Z}#o=p&X~(eAG@XM-Z1>)ByK57=g9S0}s0i|ddC zk@peW(+yPp;~J4nj6jP$%DlMLct&e$dgxeg5h7>1>cJ$#LGs;M6Uo9T@B4k~+@5~i zu{v)1=AC(j?+zOv_)?%0cd{<)3X7YeQ6&1tsTLnK8ZTr#-OE~(Sx-?8117LxR4q!- zxKzL4DE$ia{abJz;*s{!-`t$vA7s)K|L*2)JR9hjks2zkOS>RvW<8l=Xp%9qXW@{^ zdyJR0H&BVQ?}j8Pz2>{U`+N7x!CS7p;*0_k*gnpg+Xe#u%HolGxQ9(YC{1;-Q!etK zUv}GmoZE&2cxf%2q>BM@kH=`NSZ~kX_+5IfKz+J)DZe;%_|Z5pFl=_5@z$L1xfGZ^ zv*TLy^UbT-oiAVFiOrW?xN{@BvFDx!DdIj$LQ3gw!5>M{zWae+pROArWJ%@lU` zr<>RlE3fTWv8f3DHx2XEnu-iOaN#0AILW<)TEgkOd(3Uv8g9!Q=PxmJL?eMynuA7^ z5QtWGVsvffUo&t(`b95jj&%{Q<9pJQ_TFk3(@=?{ug9f z9~$!Wl^R*N`@_O98ryWe<~6<&8f3 zzoo-+d;Ao1Y`@{gBA99QSqgl|=FHM7+-WAMs;6`vw2cC&+{r5W1^$l+f--m}dpG*! z_;w7~$*VvjW|Jx}nu%~CUf&mD_{vh_x$ZGu70>0??twIo(SeR#y*^JXsDBE4uhl=@ zp4|i1`r?MN*M}m3%pyzJ@1%8*wks(pQlNGEiyiUV<)*_$Tzy|d&HQ5H46n3%lO-Wu zw*&nK`2|6LvIkbO9A>)`Nw{30`%H>_yM6QEtN|1|JqH7i82eCoJ?SLhe9CGn5!=Af zM{%#=YjZzx&5#f9|J&zg{ts(r{J&0(o@3-jZr7Z?c~lTq&Qo$}?3lY-;~BF9NwvaM zk%$u`0S_q&r8IJ6Y_SZlwODd|U1Jd?n~+9u|6vjL;B->K$9T8Zmkna_USnn-01v)c zZ8NKpc?hrT@|&G+q+BmC%PeZNWzVmR3g3i%&dL@S==$2Ip-sB25Kw&P#XI(IhnZA&qmklR`8BNmxFtlM$`dmtH1M_|uz$;+JfAh;fw@S}*K@&r=2eb`&3D>A-YY zW3VK1z?JBv}w*|DX8%WLSR+eIF9{ihCz`cLN`u=(lh~JL3y7(a*<}OLi~t zeAvMA%;zCs5BX`b=hcDf3kDu~n%)~eKwn8^$a}tPNpINBr}P4xWBS-tOh8QhBCxD% zIGE!cYuc-e#x8KtFLh<$r8{3MVcm)uSYsk{ z7*C|d2@_v4U(-q1zY417FF-jZfXd@}JXW|E#}rf}qfO?jSjl-7uSB$A9C$aP$ec#} zM9E85uU_o+;GjU2aJd822?*%;k_6kg=ejCR#!@-TnCnDY&sSd@C3M5a#;FGdtJeBZi8^IWu`Zj^ z9_yz19mEW4tI>BTi)x;IZH?!U`b2K|j$RXya7 z9>J}D?-86))VR0t5GXynr4!mJ5UX+e;qFSC&z2SKJGu$4me79VK*ALh<475(kpsnQ8Dj=k-R-?W^qgxq~frS|%+6dx%8YB}fv z%A8drTB;RHf^hoMBF#VVTwOi;N*oqH|HJO2EsY3dZ}?++MzUey4LYyYONciv?$x3i zgPqv~tK;KSS?%=a+V@?|cf)9``F$bds96jLs1oadNj|#*k)RKicfW!aam^OC>HmK5 z+Kyx&Pm5OCBdw?B`fex{u4}&Kz4yxuhptoH77Js%0#bj+cI-imq;LFlQA=uSu>;no zdFiRaYTSfWc(@^1EcXT!cvo;vc{8Y(*ql6tA&j%x+MBZFXk09c;d2acKyBg zQBoH9%Qy+_OFE`Ey!S6e`G{5Nf|U{6s0WnM!@Stjk(>jpIxsQFTZAlA;SbK2!==8j zL{|-Ck1<@m|LSKcTgO}CO%*W)$*eBF5~0qluYkqDQ&=Efgg7L}2|?ITnxa*$j<6Bu z?4hB@y@*E69{c#-ZR)LLZ*2sM^=~|g%$dI{wMmKp>YL{nRSk#k%^X71+*_&(6pOdz zR7^I}w{Q)d)r<$zGq3l16LH{*55RLg8cmm-V-f&mvA^@TYiC?t)Z}NiGOG$1FSjKvMaf7Ue9q;)FroLP&&F@6 z&WL6f5!E8-dh&(z&<#7@;s-!UU7q4Rv-;=f0&~3t5t7%1)5=?tHl-82g!O4V)1<8Z zzDMEmw?>_uDZ+S&4(qIC+7b)+t3o1O%=r!$;@jQHsGr7%0p6vIn)K|~pHw1fjL6V? z>toU6sqg|CZ~N?>9pf<|QJ;>1pQ*q^9k9;35j|915u*6!H5frJl&2}r2HJhQ!GQP? zPvg+pgQCPw7MVGF@(DRJm}k{Sm?`QXkqn@9jPl%{+SLOA2^T?S>WVmI-&e=gNyO*uo>mc+wRbk!eig+QM~l)Ruts_3>?uqpMO+E|211<&ZDb| z&c#`Ne`>I18J3m1Fg*dI1xg);_$_6htaMi6FD?U&4uZ_BBr%;kaW87!)k9lZR%b1P z+h|J^X6o#3XOS43hmvSxHH{gag9~u6YdZRl8qVu`OuEYY=)wtdpo!<17+W(SUos_k ze{yuCv>`ItP}fiUjZ6;Ybqy{rwLUu?q7>alKCb@JW_BB)~VmsulN>EKUEXH1ZIiAS+j8$Y{lR z%4bm{01znol9)+f&20eIMv7KTZ{Y@IHBSMSU`K_xzy*2ss$wlO?&` zcj;9&K@GY5d}+wqm9ey)_rUnveN&K)+VU84&2~d-TkacWERW>tdgCh7dU2%xu1BeUB>gEDCPnx(e>aZ{0VP<*sVga>SSS~*Ro z;ire)NoO2@yau6-ItO2ua56jh@SXu_!7R=1fhFxW0kxtsA#6z~!${WaBNSu?OMiAv z_H28fL_ucAmh>An!xAU6IiBqHa`yH88HV8a8PKx()xf;Xc^-HEK!1gZ+MN6n zpRZP5ifyEfCc86PSnxAdqh_dTGnk08ja0@5sry5#LOth&0WU$0SX-ai?-vNl3D4vV zb)b$BaY4yuvlvtHa76l|(!uiw496!RK+O_~8g_swmZ|NiU{1k_%YcE?lV*^gt=L8i z?ex^R(7RtZ7>O;V`-`c^L;@fLEyEk>K}m+Ctu*}OoID4f(b``AGSf1oa97Hf8zK0rc10Sr?h`Er51#^~YPA@~KSr=_hy#Y-DibM z5McXFIOy=||2DUx4gTOt@BG(v#GGN;)_BA(ut3GB@v3;O;=YzzAG5rC@`Rr0nc+Sy z$y#6OY{K(1J}r_gmu`6GHI#r&+9Z;KM``A}7sH|czH2Y928&d0v=Mf7yL;!AN~l6a z$^-Dg0P+N3g6Mm&krYGtphYVvv87eOnJZkrg@_@^e^Tfu4-qf((3H5)>&ZY4?m zkP1jmy&9c`90`F14C*@s&Rn}=c|+^Q&Az#;{5&VSb>{_0lE#Ie}2$}6Zqo)`0 zZ(9%fEFb!GUr?{L!(wXG3Fy=ci|^W8d8AOEKj0OzglAZder>#_Y{qK3)ROhiKQFtS zcr_`|<{?-ky%pk28{Hr@FAc#?ewHjkkn%47eE{O{4{-D9H@I1Wyh@UN+_Gm^#-F_T zw{u(vak>{YExwOj^VhaWzhZNbq|A$??YWYq$`f-}uVMpwjEyK8<#l#(#f{bj-}4A% zMRf}aS*%DkWoO}&iFeY9Q|@1I4ZbkB^XYOwDi3%0p#{2uZ{yg=J#z1zsCI^$orjn= zY>G0jlOQE-XP!}`kI$mc&K|8!RNY^vSF&_Vzj*3H1y;9RD%Cu0In;dx^*A;h;(0-? zie+E%EHp-&x*}8>yQK@}*5&C=`U)0-sNYaOqs@J$zt-IjC=^Ir@jGAT#?a!v5EMkN z=u+7T%0?0PMdpHF{i3IE`=QNx;#mI%Cy5)W%U70UqBL%L5F=8E`C-|H@}Lqb%X-dN zCD00%^^vJWXRc?qFlqMYrBJxm$d_E3bt?s`)bpgd;~`I8Ez@c3?M^N_71*HX3rU!n zO9)E_#fh$*)f@hh$CG~r zoAhUfQ(b&A2NlCwI6QnfXcglPivr8CH-)!-HjJ+-7S;})C3lZf;jNnUNsAf?EpT~r zbpW%Cf98d$J6qTH(k7IX5a*=(GR4A81t*XDt|c z_6)WExE_O*V?Ef4TTcrQ^3~KjdLnPV;(IGGzM{M6lqaV0V$j|&`;4t<$X0}b*DO)K zjzJ8Nu75oKd$j4rolpm>^@E{DGb;|8KfTgA94gBuuC0Aka=!OV**;LlMmMAT%YX;G z+>}1Y*6%^Mi2v$KG0slp-51S=Rl+6GYAQd@aOJWq-SdsRh-czF>_ogeTGlbYtKY}F8kIU9~yFGqEM+o)_6r)uFQGI^TByenCdttnAwF~_3nlUIem8C z6^ceXi>R%ILwl$*J0s`77D_kk<+DaTTBBPHJ7^&AyPMFl?jtxH|Z@ zz6^~~N^XNXsRygwqWHrXV5z)_v7sLXTeB-KXtxG>rwYZ~d$OI#oAABwc3ecFCc;vy z6r89B^IF|J#Jl&{i_~LVn>|9VN-9V#nO>(j_4ivej8Vvd_d|LhX(4n34>*~Qfkook z?~k2dEZl3za5)y7To@niAl9NHBeyOcJ?&#og{RkUW>{my`Z+W{0iwcd>iFnuS%i(; zH}O2d4&BLJA7+-yOIX0;6?9!!D;IJhY`-gppQntfen>W=K zat++6bImv0!TxMEc1^OViv7{1&!M(KS+I@%PlzqL`cChp722IIfuy{_I-IURIaYKf zJfN#klp!yQ#7sT#D}PWwc0lZ6=W24f7D*vS4dI3Et_V7{UCRxBaZKUzoRpwL-TH8U zuC(#Q)0J|c!wtYM7pHvA9GAlkpuqKai#GrHSx#K^5Yp$s&gaxi`*>~&cmZe6wy$Eb zlXd=57(?ID`!?go>^BU1c-ANXeIME*#T!4U}4;?Ml?aB84> zt_d@TH~m5J#Ik}>CCh3{y;x5AgB|$7X!{b6SUmwqw(63>uCWuVpoi zhy2=rqbTlxmHlu>evH4?QJ`2kYC0HRTkWOaalqcB1TXdil+mHQ{m5Mhi#5AksKs$= z)wUU$JmjU=EZn{TN|vl6qilbw|7@ZJO*>p&lb{YHK2Qv+p~`sp)OW2?zX&voJ;;5$ zIy~0(HSWLGV9`DkdyRjA6fPHk==8l8W%(-qh82+#VT@7^y(&@v%emp5qBBb2kaS%2 z&ljMJ?2^iVzw&_&u@oGYuo99vgP!_DjDlTIWzGdV1?2uh8n zGL}JJ^J2QOL_Lq{>-f2nNGCJ?Oa+XtiOs87>8Z6tbrM|1*#4g&zihS(uFoW}*i? z(+>Xa%SQBZq>+1D*?jA@7P*B|xQsrLwz<2xd_lW1bMoc)2q2T+ST3?S zzj2IRn7yg$t?yA1pCMQ5f$`8XF)~~LnS^iJ^&vR-xsmEbCrI&o<;sm(Vby7%c^`v{ z17@b#%M-_ryyxeFNtxg@4b=wiD(R3A{kZykKAmZS!#?RpgU)8t6Aq>%9BIsE+MFKV zmFU{}9FwjoIMHkMQodD1yT}mBaCBKAp{j8saFu8)DGG08j7IPz+aQ@W!yUjU;Do~h zos~t=HzB;n@Cwh^Bb>T+fgYl)x*hz2HvoRz;GET}tVD^>t>;Ysp+QdKk*C%+!$qjAVF%kM z(cw>bF$z{I;hpc`9(H#=VD3W#m-kO=O6vXLtABsUD9{vf_A%h#+%W9J_5Y)a=jL_- zafWQz8M~RHwkdd+D6(N60lXjQpS%=D3Ie-uQEY^j~T#sK8zS(FuaR66j?&HL0XSd(v%4^7r8qcg404h#dR%rt3 zh0jBp`5*e08?^s=u2aGeIiZGPuRngknB32X+n(c|jNEN{!{TSV-6u$)pmL>e4XlF^ zvo~(a9fI*wpU^LaeUwF{Z>)*viQ|^I%ozQiZ?j#4AMmsi2v3{m$0i*RmI|K~WUlzC zbWZGEQ>VnCGm?=jtc*_I+J|(8dYUV$VO?$qhA$_la(atcPi8ATK&pOHoEb~g8SiIP z51DMtl@>k}BCo7|tXREqm5R*mJkr|iW|iSX3O_E+>+WOYoFvzoi!>we@E!Q|F0RfX zK4~wcaTmPb*}qpnsSGjg*{|EUB+w=k-vZ&vFVYM0zjH<8tq@4KA-hs+!9D zRvxk>1MW0@dPxS`U-nR>(?P^ZGft1~{zCRkel?*Ns>RTb5_L1%@y_8z6D{}{^lu_Ej`mgNF3}aV zEc8jXVRVN092R2Uh~lR=g_^@gTqt{PWT*L?4pX-0tKYUwjVg~N!u;3C57k3I^1zcK zeB!s9((V%#00jUsvzOhmUl;NO&~h zB+!A}CWCz6K-WS?$@dL@TVM!?aRCC~*G`sxMIeb%$t&{5Y3S@?+#%7o5uWM#$KU*c zt<+l>7>WC{nlyROG7|ac)&JvvY5(GbNnz=Os@!_FXQ3;DMSdGMlq=KGbYo*fmD~qh z`Hm{uz1EED{iEAEJGJ|YY99sGEf-z99)0lL^ou8O>b9OI?XnO0EhOIc-+sLKc5iBN z;+~Y_LoH4nQ8rD(e4F;>5a&pWTkJZU+mu(fPNp{B3MnYz<-{E@}hQlXp&C%#5{O^SNU1M(s;=Jz}@1#WH8k3xNN^Ko%|=f#*a#S(809 zhCX$ZBi#ht@spQE8PvZ4PGH!`;na%z_tIsE-oau&DClQ2TCFDh-od!_{k9?`FgyYf z9}JwBe+&lL*Vb#!CqSmye%k5V z?+@?vkQ$$?Nah%ha4eQy$u>f+E6${|XgxYJi*dv^g&Zeq&`j+~XlLi!&Oqj)$`W7S zjnJW{jRz_5YE%!oJnIZFhZoL{)9ayh^ZYn<>UmK!0@eF^!xO8+x$C`^#4Q$&;LB7z zyb_D6+^J8lx1Il3$%1~c8&`?K%#@)J*_!7DZkM=AY3W6s_#HLp_IYK_yNA}?yX;bZ z*w&3<#X7@C-JXA;09l%-t$vn?8_?im(%<2u3hM*;*lIU>w&7a8d2c474gV9_0bt7d zTn&rDO)igS-kHd<@~GASz-w2d2w(S$OWx$4VaOFl_|*}1N$CKqc8{`$!G+N77fe1t z=VB$edaT^E$g@kb<^=>LrvCCkZKkyX(JHx*V|tv|?b#qbC(D{@NDT~kY*q!w=DUS3 z-J9?oGACDjnHO7x;}30WVROH6z|w z{9~=Y(uF2f^~0P|_iwgflV z57pl5r%ySR&8<*pI5unVse^sg^BR9}WX{^0f^9tYHA_aSw-Q2_7rPD9L&d4$x&He4 zH@PN$ys&=akI|im&A$XXM^BNO-CLoO;+rqzPh_iK!8eF@9^S(*F)1w0nWKE2z7@T@ zLr&j#?6>c+A?9b$&)^AhhV>`~1u91zZDurZbem&5JFnt z#Deo$=A;>MD=SW})Bn_;d6jv3`$V5sUdwC(|A^BjsrONxsEt&uT(|3N1zljuU{R@o zdHh?r->WXjR`2 zALi6A1(><{ll<$iVV*q7j&xkyL5pQebBD;=+vGr7$|Rf@_wb)NUCVFzJ*Tgc`^3K# zE>i8?^zCzXR(wU!?Taa6UiR4vT(Chgxo@!2_QM#5C-ZL(D(Qesl!4n)KSu)l!SINTSc_ycS4+_odYe;kfJmMLSE3<|HKon?e7QaA?l9NNK4QHbK89t z`#T-56<1xrNG|PW@!;7X!FNR&T5!+g&74e(Wk865Y;rBR|n*c7`p1m4E=Y;>1)}IL-%THMvh>}x_1Y7wqkjnx@;{OWMpPqF+ zvXN+hi_}~Bx)1Bxyy0Q}Y=e6&W5DMsdVjPOetqn_rhN}>;4Gq&pcRE6Czv(K?j z=|d)dS*)} z0s+v}^^toWoI;&53xgQzfTFgwo#*%=o!dLDThETU0i=D5G3_DiCJI`Ig?!P(3}gA_ z=|tN}v@XPBlf)k+8FILK(y-aP#r`K|x{YUznpiff;?fSo4A1g0h^!!{Z(*97pia=) zjDuXcMtzXB->>Ph`2qWT6Q|cS@8cQi#4%;Tz-v2Lb4DVCxh@%x?D!G*=mLHHCwv|8 zD3=WjLPTu+EG?ibR%H{WPqD$r_P>4^6 z#Y=Rv1=G_k;*`gi2QCbY51GY09H#pZJpR+a{&vBI)-(uZ*T{-151^GWG8WbU%+BWS zGrg1;J5v}vu*i02QDc#fQ2wOSGx$0B#>&;$YkbNms*gdN~rj&7w;h zcD~&)i7VH|k}eDb3tN?klY7ct!A?B@ilWso%_aPMK5@)_R3uz(1>|T1jJ)g!I|edi zl(Cx*=lwRMvPPCsr!_4l9Q9hEr{PJ?nw3FRlMSc2agt7^{b4vl zptlhe^#0ju4dl7s956xu2>=)_X&=#h_(_6hb>AbH8!GQQS!Qk0z+|q01r*V-Hi;=V z-~VaF^lpm(8#(C$w)*T^P&p{RBQQObBMrG?l{mAyaC8R9>h}zSCfP|+!g8i}YBC5$ zhb(zdX7?demiU;K%SMokdrzn2hLh!2#7ke$apDpLux4D-ltjdQ z_-cEEM=l z)uQKpmS{soT(l!x`-)!JKtXB6wHBweb|l7x=b+(;PurfDgN>m#i`eDma9W(5Z-F5l zFC6WMY_G1jO4G7?R$rqYdcY_rZq>)DFV0H4K4-v>P?|Gbhy*rZ+>W+xqc&p9Agy9B zct&#EHEaJ8Qw?qS*}vT3e` zPj6SAh;K=cLMb8l%J6+bfyvGTU1RV^U*8x-_c?FF1~SJTfyi9+@=uIC*dQAc|ixRJX!_kQd=z;iJ7xz zt9US240wE#oP^(R?55U1i&fB$`OaDp<@%|4EW}cxGSI_5M`c*vPUXD?&Tg|vO$YXZ z?jDVQNxEy&i^wptvO2qc0BUs#=yW`;bSHa%1VP#sKwN{S^4o{{K zApVE7n0oD(ob;6iPDn!rW+I;Y^93LK^nsQo4s_~Efq*9_HrNjtBMeb{uj~BIW^T)) z8xNWXt3~WX?Q_(NROY&__Xap->Y&Z1$MytGX!y-p|b-SLN$I zIVVPf=N}IeaWvLnLH9F57-ak?JYeQ2Px} zaj{+W{l=v4?BsmiNtz7IpN{UQTjqWz`qdQmeOxuRLii{Qo8w_{Nm?KdlH&qeu%5!%6!VCJQgbRyr4NB z?EFd8=IN`f;d9UAAn)c#cfvTItm$eU-{Zttt{PRk-%@Fa!@p$PAlt zC26-z$H-wypSHGfuC5oz1fcV``#77fL8XK1hQYKuVdt+l$irz13!yh5Luy^`a{F@* zQ{n?{lHwOjjTa8TTzS;ZCsN|;9m@=@jODp~J#n94%00vS`5d^L8{UmjDAN%9)RSaTqlPYjDX^n8z~ zIDs|~MmtGQldw5axQkD1ekQkK2j{Sx{lh{BiCV9q7PM!tfZp&-s#2iYS*K1xE8Oe^XXr~vUja17Y=nn@W2MWvKVP@dTGDC|RzDBi1YtHk=873WvCF0WQ4q;t=$&Pf z@wx-}`~e%z%#mxp0uGF&s2wBQ7Ezy?L{cIt*WAt*;9DUYd_#hfZ3U4H)6tdw8@1cI z&Y2s1Ckn1Cv{UlVr>H;212()Dbm8SDT%1SJmBp(8p)T4Hgl59<59H)BCZDZTA<;;> z^={rLUN~CEE)k3(H}9(J^$f{1p)Z)lq|{0aY^m^c!n*#d#Qir6(lId0=@vq&+M1p7 zbs_cF`v81vnK7_puZ%3(9s$#|&J2{N=4 zwP5E|itjQGpI;LewMy`qPZU>GpBUu2BgCdgRUt9BW~?UK^D;;@T3Nkx(9 z3606ePZ}{V7%}8gq&$A+^xLe@irC`HdyoUwDbWIHTx!wjLjaz?T+ZKD8;C8sw8X`qjo& zy-ZoY!aDb;9~9*1*WE1r7&le$W!k0+ybEGlXoAf;e5ACnfdB&=i`=G;eVw^*&$Xj_ z93!F|;ko%w;t>>(Rb-pS%_p)hsIkzq;&xnw9|U!DJ>)Xaki+RjXL0|izSRpq@hSp~ z2Uv_E68C4gM+00Zu6pWFD+k z1m53j|CKKAvM_8lNMP)clGUux@M}7sTSPznYP7?AO@%{|uL3L(L#m@8!SUVQdzN#K zdA{|Jj_`5lVrfQ)2!*%JP9i4={ ziJ=~?Obl83>Mk-~lAF`}Hk(=VOX}i`-_b8?Elg~EeAkcB!4QO_|L>5srm#tPx`Q6% z6FgG6f78stugX^vGF>^Vb+yB#O*}u_K7VJOI>}P56wpi_MLmjRZxRPvw_g2#3EO+g z05a3RZbj;=TYGjNV?a6e&1+fkPeuPTvT&BuMS@#_w94aKT2;eHo%1g#9o#DnD1;df z(|cE6J9;a-S85fxrLrknqjCLW?KTGvM!9IU)}yk?%z<;v#;7iEzHy}2R?F_`W7-r- zsQ+|5V^kb}RL6o7WhCK!0(E*{bAZ1i{Q-aLN6i@IEBwl{Fb{K#?4FIlnwM6ApC8&gc7+xIoJMUw zX{UaclR5?}oDTP#HF7?@bvqx|`8D@m15WaGTElC78d<8rq1TVayw;2e8{B@! zKOH^Hn%47og8Z{UKjA^oOw5F&$&N6898T`2lLd(uYW4;_cltaD><04$w0Sm=_}6?d zdCmC8=ApUgc4|IEp`TnZ8n?D7bp8y5!l&5N;od#gt26s&tX9jHi1t~!+m5<0W}X#X+hrD4Pp4#02Z;7PX&z= z0nYi{L~k>D^rncQXo&Om@QkuyWMyRA@`zYnLTZ;1iW%< zvFzH(WkC;FJa~?c?9IhoDb~MoLu_0i%*8a4@{a509<{m|TCd;5{!V;h(B;ciTewNg zCj5F)t$kCtzA@Z-FWhlsCsj*C8)a!;5PuV+!vA?QwvHh7&0VXDz1qj?7lBgnXsRF+(RirD^>~~wuH=qUk0`5 z{PDvk3}pymC2({-zvTlP=yq@Rq3MR_oD_EpKc9~lVQ#l4$Gl3rubYo|;(rTXe;D`N z7%kEvH_iSzmO{Fy=-E#`E8~!Zj){j>^o$?m{%#X}vV$3!H+V+v6&T+B5l6EhAOh_t2Kh$#E_6g+B3u>;(vPg(>jN zRiTh*6BZ{SiAxGaQ=5u2zVb{|Cx*@p0z1Ch#cp#})7BUES>Lyu`jyixX@`-a{t$j- z$+Vxj-|LB7sP^AHDi!6>8GE04FW;PsXN;oJNhN7kx?NMz@zsI%@M93-L0e0~GFOc`A4!1}Ch7vLuWd+khDu+oU7Y#}Tx+?tXVdn~&Xv%Q8*QUGr*sXYEApX}p z#aOsnsIB=eW9>VpuX0a}%(qk*Z6I&Ta&y^kIi8@jR8%OAMSd*bZ7m@z7{n41Lf3rM z*M8b}S3E6*)0`}K7)Ymr9**l>G!LSXE2@h8nt9?lLGIv=wf*^-ZAHa|E9d819eEdWz5n&RoH#-6wI-W5Kw{2N5CV>{ji^}j z9>|uyhPvg^>DAnPvqLun6;_zT*5&A^bEPCFsMXq{VRqJh`6xnd7E!{xJ?<%RjsLBt z7G8&Qr~c#yDGy=4(9zg<%->AQj|HPj&}plkWS7Hwi(*&>yTYvpFhN|jrtpmS{RG-X zjl=GyA)nBZL&%+YX-$r;x9P-v?a&UGC@q4zX@gFDLynz`JltP_^=;FKoqY!DC-rQ; zPRAr-;t?3y;8uNg$X_j?|0-m=*I63M#U#=)`2RZr1`?n>lk?q`=Cj~~@{{*9f12^{ zUg(X;K6657O^`a5-C8le4t;OEFcU0rY%KG%vHz1vUDRf`!gm+b^OeKEjHU`L!u(VN ze@C^!-YM73vyQ3aut}u$oKX)e!($8T)BMAbC|pjyJkajC!i48T|Fe;JAU_;1JNlj# zirn6Y0-fTCyIMKVtFVVdZ)8`3jBMd1D_C=DAjf{f?xR=k8UkB7X1hc) z$IsEwj|)wt){^xWI3^((AJpUKv_RPs6NFs*6Hxx817toKaNF`*rZ3Tgw4d&I>+cgs zAb|<8*SZXV_YA{rgj~f`~Ugx{!4}$;xb}_pKnZ=)Weow!ZKWASwU$FW=!*_qiD%#Vu=swe>GU<|z9??)MjnnRb z$4dT-T6sIO>-A4TS&CH3=&V}J+D_O#?tUC%gT}0Wp`}0FKSWnx zgT=JorES&GgiF&eGR%YIMEmZev>v0UR>dTmlU{DrDWFhr1^s1AULyItcVcl7b=zZ# zMi|v%{C|~&_FRVtf6R)S{Yy0(%Jfo~ic94FW-6O$t-8H|CxqwwQQw`g>RK(TR~6U` z_JShvoP1!qsIi^$+8o;*U3seS2DyuS-0C!Ha@hY83Q)NHsF0e^os`F3RMbh@4^1?3 zY0DJxD+EB`qT`tX-G>HuSHy=wbM8Fa&WuB>O`|$GrTkU%O`LEwHY@6HoZ|nT_f>*Pgz9_c7SaoLpuaM-KC8_%k z*Mz=V{q}ENImeA!PuB(Z3mxvua=&?4xWl=j`R`fL+MKoGR=MoIj4ThFWZixI*nCWf zF|zRO2X)@QPr&N36&LImUi*FT&ZD|tpF;%Sd;RUZbx;0D7|XBa_jw+!V)`pQ9k@s= z{ILC>;H>TY7kaX?+x9(@l6?VO*vFy1zRkd7*{MG>P5N4^6}MbBS-Bwi&B@)hprnH;$Kj_HkjyFr!G)~WH8lcG*ryOdYdUJ(7boB8j@%bG60 zmSp%b{cM}Peq|3HYee+Nb39yR)3@^OK6SwxhaX4GykxQK(}`}|O2mS6q`9X94}fcD znrH7lBM3~*Ux7p3Nb4(rIn4chTfske#Vu}51){%JHNCb1^&2#p{!R`TS+alqw*RZT z=3lqC^XOi_ocG+>+F$Nhn72%PTp2O}(E)&acl zRmE1U6aXXlOWU;O^4Y(4fvg0kG84H&6O9$;?3!k0Rqrm$op$w6-0SPt?>@Ts*y8tJ zflXoNXYS@6pZ0t2*Thmz0RhZq2cW7%A6Vu7pVYuC=eNLKAz~wN)I{+hJW7C#{Vi^v z4eGRd6LaHopP9c}>ymGBRW>0+^rO48g!Rl@xAg7#4*uUzDSG#juk^+LPW^eLSHW}mmBfVC_0idtFLF-0|L2y| zAj^l`Py8-!O})7qcvPK<_IkkNG*OWWhA_gUl9%(E6YDeF!+J7;h+i3K|wxmFne`qQ3X=F z7nHH9iV>SF@VLSY*t$R-=!d%w&0`?$D^M1@K6^8K9}lM8@XQwA1@xSVur*O?*Mq9- zdvn0+Lh(5OSOhl)0X?^L#tlw9eu9bP0l+XkKL7sTj diff --git a/python/custreamz/README.md b/python/custreamz/README.md index a1d98425d66..88657ec0d50 100644 --- a/python/custreamz/README.md +++ b/python/custreamz/README.md @@ -1,4 +1,4 @@ -#
 custreamz - GPU Accelerated Streaming
+# custreamz - GPU Accelerated Streaming Built as an extension to [python streamz](https://github.com/python-streamz/streamz), cuStreamz provides GPU accelerated abstractions for streaming data. CuStreamz can be used along side python streamz or as a standalone library for ingesting streaming data to cudf dataframes. From ded5130b60ddd1999c52089acb972d8b525b1b49 Mon Sep 17 00:00:00 2001 From: Alexander Ocsa Date: Fri, 24 Nov 2023 09:58:06 -0500 Subject: [PATCH 39/44] Introduce Comprehensive Pathological Unit Tests for Issue #14409 (#14459) This PR addresses the issue at #14409. I would like to propose the addition of unit tests that involve scenarios like having 100 or 1000 elements in a tree, reaching 100 levels of depth, with diferent data types and similar stress tests. The purpose of these tests is to conduct comprehensive testing and stress the Abstract Syntax Tree (AST), ultimately aiding in the identification and resolution of any potential issues. By introducing these pathological tests, we aim to ensure the robustness and reliability of our codebase. These tests can help us uncover edge cases and performance bottlenecks that might otherwise go unnoticed. Authors: - Alexander Ocsa (https://github.com/aocsa) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14459 --- cpp/tests/ast/transform_tests.cpp | 74 +++++++++++++++++++++++++++++++ 1 file changed, 74 insertions(+) diff --git a/cpp/tests/ast/transform_tests.cpp b/cpp/tests/ast/transform_tests.cpp index 624a781c5b9..0476cb17693 100644 --- a/cpp/tests/ast/transform_tests.cpp +++ b/cpp/tests/ast/transform_tests.cpp @@ -343,6 +343,80 @@ TEST_F(TransformTest, ImbalancedTreeArithmeticDeep) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity); } +TEST_F(TransformTest, DeeplyNestedArithmeticLogicalExpression) +{ + // Test logic for deeply nested arithmetic and logical expressions. + constexpr int64_t left_depth_level = 100; + constexpr int64_t right_depth_level = 75; + + auto generate_ast_expr = [](int64_t depth_level, + cudf::ast::column_reference col_ref, + cudf::ast::ast_operator root_operator, + cudf::ast::ast_operator arithmetic_operator, + bool nested_left_tree) { + // Note that a std::list is required here because of its guarantees against reference + // invalidation when items are added or removed. References to items in a std::vector are not + // safe if the vector must re-allocate. + auto expressions = std::list(); + + auto op = arithmetic_operator; + expressions.push_back(cudf::ast::operation(op, col_ref, col_ref)); + + for (int64_t i = 0; i < depth_level - 1; i++) { + if (i == depth_level - 2) { + op = root_operator; + } else { + op = arithmetic_operator; + } + if (nested_left_tree) { + expressions.push_back(cudf::ast::operation(op, expressions.back(), col_ref)); + } else { + expressions.push_back(cudf::ast::operation(op, col_ref, expressions.back())); + } + } + return expressions; + }; + + auto c_0 = column_wrapper{0, 0, 0}; + auto c_1 = column_wrapper{0, 0, 0}; + auto table = cudf::table_view{{c_0, c_1}}; + + auto col_ref_0 = cudf::ast::column_reference(0); + auto col_ref_1 = cudf::ast::column_reference(1); + + auto left_expression = generate_ast_expr(left_depth_level, + col_ref_0, + cudf::ast::ast_operator::LESS, + cudf::ast::ast_operator::ADD, + false); + auto right_expression = generate_ast_expr(right_depth_level, + col_ref_1, + cudf::ast::ast_operator::EQUAL, + cudf::ast::ast_operator::SUB, + true); + + auto expression_tree = cudf::ast::operation( + cudf::ast::ast_operator::LOGICAL_OR, left_expression.back(), right_expression.back()); + + // Expression: + // OR(<(+(+(+(+($0, $0), $0), $0), $0), $0), ==($1, -($1, -($1, -($1, -($1, $1)))))) + // ... + // OR(<($L, $0), ==($1, $R)) + // true + // + // Breakdown: + // - Left Operand ($L): (+(+(+(+($0, $0), $0), $0), $0), $0) + // - Right Operand ($R): -($1, -($1, -($1, -($1, $1)))) + // Explanation: + // If all $1 values and $R values are zeros, the result is true because of the equality check + // combined with the OR operator in OR(<($L, $0), ==($1, $R)). + + auto result = cudf::compute_column(table, expression_tree); + auto expected = column_wrapper{true, true, true}; + + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, result->view(), verbosity); +} + TEST_F(TransformTest, MultiLevelTreeComparator) { auto c_0 = column_wrapper{3, 20, 1, 50}; From db6745b5909233f0090d617c2eadb58a39c1348c Mon Sep 17 00:00:00 2001 From: Lawrence Mitchell Date: Fri, 24 Nov 2023 17:13:59 +0000 Subject: [PATCH 40/44] Expunge as_frame conversions in Column algorithms (#14491) Previously a number of algorithms on Columns first converted to a single column frame and called a frame-based algorithm (which calls directly into libcudf using the column we first thought of). This is unnecessary since we already have the column to hand when calling the same algorithm at the column level. Moreover, in many cases where the algorithm is a user-facing API, the frame-based approach does more work (for example conversions and dtype matching). By removing this round trip we reduce some (unnecessary) overhead, and also make the memory footprint and behaviour of column-based methods more transparent. - Closes #13565 Authors: - Lawrence Mitchell (https://github.com/wence-) Approvers: - GALI PREM SAGAR (https://github.com/galipremsagar) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cudf/pull/14491 --- python/cudf/cudf/core/column/column.py | 27 ++++++++++++------- python/cudf/cudf/core/column/decimal.py | 14 ++++------ .../cudf/cudf/core/column/numerical_base.py | 11 +++----- python/cudf/cudf/core/indexed_frame.py | 5 +++- python/cudf/cudf/tests/test_search.py | 5 +++- 5 files changed, 34 insertions(+), 28 deletions(-) diff --git a/python/cudf/cudf/core/column/column.py b/python/cudf/cudf/core/column/column.py index b4f65693d85..a4fa31e5b25 100644 --- a/python/cudf/cudf/core/column/column.py +++ b/python/cudf/cudf/core/column/column.py @@ -943,14 +943,14 @@ def is_unique(self) -> bool: @property def is_monotonic_increasing(self) -> bool: - return not self.has_nulls() and self.as_frame()._is_sorted( - ascending=None, null_position=None + return not self.has_nulls() and libcudf.sort.is_sorted( + [self], [True], None ) @property def is_monotonic_decreasing(self) -> bool: - return not self.has_nulls() and self.as_frame()._is_sorted( - ascending=[False], null_position=None + return not self.has_nulls() and libcudf.sort.is_sorted( + [self], [False], None ) def sort_values( @@ -1134,8 +1134,8 @@ def apply_boolean_mask(self, mask) -> ColumnBase: def argsort( self, ascending: bool = True, na_position: str = "last" ) -> "cudf.core.column.NumericalColumn": - return self.as_frame()._get_sorted_inds( - ascending=ascending, na_position=na_position + return libcudf.sort.order_by( + [self], [ascending], na_position, stable=True ) def __arrow_array__(self, type=None): @@ -1161,10 +1161,17 @@ def searchsorted( side: str = "left", ascending: bool = True, na_position: str = "last", - ): - values = as_column(value).as_frame() - return self.as_frame().searchsorted( - values, side, ascending=ascending, na_position=na_position + ) -> Self: + if not isinstance(value, ColumnBase) or value.dtype != self.dtype: + raise ValueError( + "Column searchsorted expects values to be column of same dtype" + ) + return libcudf.search.search_sorted( + [self], + [value], + side=side, + ascending=ascending, + na_position=na_position, ) def unique(self) -> ColumnBase: diff --git a/python/cudf/cudf/core/column/decimal.py b/python/cudf/cudf/core/column/decimal.py index 5a823c5f7c3..972280b041f 100644 --- a/python/cudf/cudf/core/column/decimal.py +++ b/python/cudf/cudf/core/column/decimal.py @@ -10,7 +10,6 @@ import cudf from cudf import _lib as libcudf -from cudf._lib.quantiles import quantile as cpp_quantile from cudf._lib.strings.convert.convert_fixed_point import ( from_decimal as cpp_from_decimal, ) @@ -192,15 +191,12 @@ def _decimal_quantile( ) -> ColumnBase: quant = [float(q)] if not isinstance(q, (Sequence, np.ndarray)) else q # get sorted indices and exclude nulls - sorted_indices = self.as_frame()._get_sorted_inds( - ascending=True, na_position="first" + indices = libcudf.sort.order_by( + [self], [True], "first", stable=True + ).slice(self.null_count, len(self)) + result = libcudf.quantiles.quantile( + self, quant, interpolation, indices, exact ) - sorted_indices = sorted_indices[self.null_count :] - - result = cpp_quantile( - self, quant, interpolation, sorted_indices, exact - ) - return result._with_type_metadata(self.dtype) def as_numerical_column( diff --git a/python/cudf/cudf/core/column/numerical_base.py b/python/cudf/cudf/core/column/numerical_base.py index f3c5f99ae1e..c45a9c7fd5d 100644 --- a/python/cudf/cudf/core/column/numerical_base.py +++ b/python/cudf/cudf/core/column/numerical_base.py @@ -182,15 +182,12 @@ def _numeric_quantile( self, q: np.ndarray, interpolation: str, exact: bool ) -> NumericalBaseColumn: # get sorted indices and exclude nulls - sorted_indices = self.as_frame()._get_sorted_inds( - ascending=True, na_position="first" - ) - sorted_indices = sorted_indices.slice( - self.null_count, len(sorted_indices) - ) + indices = libcudf.sort.order_by( + [self], [True], "first", stable=True + ).slice(self.null_count, len(self)) return libcudf.quantiles.quantile( - self, q, interpolation, sorted_indices, exact + self, q, interpolation, indices, exact ) def cov(self, other: NumericalBaseColumn) -> float: diff --git a/python/cudf/cudf/core/indexed_frame.py b/python/cudf/cudf/core/indexed_frame.py index 4211a8c24bf..5ef6dc2a6e3 100644 --- a/python/cudf/cudf/core/indexed_frame.py +++ b/python/cudf/cudf/core/indexed_frame.py @@ -3260,8 +3260,11 @@ def _first_or_last( # is on the end of the offset. See pandas gh29623 for detail. to_search = to_search - pd_offset.base return self.loc[:to_search] + needle = as_column(to_search, dtype=self._index.dtype) end_point = int( - self._index._column.searchsorted(to_search, side=side)[0] + self._index._column.searchsorted( + needle, side=side + ).element_indexing(0) ) return slice_func(end_point) diff --git a/python/cudf/cudf/tests/test_search.py b/python/cudf/cudf/tests/test_search.py index b0eacb1a709..17cf3cf8141 100644 --- a/python/cudf/cudf/tests/test_search.py +++ b/python/cudf/cudf/tests/test_search.py @@ -38,7 +38,10 @@ def test_searchsorted(side, obj_class, vals_class): pvals = vals.to_pandas() expect = psr.searchsorted(pvals, side) - got = sr.searchsorted(vals, side) + if obj_class == "column": + got = sr.searchsorted(vals._column, side) + else: + got = sr.searchsorted(vals, side) assert_eq(expect, cupy.asnumpy(got)) From c8d481e24a8cf6054cb9400213df00a4b42a1566 Mon Sep 17 00:00:00 2001 From: Ashwin Srinath <3190405+shwina@users.noreply.github.com> Date: Mon, 27 Nov 2023 12:34:55 -0500 Subject: [PATCH 41/44] Enable `pd.Timestamp` objects to be picklable when `cudf.pandas` is active (#14474) Closes https://github.com/rapidsai/cudf/issues/14471. This PR makes `Timestamp` objects picklable by registering a custom reducer for `pd.Timestamp` objects when `cudf.pandas` is active. Authors: - Ashwin Srinath (https://github.com/shwina) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) URL: https://github.com/rapidsai/cudf/pull/14474 --- python/cudf/cudf/pandas/_wrappers/pandas.py | 31 ++++++++++++++++++- .../cudf_pandas_tests/test_cudf_pandas.py | 4 +++ 2 files changed, 34 insertions(+), 1 deletion(-) diff --git a/python/cudf/cudf/pandas/_wrappers/pandas.py b/python/cudf/cudf/pandas/_wrappers/pandas.py index 71daf1e6f0d..193ef404a8c 100644 --- a/python/cudf/cudf/pandas/_wrappers/pandas.py +++ b/python/cudf/cudf/pandas/_wrappers/pandas.py @@ -1,7 +1,8 @@ # SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. # All rights reserved. # SPDX-License-Identifier: Apache-2.0 - +import copyreg +import pickle import sys import pandas as pd @@ -1304,3 +1305,31 @@ def _df_query_method(self, *args, local_dict=None, global_dict=None, **kwargs): _Unusable, typ, ) + +# timestamps and timedeltas are not proxied, but non-proxied +# pandas types are currently not picklable. Thus, we define +# custom reducer/unpicker functions for these types: +def _reduce_obj(obj): + from cudf.pandas.module_accelerator import disable_module_accelerator + + with disable_module_accelerator(): + # args can contain objects that are unpicklable + # when the module accelerator is disabled + # (freq is of a proxy type): + pickled_args = pickle.dumps(obj.__reduce__()) + + return _unpickle_obj, (pickled_args,) + + +def _unpickle_obj(pickled_args): + from cudf.pandas.module_accelerator import disable_module_accelerator + + with disable_module_accelerator(): + unpickler, args = pickle.loads(pickled_args) + obj = unpickler(*args) + return obj + + +copyreg.dispatch_table[pd.Timestamp] = _reduce_obj +# same reducer/unpickler can be used for Timedelta: +copyreg.dispatch_table[pd.Timedelta] = _reduce_obj diff --git a/python/cudf/cudf_pandas_tests/test_cudf_pandas.py b/python/cudf/cudf_pandas_tests/test_cudf_pandas.py index 2500ba07bd9..0dbf2c305e5 100644 --- a/python/cudf/cudf_pandas_tests/test_cudf_pandas.py +++ b/python/cudf/cudf_pandas_tests/test_cudf_pandas.py @@ -1076,6 +1076,10 @@ def test_np_array_of_timestamps(): xpd.Index(["a", 2, 3]), # Other types xpd.tseries.offsets.BDay(5), + xpd.Timestamp("2001-01-01"), + xpd.Timestamp("2001-01-01", freq="D"), + xpd.Timedelta("1 days"), + xpd.Timedelta(1, "D"), ], ) def test_pickle(obj): From dba69341b45f5c925630d5e344a7a5e354318510 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Mon, 27 Nov 2023 12:50:58 -0800 Subject: [PATCH 42/44] Some doxygen improvements (#14469) This PR fixes some errors in the doxygen docs and adds groups for some items that were previously missing altogether. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - David Wendt (https://github.com/davidwendt) URL: https://github.com/rapidsai/cudf/pull/14469 --- cpp/include/cudf/ast/expressions.hpp | 6 ++++++ cpp/include/cudf/column/column_device_view.cuh | 2 ++ cpp/include/cudf/dictionary/update_keys.hpp | 12 ++++++------ cpp/include/cudf/io/types.hpp | 7 +++++++ cpp/include/cudf/lists/lists_column_view.hpp | 6 +++--- cpp/include/cudf/strings/strings_column_view.hpp | 4 ++-- cpp/include/cudf/structs/structs_column_view.hpp | 2 +- cpp/include/cudf/utilities/default_stream.hpp | 8 +++++++- cpp/include/cudf/utilities/span.hpp | 8 ++++++++ cpp/include/doxygen_groups.h | 3 +++ 10 files changed, 45 insertions(+), 13 deletions(-) diff --git a/cpp/include/cudf/ast/expressions.hpp b/cpp/include/cudf/ast/expressions.hpp index c5172486fa6..26916e49012 100644 --- a/cpp/include/cudf/ast/expressions.hpp +++ b/cpp/include/cudf/ast/expressions.hpp @@ -25,6 +25,11 @@ namespace cudf { namespace ast { +/** + * @addtogroup expressions + * @{ + * @file + */ // Forward declaration. namespace detail { @@ -544,6 +549,7 @@ class column_name_reference : public expression { std::string column_name; }; +/** @} */ // end of group } // namespace ast } // namespace cudf diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index b1ff0bbaea7..daee443a5f3 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -52,6 +52,8 @@ namespace cudf { * If used at compile-time, this indicator can tell the optimizer * to include or exclude any null-checking clauses. * + * @ingroup column_classes + * */ struct nullate { struct YES : std::bool_constant {}; diff --git a/cpp/include/cudf/dictionary/update_keys.hpp b/cpp/include/cudf/dictionary/update_keys.hpp index 81728e1ff73..40504c22edd 100644 --- a/cpp/include/cudf/dictionary/update_keys.hpp +++ b/cpp/include/cudf/dictionary/update_keys.hpp @@ -46,9 +46,9 @@ namespace dictionary { * Null entries from the input column are copied to the output column. * No new null entries are created by this operation. * - * @throw cudf_logic_error if the new_keys type does not match the keys type in + * @throw cudf::logic_error if the new_keys type does not match the keys type in * the dictionary_column. - * @throw cudf_logic_error if the new_keys contain nulls. + * @throw cudf::logic_error if the new_keys contain nulls. * * @param dictionary_column Existing dictionary column. * @param new_keys New keys to incorporate into the dictionary_column. @@ -78,9 +78,9 @@ std::unique_ptr add_keys( * @endcode * Note that "a" has been removed so output row[2] becomes null. * - * @throw cudf_logic_error if the keys_to_remove type does not match the keys type in + * @throw cudf::logic_error if the keys_to_remove type does not match the keys type in * the dictionary_column. - * @throw cudf_logic_error if the keys_to_remove contain nulls. + * @throw cudf::logic_error if the keys_to_remove contain nulls. * * @param dictionary_column Existing dictionary column. * @param keys_to_remove The keys to remove from the dictionary_column. @@ -134,9 +134,9 @@ std::unique_ptr remove_unused_keys( * d2 is now {keys=["b", "c", "d"], indices=[1, x, 0, 1, 0], valids=[1, 0, 1, 1, 1]} * @endcode * - * @throw cudf_logic_error if the keys type does not match the keys type in + * @throw cudf::logic_error if the keys type does not match the keys type in * the dictionary_column. - * @throw cudf_logic_error if the keys contain nulls. + * @throw cudf::logic_error if the keys contain nulls. * * @param dictionary_column Existing dictionary column. * @param keys New keys to use for the output column. Must not contain nulls. diff --git a/cpp/include/cudf/io/types.hpp b/cpp/include/cudf/io/types.hpp index 50119e60882..d790f869e22 100644 --- a/cpp/include/cudf/io/types.hpp +++ b/cpp/include/cudf/io/types.hpp @@ -44,6 +44,12 @@ class datasource; namespace cudf { //! IO interfaces namespace io { +/** + * @addtogroup io_types + * @{ + * @file + */ + /** * @brief Compression algorithms */ @@ -938,5 +944,6 @@ class reader_column_schema { [[nodiscard]] size_t get_num_children() const { return children.size(); } }; +/** @} */ // end of group } // namespace io } // namespace cudf diff --git a/cpp/include/cudf/lists/lists_column_view.hpp b/cpp/include/cudf/lists/lists_column_view.hpp index 8c6368eacb6..57a4f724c2d 100644 --- a/cpp/include/cudf/lists/lists_column_view.hpp +++ b/cpp/include/cudf/lists/lists_column_view.hpp @@ -83,7 +83,7 @@ class lists_column_view : private column_view { /** * @brief Returns the internal column of offsets * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * @return The internal column of offsets */ [[nodiscard]] column_view offsets() const; @@ -91,7 +91,7 @@ class lists_column_view : private column_view { /** * @brief Returns the internal child column * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * @return The internal child column */ [[nodiscard]] column_view child() const; @@ -104,7 +104,7 @@ class lists_column_view : private column_view { * the child columns when recursing. Most functions operating in a recursive manner * on lists columns should be using `get_sliced_child()` instead of `child()`. * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * @param stream CUDA stream used for device memory operations and kernel launches * @return A sliced child column view */ diff --git a/cpp/include/cudf/strings/strings_column_view.hpp b/cpp/include/cudf/strings/strings_column_view.hpp index f1aa8e49f00..e27d32fceb9 100644 --- a/cpp/include/cudf/strings/strings_column_view.hpp +++ b/cpp/include/cudf/strings/strings_column_view.hpp @@ -80,7 +80,7 @@ class strings_column_view : private column_view { /** * @brief Returns the internal column of offsets * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * @return The offsets column */ [[nodiscard]] column_view offsets() const; @@ -106,7 +106,7 @@ class strings_column_view : private column_view { /** * @brief Returns the internal column of chars * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * @return The chars column */ [[nodiscard]] column_view chars() const; diff --git a/cpp/include/cudf/structs/structs_column_view.hpp b/cpp/include/cudf/structs/structs_column_view.hpp index 6a9f2890177..4a50488ef00 100644 --- a/cpp/include/cudf/structs/structs_column_view.hpp +++ b/cpp/include/cudf/structs/structs_column_view.hpp @@ -87,7 +87,7 @@ class structs_column_view : public column_view { * the child columns when recursing. Most functions operating in a recursive manner * on struct columns should be using `get_sliced_child()` instead of `child()`. * - * @throw cudf::logic error if this is an empty column + * @throw cudf::logic_error if this is an empty column * * @param index The index of the child column to return * @param stream The stream on which to perform the operation. Uses the default CUDF diff --git a/cpp/include/cudf/utilities/default_stream.hpp b/cpp/include/cudf/utilities/default_stream.hpp index 1eec3b994d0..aacab996e8a 100644 --- a/cpp/include/cudf/utilities/default_stream.hpp +++ b/cpp/include/cudf/utilities/default_stream.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -22,6 +22,11 @@ #include namespace cudf { +/** + * @addtogroup default_stream + * @{ + * @file + */ /** * @brief Get the current default stream @@ -37,4 +42,5 @@ rmm::cuda_stream_view const get_default_stream(); */ bool is_ptds_enabled(); +/** @} */ // end of group } // namespace cudf diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp index 3e5f6e3e97a..a6190ee02cb 100644 --- a/cpp/include/cudf/utilities/span.hpp +++ b/cpp/include/cudf/utilities/span.hpp @@ -185,6 +185,13 @@ class span_base { } // namespace detail +/** + * @addtogroup utility_span + * @{ + * @file + * @brief APIs for labeling values by bin. + */ + // ===== host_span ================================================================================= template @@ -318,6 +325,7 @@ struct device_span : public cudf::detail::span_base Date: Mon, 27 Nov 2023 11:20:27 -1000 Subject: [PATCH 43/44] REF: Remove **kwargs from to_pandas, raise if nullable is not implemented (#14438) * Remove unnecessary **kwargs from signature * Typing and improve error message if nullable is not implemented Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14438 --- python/cudf/cudf/core/_base_index.py | 2 +- python/cudf/cudf/core/column/categorical.py | 5 +++- python/cudf/cudf/core/column/column.py | 13 ++++++---- python/cudf/cudf/core/column/datetime.py | 17 ++++++++---- python/cudf/cudf/core/column/interval.py | 4 ++- python/cudf/cudf/core/column/numerical.py | 4 +-- python/cudf/cudf/core/column/string.py | 4 +-- python/cudf/cudf/core/column/struct.py | 10 +++---- python/cudf/cudf/core/column/timedelta.py | 5 +++- python/cudf/cudf/core/dataframe.py | 2 +- python/cudf/cudf/core/index.py | 17 +++++++----- python/cudf/cudf/core/multiindex.py | 2 +- python/cudf/cudf/core/series.py | 6 ++++- python/cudf/cudf/io/json.py | 16 +++++++++++- python/cudf/cudf/testing/testing.py | 8 ++++-- .../cudf/tests/dataframe/test_conversion.py | 12 +++++++-- .../cudf/cudf/tests/series/test_conversion.py | 3 ++- .../cudf/tests/series/test_datetimelike.py | 13 ++++++++++ python/cudf/cudf/tests/test_categorical.py | 4 +-- python/cudf/cudf/tests/test_dataframe.py | 6 ++--- python/cudf/cudf/tests/test_index.py | 26 +++++++++++++++---- python/cudf/cudf/tests/test_list.py | 6 ++--- python/cudf/cudf/tests/test_parquet.py | 10 +++++-- python/cudf/cudf/tests/test_udf_masked_ops.py | 6 ++--- python/dask_cudf/dask_cudf/core.py | 6 ++--- .../dask_cudf/tests/test_accessor.py | 4 +-- 26 files changed, 147 insertions(+), 64 deletions(-) diff --git a/python/cudf/cudf/core/_base_index.py b/python/cudf/cudf/core/_base_index.py index 3616ec1b542..8387ef96dfa 100644 --- a/python/cudf/cudf/core/_base_index.py +++ b/python/cudf/cudf/core/_base_index.py @@ -850,7 +850,7 @@ def notna(self): """ raise NotImplementedError - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False): """ Convert to a Pandas Index. diff --git a/python/cudf/cudf/core/column/categorical.py b/python/cudf/cudf/core/column/categorical.py index bab07624dfa..689cb02ed86 100644 --- a/python/cudf/cudf/core/column/categorical.py +++ b/python/cudf/cudf/core/column/categorical.py @@ -974,8 +974,11 @@ def __cuda_array_interface__(self) -> Mapping[str, Any]: ) def to_pandas( - self, index: Optional[pd.Index] = None, **kwargs + self, *, index: Optional[pd.Index] = None, nullable: bool = False ) -> pd.Series: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") + if self.categories.dtype.kind == "f": new_mask = bools_to_mask(self.notnull()) col = column.build_categorical_column( diff --git a/python/cudf/cudf/core/column/column.py b/python/cudf/cudf/core/column/column.py index a4fa31e5b25..015078aea17 100644 --- a/python/cudf/cudf/core/column/column.py +++ b/python/cudf/cudf/core/column/column.py @@ -203,17 +203,20 @@ def __repr__(self): ) def to_pandas( - self, index: Optional[pd.Index] = None, **kwargs + self, + *, + index: Optional[pd.Index] = None, + nullable: bool = False, ) -> pd.Series: """Convert object to pandas type. The default implementation falls back to PyArrow for the conversion. """ # This default implementation does not handle nulls in any meaningful - # way, but must consume the parameter to avoid passing it to PyArrow - # (which does not recognize it). - kwargs.pop("nullable", None) - pd_series = self.to_arrow().to_pandas(**kwargs) + # way + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") + pd_series = self.to_arrow().to_pandas() if index is not None: pd_series.index = index diff --git a/python/cudf/cudf/core/column/datetime.py b/python/cudf/cudf/core/column/datetime.py index 33354e1c3bc..7980b58ab8b 100644 --- a/python/cudf/cudf/core/column/datetime.py +++ b/python/cudf/cudf/core/column/datetime.py @@ -308,10 +308,12 @@ def day_of_year(self) -> ColumnBase: def to_pandas( self, + *, index: Optional[pd.Index] = None, nullable: bool = False, - **kwargs, - ) -> "cudf.Series": + ) -> pd.Series: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") # `copy=True` workaround until following issue is fixed: # https://issues.apache.org/jira/browse/ARROW-9772 @@ -686,13 +688,18 @@ def __init__( def to_pandas( self, + *, index: Optional[pd.Index] = None, nullable: bool = False, - **kwargs, - ) -> "cudf.Series": - return self._local_time.to_pandas().dt.tz_localize( + ) -> pd.Series: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") + series = self._local_time.to_pandas().dt.tz_localize( self.dtype.tz, ambiguous="NaT", nonexistent="NaT" ) + if index is not None: + series.index = index + return series def to_arrow(self): return pa.compute.assume_timezone( diff --git a/python/cudf/cudf/core/column/interval.py b/python/cudf/cudf/core/column/interval.py index d4855def832..b550e272b2c 100644 --- a/python/cudf/cudf/core/column/interval.py +++ b/python/cudf/cudf/core/column/interval.py @@ -126,13 +126,15 @@ def as_interval_column(self, dtype, **kwargs): raise ValueError("dtype must be IntervalDtype") def to_pandas( - self, index: Optional[pd.Index] = None, **kwargs + self, *, index: Optional[pd.Index] = None, nullable: bool = False ) -> pd.Series: # Note: This does not handle null values in the interval column. # However, this exact sequence (calling __from_arrow__ on the output of # self.to_arrow) is currently the best known way to convert interval # types into pandas (trying to convert the underlying numerical columns # directly is problematic), so we're stuck with this for now. + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") return pd.Series( self.dtype.to_pandas().__from_arrow__(self.to_arrow()), index=index ) diff --git a/python/cudf/cudf/core/column/numerical.py b/python/cudf/cudf/core/column/numerical.py index 4ea49c8c7c0..cdecf44cc8f 100644 --- a/python/cudf/cudf/core/column/numerical.py +++ b/python/cudf/cudf/core/column/numerical.py @@ -679,9 +679,9 @@ def _with_type_metadata(self: ColumnBase, dtype: Dtype) -> ColumnBase: def to_pandas( self, + *, index: Optional[pd.Index] = None, nullable: bool = False, - **kwargs, ) -> pd.Series: if nullable and self.dtype in np_dtypes_to_pandas_dtypes: pandas_nullable_dtype = np_dtypes_to_pandas_dtypes[self.dtype] @@ -691,7 +691,7 @@ def to_pandas( elif str(self.dtype) in NUMERIC_TYPES and not self.has_nulls(): pd_series = pd.Series(self.values_host, copy=False) else: - pd_series = self.to_arrow().to_pandas(**kwargs) + pd_series = self.to_arrow().to_pandas() if index is not None: pd_series.index = index diff --git a/python/cudf/cudf/core/column/string.py b/python/cudf/cudf/core/column/string.py index eb86f555432..a5e9bbf81f2 100644 --- a/python/cudf/cudf/core/column/string.py +++ b/python/cudf/cudf/core/column/string.py @@ -5756,15 +5756,15 @@ def values(self) -> cupy.ndarray: def to_pandas( self, + *, index: Optional[pd.Index] = None, nullable: bool = False, - **kwargs, ) -> pd.Series: if nullable: pandas_array = pd.StringDtype().__from_arrow__(self.to_arrow()) pd_series = pd.Series(pandas_array, copy=False) else: - pd_series = self.to_arrow().to_pandas(**kwargs) + pd_series = self.to_arrow().to_pandas() if index is not None: pd_series.index = index diff --git a/python/cudf/cudf/core/column/struct.py b/python/cudf/cudf/core/column/struct.py index 0bb21f4c25a..35fb18cc5c7 100644 --- a/python/cudf/cudf/core/column/struct.py +++ b/python/cudf/cudf/core/column/struct.py @@ -59,16 +59,14 @@ def to_arrow(self): ) def to_pandas( - self, index: Optional[pd.Index] = None, **kwargs + self, *, index: Optional[pd.Index] = None, nullable: bool = False ) -> pd.Series: # We cannot go via Arrow's `to_pandas` because of the following issue: # https://issues.apache.org/jira/browse/ARROW-12680 + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") - pd_series = pd.Series(self.to_arrow().tolist(), dtype="object") - - if index is not None: - pd_series.index = index - return pd_series + return pd.Series(self.to_arrow().tolist(), dtype="object", index=index) @cached_property def memory_usage(self): diff --git a/python/cudf/cudf/core/column/timedelta.py b/python/cudf/cudf/core/column/timedelta.py index 13bb97b9a89..572b3b894dc 100644 --- a/python/cudf/cudf/core/column/timedelta.py +++ b/python/cudf/cudf/core/column/timedelta.py @@ -144,7 +144,7 @@ def to_arrow(self) -> pa.Array: ) def to_pandas( - self, index=None, nullable: bool = False, **kwargs + self, *, index: Optional[pd.Index] = None, nullable: bool = False ) -> pd.Series: # `copy=True` workaround until following issue is fixed: # https://issues.apache.org/jira/browse/ARROW-9772 @@ -152,6 +152,9 @@ def to_pandas( # Pandas only supports `timedelta64[ns]` dtype # and conversion to this type is necessary to make # arrow to pandas conversion happen for large values. + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") + return pd.Series( self.astype("timedelta64[ns]").to_arrow(), copy=True, diff --git a/python/cudf/cudf/core/dataframe.py b/python/cudf/cudf/core/dataframe.py index 43ae9b9e81e..785f3d98712 100644 --- a/python/cudf/cudf/core/dataframe.py +++ b/python/cudf/cudf/core/dataframe.py @@ -5128,7 +5128,7 @@ def describe( ) @_cudf_nvtx_annotate - def to_pandas(self, nullable=False, **kwargs): + def to_pandas(self, *, nullable: bool = False) -> pd.DataFrame: """ Convert to a Pandas DataFrame. diff --git a/python/cudf/cudf/core/index.py b/python/cudf/cudf/core/index.py index 277b5d3bb17..98d537b2a0f 100644 --- a/python/cudf/cudf/core/index.py +++ b/python/cudf/cudf/core/index.py @@ -501,7 +501,9 @@ def dtype(self): return _maybe_convert_to_default_type(dtype) @_cudf_nvtx_annotate - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False) -> pd.RangeIndex: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") return pd.RangeIndex( start=self._start, stop=self._stop, @@ -1567,7 +1569,7 @@ def _clean_nulls_from_index(self): def any(self): return self._values.any() - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False) -> pd.Index: return pd.Index( self._values.to_pandas(nullable=nullable), name=self.name ) @@ -2508,7 +2510,9 @@ def isocalendar(self): return cudf.core.tools.datetimes._to_iso_calendar(self) @_cudf_nvtx_annotate - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False) -> pd.DatetimeIndex: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") # TODO: no need to convert to nanos with Pandas 2.x if isinstance(self.dtype, pd.DatetimeTZDtype): nanos = self._values.astype( @@ -2832,11 +2836,12 @@ def __getitem__(self, index): return value @_cudf_nvtx_annotate - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False) -> pd.TimedeltaIndex: + if nullable: + raise NotImplementedError(f"{nullable=} is not implemented.") return pd.TimedeltaIndex( self._values.to_pandas(), name=self.name, - unit=self._values.time_unit, ) @property # type: ignore @@ -3300,7 +3305,7 @@ def __init__(self, values, copy=False, **kwargs): super().__init__(values, **kwargs) @_cudf_nvtx_annotate - def to_pandas(self, nullable=False): + def to_pandas(self, *, nullable: bool = False) -> pd.Index: return pd.Index( self.to_numpy(na_value=None), name=self.name, diff --git a/python/cudf/cudf/core/multiindex.py b/python/cudf/cudf/core/multiindex.py index d0c8a513686..820f6e7769c 100644 --- a/python/cudf/cudf/core/multiindex.py +++ b/python/cudf/cudf/core/multiindex.py @@ -1592,7 +1592,7 @@ def droplevel(self, level=-1): return mi @_cudf_nvtx_annotate - def to_pandas(self, nullable=False, **kwargs): + def to_pandas(self, *, nullable: bool = False) -> pd.MultiIndex: result = self.to_frame( index=False, name=list(range(self.nlevels)) ).to_pandas(nullable=nullable) diff --git a/python/cudf/cudf/core/series.py b/python/cudf/cudf/core/series.py index 04a7ed3abf7..b3782aae24a 100644 --- a/python/cudf/cudf/core/series.py +++ b/python/cudf/cudf/core/series.py @@ -2064,7 +2064,9 @@ def any(self, axis=0, bool_only=None, skipna=True, level=None, **kwargs): return super().any(axis, skipna, level, **kwargs) @_cudf_nvtx_annotate - def to_pandas(self, index=True, nullable=False, **kwargs): + def to_pandas( + self, *, index: bool = True, nullable: bool = False + ) -> pd.Series: """ Convert to a Pandas Series. @@ -2126,6 +2128,8 @@ def to_pandas(self, index=True, nullable=False, **kwargs): """ if index is True: index = self.index.to_pandas() + else: + index = None # type: ignore[assignment] s = self._column.to_pandas(index=index, nullable=nullable) s.name = self.name return s diff --git a/python/cudf/cudf/io/json.py b/python/cudf/cudf/io/json.py index efac24aee17..ae2f0203642 100644 --- a/python/cudf/cudf/io/json.py +++ b/python/cudf/cudf/io/json.py @@ -179,6 +179,13 @@ def read_json( return df +def maybe_return_nullable_pd_obj(cudf_obj): + try: + return cudf_obj.to_pandas(nullable=True) + except NotImplementedError: + return cudf_obj.to_pandas(nullable=False) + + @ioutils.doc_to_json() def to_json( cudf_val, @@ -228,7 +235,14 @@ def to_json( return path_or_buf.read() elif engine == "pandas": warnings.warn("Using CPU via Pandas to write JSON dataset") - pd_value = cudf_val.to_pandas(nullable=True) + if isinstance(cudf_val, cudf.DataFrame): + pd_data = { + col: maybe_return_nullable_pd_obj(series) + for col, series in cudf_val.items() + } + pd_value = pd.DataFrame(pd_data) + else: + pd_value = maybe_return_nullable_pd_obj(cudf_val) return pd.io.json.to_json( path_or_buf, pd_value, diff --git a/python/cudf/cudf/testing/testing.py b/python/cudf/cudf/testing/testing.py index a9c54ddcaa1..a6262e99481 100644 --- a/python/cudf/cudf/testing/testing.py +++ b/python/cudf/cudf/testing/testing.py @@ -271,8 +271,12 @@ def assert_column_equal( left = left.astype(left.categories.dtype) right = right.astype(right.categories.dtype) if not columns_equal: - ldata = str([val for val in left.to_pandas(nullable=True)]) - rdata = str([val for val in right.to_pandas(nullable=True)]) + try: + ldata = str([val for val in left.to_pandas(nullable=True)]) + rdata = str([val for val in right.to_pandas(nullable=True)]) + except NotImplementedError: + ldata = str([val for val in left.to_pandas(nullable=False)]) + rdata = str([val for val in right.to_pandas(nullable=False)]) try: diff = 0 for i in range(left.size): diff --git a/python/cudf/cudf/tests/dataframe/test_conversion.py b/python/cudf/cudf/tests/dataframe/test_conversion.py index 3673ea827f9..fa7e5ec1d4c 100644 --- a/python/cudf/cudf/tests/dataframe/test_conversion.py +++ b/python/cudf/cudf/tests/dataframe/test_conversion.py @@ -1,5 +1,6 @@ # Copyright (c) 2023, NVIDIA CORPORATION. import pandas as pd +import pytest import cudf from cudf.testing._utils import assert_eq @@ -26,6 +27,9 @@ def test_convert_dtypes(): "category", "datetime64[ns]", ] + nullable_columns = list("abcdef") + non_nullable_columns = list(set(data.keys()).difference(nullable_columns)) + df = pd.DataFrame( { k: pd.Series(v, dtype=d) @@ -33,6 +37,10 @@ def test_convert_dtypes(): } ) gdf = cudf.DataFrame.from_pandas(df) - expect = df.convert_dtypes() - got = gdf.convert_dtypes().to_pandas(nullable=True) + expect = df[nullable_columns].convert_dtypes() + got = gdf[nullable_columns].convert_dtypes().to_pandas(nullable=True) assert_eq(expect, got) + + with pytest.raises(NotImplementedError): + # category and datetime64[ns] are not nullable + gdf[non_nullable_columns].convert_dtypes().to_pandas(nullable=True) diff --git a/python/cudf/cudf/tests/series/test_conversion.py b/python/cudf/cudf/tests/series/test_conversion.py index 08124a9a98e..43ac35e41a6 100644 --- a/python/cudf/cudf/tests/series/test_conversion.py +++ b/python/cudf/cudf/tests/series/test_conversion.py @@ -26,7 +26,8 @@ def test_convert_dtypes(data, dtype): # because we don't have distinct nullable types, we check that we # get the same result if we convert to nullable pandas types: - got = gs.convert_dtypes().to_pandas(nullable=True) + nullable = dtype not in ("category", "datetime64[ns]") + got = gs.convert_dtypes().to_pandas(nullable=nullable) assert_eq(expect, got) diff --git a/python/cudf/cudf/tests/series/test_datetimelike.py b/python/cudf/cudf/tests/series/test_datetimelike.py index 85da985940f..df68eaca399 100644 --- a/python/cudf/cudf/tests/series/test_datetimelike.py +++ b/python/cudf/cudf/tests/series/test_datetimelike.py @@ -180,6 +180,19 @@ def test_convert_edge_cases(data, original_timezone, target_timezone): assert_eq(expect, got) +def test_to_pandas_index_true_timezone(): + data = [ + "2008-05-12", + "2008-12-12", + "2009-05-12", + ] + dti = cudf.DatetimeIndex(data).tz_localize("UTC") + ser = cudf.Series(dti, index=list("abc")) + result = ser.to_pandas(index=True) + expected = pd.Series(pd.to_datetime(data, utc=True), index=list("abc")) + assert_eq(result, expected) + + def test_tz_aware_attributes_local(): data = [ "2008-05-12 13:50:00", diff --git a/python/cudf/cudf/tests/test_categorical.py b/python/cudf/cudf/tests/test_categorical.py index afbc0dc6c17..49eeff01bee 100644 --- a/python/cudf/cudf/tests/test_categorical.py +++ b/python/cudf/cudf/tests/test_categorical.py @@ -227,7 +227,7 @@ def test_df_cat_set_index(): df["b"] = np.arange(len(df)) got = df.set_index("a") - pddf = df.to_pandas(nullable_pd_dtype=False) + pddf = df.to_pandas() expect = pddf.set_index("a") assert_eq(got, expect) @@ -239,7 +239,7 @@ def test_df_cat_sort_index(): df["b"] = np.arange(len(df)) got = df.set_index("a").sort_index() - expect = df.to_pandas(nullable_pd_dtype=False).set_index("a").sort_index() + expect = df.to_pandas().set_index("a").sort_index() assert_eq(got, expect) diff --git a/python/cudf/cudf/tests/test_dataframe.py b/python/cudf/cudf/tests/test_dataframe.py index 74165731683..97c89217f9f 100644 --- a/python/cudf/cudf/tests/test_dataframe.py +++ b/python/cudf/cudf/tests/test_dataframe.py @@ -2163,9 +2163,9 @@ def test_dataframe_transpose(nulls, num_cols, num_rows, dtype): got_property = gdf.T expect = pdf.transpose() - - assert_eq(expect, got_function.to_pandas(nullable=True)) - assert_eq(expect, got_property.to_pandas(nullable=True)) + nullable = dtype not in DATETIME_TYPES + assert_eq(expect, got_function.to_pandas(nullable=nullable)) + assert_eq(expect, got_property.to_pandas(nullable=nullable)) @pytest.mark.parametrize("num_cols", [1, 2, 10]) diff --git a/python/cudf/cudf/tests/test_index.py b/python/cudf/cudf/tests/test_index.py index 087b93f1a02..c393522c28b 100644 --- a/python/cudf/cudf/tests/test_index.py +++ b/python/cudf/cudf/tests/test_index.py @@ -2750,11 +2750,14 @@ def test_index_error_list_index(): ) def test_index_hasnans(data): gs = cudf.Index(data, nan_as_null=False) - ps = gs.to_pandas(nullable=True) - - # Check type to avoid mixing Python bool and NumPy bool - assert isinstance(gs.hasnans, bool) - assert gs.hasnans == ps.hasnans + if isinstance(gs, cudf.RangeIndex): + with pytest.raises(NotImplementedError): + gs.to_pandas(nullable=True) + else: + ps = gs.to_pandas(nullable=True) + # Check type to avoid mixing Python bool and NumPy bool + assert isinstance(gs.hasnans, bool) + assert gs.hasnans == ps.hasnans @pytest.mark.parametrize( @@ -2949,3 +2952,16 @@ def test_index_getitem_from_int(idx): def test_index_getitem_from_nonint_raises(idx): with pytest.raises(ValueError): cudf.Index([1, 2])[idx] + + +@pytest.mark.parametrize( + "idx", + [ + cudf.RangeIndex(1), + cudf.DatetimeIndex(np.array([1, 2], dtype="datetime64[ns]")), + cudf.TimedeltaIndex(np.array([1, 2], dtype="timedelta64[ns]")), + ], +) +def test_index_to_pandas_nullable_notimplemented(idx): + with pytest.raises(NotImplementedError): + idx.to_pandas(nullable=True) diff --git a/python/cudf/cudf/tests/test_list.py b/python/cudf/cudf/tests/test_list.py index 8149188ba93..7ae7ae34b97 100644 --- a/python/cudf/cudf/tests/test_list.py +++ b/python/cudf/cudf/tests/test_list.py @@ -90,10 +90,8 @@ def test_leaves(data): def test_list_to_pandas_nullable_true(): df = cudf.DataFrame({"a": cudf.Series([[1, 2, 3]])}) - actual = df.to_pandas(nullable=True) - expected = pd.DataFrame({"a": pd.Series([[1, 2, 3]])}) - - assert_eq(actual, expected) + with pytest.raises(NotImplementedError): + df.to_pandas(nullable=True) def test_listdtype_hash(): diff --git a/python/cudf/cudf/tests/test_parquet.py b/python/cudf/cudf/tests/test_parquet.py index af4d0294293..b3a06dbd742 100644 --- a/python/cudf/cudf/tests/test_parquet.py +++ b/python/cudf/cudf/tests/test_parquet.py @@ -2593,7 +2593,8 @@ def test_parquet_writer_decimal(decimal_type, data): gdf.to_parquet(buff) got = pd.read_parquet(buff, use_nullable_dtypes=True) - assert_eq(gdf.to_pandas(nullable=True), got) + assert_eq(gdf["val"].to_pandas(nullable=True), got["val"]) + assert_eq(gdf["dec_val"].to_pandas(), got["dec_val"]) def test_parquet_writer_column_validation(): @@ -2627,6 +2628,11 @@ def test_parquet_writer_nulls_pandas_read(tmpdir, pdf): got = pd.read_parquet(fname) nullable = num_rows > 0 + if nullable: + gdf = gdf.drop(columns="col_datetime64[ms]") + gdf = gdf.drop(columns="col_datetime64[us]") + got = got.drop(columns="col_datetime64[ms]") + got = got.drop(columns="col_datetime64[us]") assert_eq(gdf.to_pandas(nullable=nullable), got) @@ -2671,7 +2677,7 @@ def test_parquet_reader_one_level_list(datadir): fname = datadir / "one_level_list.parquet" expect = pd.read_parquet(fname) - got = cudf.read_parquet(fname).to_pandas(nullable=True) + got = cudf.read_parquet(fname) assert_eq(expect, got) diff --git a/python/cudf/cudf/tests/test_udf_masked_ops.py b/python/cudf/cudf/tests/test_udf_masked_ops.py index 85531f8fae8..ad0c961a749 100644 --- a/python/cudf/cudf/tests/test_udf_masked_ops.py +++ b/python/cudf/cudf/tests/test_udf_masked_ops.py @@ -64,9 +64,9 @@ def substr(request): return request.param -def run_masked_udf_test(func, data, args=(), **kwargs): +def run_masked_udf_test(func, data, args=(), nullable=True, **kwargs): gdf = data - pdf = data.to_pandas(nullable=True) + pdf = data.to_pandas(nullable=nullable) expect = pdf.apply(func, args=args, axis=1) obtain = gdf.apply(func, args=args, axis=1) @@ -184,7 +184,7 @@ def func(row): ) gdf["a"] = gdf["a"].astype(dtype_l) gdf["b"] = gdf["b"].astype(dtype_r) - run_masked_udf_test(func, gdf, check_dtype=False) + run_masked_udf_test(func, gdf, nullable=False, check_dtype=False) @pytest.mark.parametrize("op", comparison_ops) diff --git a/python/dask_cudf/dask_cudf/core.py b/python/dask_cudf/dask_cudf/core.py index 17650c9b70d..ecdc566037d 100644 --- a/python/dask_cudf/dask_cudf/core.py +++ b/python/dask_cudf/dask_cudf/core.py @@ -56,10 +56,8 @@ def __repr__(self): @_dask_cudf_nvtx_annotate def to_dask_dataframe(self, **kwargs): """Create a dask.dataframe object from a dask_cudf object""" - nullable_pd_dtype = kwargs.get("nullable_pd_dtype", False) - return self.map_partitions( - M.to_pandas, nullable_pd_dtype=nullable_pd_dtype - ) + nullable = kwargs.get("nullable", False) + return self.map_partitions(M.to_pandas, nullable=nullable) concat = dd.concat diff --git a/python/dask_cudf/dask_cudf/tests/test_accessor.py b/python/dask_cudf/dask_cudf/tests/test_accessor.py index bea0cbb48ae..7ed5d797822 100644 --- a/python/dask_cudf/dask_cudf/tests/test_accessor.py +++ b/python/dask_cudf/dask_cudf/tests/test_accessor.py @@ -259,13 +259,13 @@ def test_categorical_categories(): {"a": ["a", "b", "c", "d", "e", "e", "a", "d"], "b": range(8)} ) df["a"] = df["a"].astype("category") - pdf = df.to_pandas(nullable_pd_dtype=False) + pdf = df.to_pandas(nullable=False) ddf = dgd.from_cudf(df, 2) dpdf = dd.from_pandas(pdf, 2) dd.assert_eq( - ddf.a.cat.categories.to_series().to_pandas(nullable_pd_dtype=False), + ddf.a.cat.categories.to_series().to_pandas(nullable=False), dpdf.a.cat.categories.to_series(), check_index=False, ) From 5e58e71836fd69ead04fbed5fdccb5e2e2c4d95c Mon Sep 17 00:00:00 2001 From: Matthew Roeschke <10647082+mroeschke@users.noreply.github.com> Date: Mon, 27 Nov 2023 16:25:19 -1000 Subject: [PATCH 44/44] REF: Make DataFrame.from_pandas process by column (#14483) Also encountered a bug where `cudf.Index.from_pandas` would return an `cudf.Index[int64]` from a `pandas.RangeIndex` Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Lawrence Mitchell (https://github.com/wence-) URL: https://github.com/rapidsai/cudf/pull/14483 --- python/cudf/cudf/core/_base_index.py | 16 ++++++-- python/cudf/cudf/core/dataframe.py | 55 +++++++++++----------------- python/cudf/cudf/tests/test_index.py | 7 ++++ 3 files changed, 40 insertions(+), 38 deletions(-) diff --git a/python/cudf/cudf/core/_base_index.py b/python/cudf/cudf/core/_base_index.py index 8387ef96dfa..fcfe8a21f05 100644 --- a/python/cudf/cudf/core/_base_index.py +++ b/python/cudf/cudf/core/_base_index.py @@ -1873,10 +1873,18 @@ def from_pandas(cls, index, nan_as_null=no_default): if not isinstance(index, pd.Index): raise TypeError("not a pandas.Index") - - ind = cudf.Index(column.as_column(index, nan_as_null=nan_as_null)) - ind.name = index.name - return ind + if isinstance(index, pd.RangeIndex): + return cudf.RangeIndex( + start=index.start, + stop=index.stop, + step=index.step, + name=index.name, + ) + else: + return cudf.Index( + column.as_column(index, nan_as_null=nan_as_null), + name=index.name, + ) @property def _constructor_expanddim(self): diff --git a/python/cudf/cudf/core/dataframe.py b/python/cudf/cudf/core/dataframe.py index 785f3d98712..4a31866a940 100644 --- a/python/cudf/cudf/core/dataframe.py +++ b/python/cudf/cudf/core/dataframe.py @@ -5245,30 +5245,20 @@ def from_pandas(cls, dataframe, nan_as_null=no_default): if not dataframe.columns.is_unique: raise ValueError("Duplicate column names are not allowed") - # Set columns - data = {} - for col_name, col_value in dataframe.items(): - # necessary because multi-index can return multiple - # columns for a single key - if len(col_value.shape) == 1: - data[col_name] = column.as_column( - col_value.array, nan_as_null=nan_as_null - ) - else: - vals = col_value.values.T - if vals.shape[0] == 1: - data[col_name] = column.as_column( - vals.flatten(), nan_as_null=nan_as_null - ) - else: - if isinstance(col_name, tuple): - col_name = str(col_name) - for idx in range(len(vals.shape)): - data[col_name] = column.as_column( - vals[idx], nan_as_null=nan_as_null - ) - - index = cudf.from_pandas(dataframe.index, nan_as_null=nan_as_null) + data = { + col_name: column.as_column( + col_value.array, nan_as_null=nan_as_null + ) + for col_name, col_value in dataframe.items() + } + if isinstance(dataframe.index, pd.MultiIndex): + index = cudf.MultiIndex.from_pandas( + dataframe.index, nan_as_null=nan_as_null + ) + else: + index = cudf.Index.from_pandas( + dataframe.index, nan_as_null=nan_as_null + ) df = cls._from_data(data, index) df._data._level_names = tuple(dataframe.columns.names) @@ -5279,13 +5269,14 @@ def from_pandas(cls, dataframe, nan_as_null=no_default): df.columns = dataframe.columns return df + elif hasattr(dataframe, "__dataframe__"): + # TODO: Probably should be handled in the constructor as + # this isn't pandas specific + return from_dataframe(dataframe, allow_copy=True) else: - try: - return from_dataframe(dataframe, allow_copy=True) - except Exception: - raise TypeError( - f"Could not construct DataFrame from {type(dataframe)}" - ) + raise TypeError( + f"Could not construct DataFrame from {type(dataframe)}" + ) @classmethod @_cudf_nvtx_annotate @@ -7915,10 +7906,6 @@ def from_pandas(obj, nan_as_null=no_default): return ret elif isinstance(obj, pd.MultiIndex): return MultiIndex.from_pandas(obj, nan_as_null=nan_as_null) - elif isinstance(obj, pd.RangeIndex): - return cudf.core.index.RangeIndex( - start=obj.start, stop=obj.stop, step=obj.step, name=obj.name - ) elif isinstance(obj, pd.Index): return cudf.Index.from_pandas(obj, nan_as_null=nan_as_null) elif isinstance(obj, pd.CategoricalDtype): diff --git a/python/cudf/cudf/tests/test_index.py b/python/cudf/cudf/tests/test_index.py index c393522c28b..7b859fefe9f 100644 --- a/python/cudf/cudf/tests/test_index.py +++ b/python/cudf/cudf/tests/test_index.py @@ -2954,6 +2954,13 @@ def test_index_getitem_from_nonint_raises(idx): cudf.Index([1, 2])[idx] +def test_from_pandas_rangeindex_return_rangeindex(): + pidx = pd.RangeIndex(start=3, stop=9, step=3, name="a") + result = cudf.Index.from_pandas(pidx) + expected = cudf.RangeIndex(start=3, stop=9, step=3, name="a") + assert_eq(result, expected, exact=True) + + @pytest.mark.parametrize( "idx", [