From 54f29b3e2b8f020a9d5e89f94002fdcc5273d124 Mon Sep 17 00:00:00 2001
From: ahmed
Date: Wed, 23 Feb 2022 00:07:35 -0500
Subject: [PATCH] RXMesh v0.2.1 (#13)
* Integrate Polyscope as a mesh viewer + demo
* Implement user-defined binary reduction operations on attributes
* Use `memcpy_async` for reading from global memory to shared memory
* API for allowing multiple queries in a single kernel
* Add validation function to make sure that the internal data structure is consistent (useful for dynamic changes support in RXMesh)
* Initial support for edge flip
* Update README with better documentation on how to use RXMesh for writing applications
---
.github/workflows/Ubuntu.yml | 11 +-
.github/workflows/Windows.yml | 9 +-
.gitignore | 4 +-
CMakeLists.txt | 29 +-
README.md | 253 +-
apps/Filtering/filtering.cu | 19 +-
apps/Filtering/filtering_rxmesh.cuh | 22 +-
apps/Geodesic/geodesic.cu | 20 +-
apps/Geodesic/geodesic_ptp_openmesh.h | 20 +-
apps/Geodesic/geodesic_ptp_rxmesh.h | 13 +-
apps/MCF/mcf.cu | 17 +-
apps/MCF/mcf_rxmesh.h | 11 +-
apps/VertexNormal/vertex_normal.cu | 5 +-
assets/polyscope_dragon.PNG | Bin 0 -> 1234302 bytes
include/rxmesh/attribute.h | 129 +-
include/rxmesh/context.h | 88 +-
include/rxmesh/iterator.cuh | 14 +-
include/rxmesh/kernels/attribute.cuh | 63 +-
include/rxmesh/kernels/collective.cuh | 3 +-
include/rxmesh/kernels/edge_flip.cuh | 190 +
include/rxmesh/kernels/loader.cuh | 297 +-
include/rxmesh/kernels/query_dispatcher.cuh | 45 +-
include/rxmesh/kernels/rxmesh_queries.cuh | 54 +-
include/rxmesh/kernels/update_dispatcher.cuh | 46 +
include/rxmesh/reduce_handle.h | 92 +-
include/rxmesh/rxmesh.cpp | 15 +-
include/rxmesh/rxmesh.h | 14 +-
include/rxmesh/rxmesh_dynamic.cu | 370 +
include/rxmesh/rxmesh_dynamic.h | 119 +
include/rxmesh/rxmesh_static.h | 354 +-
include/rxmesh/types.h | 10 +-
include/rxmesh/util/cuda_query.h | 2 +-
include/rxmesh/util/macros.h | 5 +-
include/rxmesh/util/util.h | 39 +
include/rxmesh/util/vector.h | 2 +-
input/diamond.obj | 6 +
input/dragon.obj | 30016 +++++++++++++++++
tests/CMakeLists.txt | 1 +
tests/Polyscope_test/CMakeLists.txt | 27 +
tests/Polyscope_test/test_polyscope.cu | 96 +
tests/RXMesh_test/CMakeLists.txt | 3 +
tests/RXMesh_test/edge_flip.cuh | 22 +
tests/RXMesh_test/rxmesh_test_main.cu | 4 +
tests/RXMesh_test/test_attribute.cuh | 85 +-
tests/RXMesh_test/test_edge_flip.h | 30 +
tests/RXMesh_test/test_for_each.h | 10 +-
tests/RXMesh_test/test_higher_queries.h | 2 +-
tests/RXMesh_test/test_queries.h | 8 +-
tests/RXMesh_test/test_util.cu | 27 +-
tests/RXMesh_test/test_validate.h | 13 +
50 files changed, 32232 insertions(+), 502 deletions(-)
create mode 100644 assets/polyscope_dragon.PNG
create mode 100644 include/rxmesh/kernels/edge_flip.cuh
create mode 100644 include/rxmesh/kernels/update_dispatcher.cuh
create mode 100644 include/rxmesh/rxmesh_dynamic.cu
create mode 100644 include/rxmesh/rxmesh_dynamic.h
create mode 100644 input/diamond.obj
create mode 100644 input/dragon.obj
create mode 100644 tests/Polyscope_test/CMakeLists.txt
create mode 100644 tests/Polyscope_test/test_polyscope.cu
create mode 100644 tests/RXMesh_test/edge_flip.cuh
create mode 100644 tests/RXMesh_test/test_edge_flip.h
create mode 100644 tests/RXMesh_test/test_validate.h
diff --git a/.github/workflows/Ubuntu.yml b/.github/workflows/Ubuntu.yml
index c618db96..68b51fd4 100644
--- a/.github/workflows/Ubuntu.yml
+++ b/.github/workflows/Ubuntu.yml
@@ -8,7 +8,9 @@ jobs:
id: cuda-toolkit
with:
cuda: '11.2.2'
- linux-local-args: '["--toolkit"]'
+ linux-local-args: '["--toolkit"]'
+ - run: sudo apt-get update
+ - run: sudo apt-get install -y xorg-dev libglu1-mesa-dev freeglut3-dev mesa-common-dev
- run: nvcc -V
- name: Checkout
uses: actions/checkout@v2
@@ -19,12 +21,7 @@ jobs:
run: cmake ../
- name: Run make
working-directory: ${{github.workspace}}/build
- run: |
- make RXMesh_test -j 99
- make Geodesic -j 99
- make MCF -j 99
- make VertexNormal -j 99
- # make Filtering -j 99
+ run: make -j 4
#- name: Run Test
# working-directory: ${{github.workspace}}/build
# run: ctest --no-compress-output -T Test -C Release --output-on-failure
diff --git a/.github/workflows/Windows.yml b/.github/workflows/Windows.yml
index 575fbc68..08ad1b7a 100644
--- a/.github/workflows/Windows.yml
+++ b/.github/workflows/Windows.yml
@@ -2,7 +2,7 @@ name: Windows
on: [push, pull_request, workflow_dispatch]
jobs:
WindowsRun:
- runs-on: windows-latest
+ runs-on: windows-2019
steps:
- uses: Jimver/cuda-toolkit@v0.2.4
id: cuda-toolkit
@@ -18,12 +18,7 @@ jobs:
working-directory: ${{github.workspace}}/build
run: cmake ../
- name: Run VS
- run: |
- cmake --build ${{github.workspace}}/build --target RXMesh_test --config Release -j 99
- cmake --build ${{github.workspace}}/build --target Geodesic --config Release -j 99
- cmake --build ${{github.workspace}}/build --target MCF --config Release -j 99
- cmake --build ${{github.workspace}}/build --target VertexNormal --config Release -j 99
- # cmake --build ${{github.workspace}}/build --target Filtering --config Release -j 99
+ run: cmake --build ${{github.workspace}}/build --clean-first --config Release -j 4
#- name: Run Test
# working-directory: ${{github.workspace}}/build
# run: ctest --no-compress-output -T Test -C Release --output-on-failure
diff --git a/.gitignore b/.gitignore
index 6b9aad91..ba62169c 100644
--- a/.gitignore
+++ b/.gitignore
@@ -1,5 +1,7 @@
output/
-input/
+input/*
+!input/dragon.obj
+!input/diamond.obj
build/
include/rxmesh/util/git_sha1.cpp
.vscode/
\ No newline at end of file
diff --git a/CMakeLists.txt b/CMakeLists.txt
index ea45e707..d8292893 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -5,9 +5,17 @@ if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.18)
endif()
project(RXMesh
- VERSION 0.2.0
+ VERSION 0.2.1
LANGUAGES C CXX CUDA)
+if (WIN32)
+ set(USE_POLYSCOPE "ON" CACHE BOOL "Enable Ployscope for visualization")
+ message(STATUS "Polyscope is enabled")
+else()
+ set(USE_POLYSCOPE "OFF" CACHE BOOL "Enable Ployscope for visualization")
+ message(STATUS "Polyscope is disabled")
+endif()
+
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
@@ -49,6 +57,15 @@ FetchContent_Declare(spdlog
)
FetchContent_Populate(spdlog)
+# polyscope
+if(USE_POLYSCOPE)
+FetchContent_Declare(polyscope
+ GIT_REPOSITORY https://github.com/Ahdhn/polyscope.git
+ GIT_TAG glm_patch #v1.3.0 with patch for glm
+)
+FetchContent_MakeAvailable(polyscope)
+endif()
+
# Auto-detect GPU architecture, sets ${CUDA_ARCHS}
include("cmake/AutoDetectCudaArch.cmake")
@@ -73,6 +90,9 @@ target_compile_definitions(RXMesh_header_lib
INTERFACE INPUT_DIR=${CMAKE_CURRENT_SOURCE_DIR}/input/
INTERFACE OUTPUT_DIR=${CMAKE_CURRENT_SOURCE_DIR}/output/
)
+if (USE_POLYSCOPE)
+ target_compile_definitions(RXMesh_header_lib INTERFACE USE_POLYSCOPE)
+endif()
target_include_directories( RXMesh_header_lib
INTERFACE "include"
INTERFACE "${rapidjson_SOURCE_DIR}/include"
@@ -101,7 +121,8 @@ set(cuda_flags
-use_fast_math
$<$:-O3>
--expt-relaxed-constexpr
- #-Xptxas -warn-spills -res-usage
+ -Xptxas -warn-spills -res-usage
+ --ptxas-options=-v
#-G
)
@@ -114,6 +135,10 @@ target_compile_options(developer_flags INTERFACE
target_link_libraries(RXMesh_header_lib INTERFACE $)
+if (USE_POLYSCOPE)
+ target_link_libraries(RXMesh_header_lib INTERFACE polyscope)
+endif()
+
#OpenMP
find_package(OpenMP)
if(OpenMP_CXX_FOUND)
diff --git a/README.md b/README.md
index 4afaa1e9..9afe82f7 100644
--- a/README.md
+++ b/README.md
@@ -4,17 +4,28 @@
+## **Contents**
+- [**About**](#about)
+- [**Compilation**](#compilation)
+ * [**Dependencies**](#dependencies)
+- [**Organization**](#organization)
+- [**Programming Model**](#programming-model)
+ * [**Structures**](#structures)
+ * [**Computation**](#computation)
+ * [**Viewer**](#viewer)
+- [**Replicability**](#replicability)
+- [**Bibtex**](#bibtex)
+
## **About**
-RXMesh is a surface triangle mesh data structure and programming model for processing static meshes on the GPU. RXMesh aims at provides a high-performance, generic, and compact data structure that can handle meshes regardless of their quality (e.g., non-manifold). The programming model helps to hide the complexity of the data structure and provides an intuitive access model for different use cases. For more details, please check out our paper:
+RXMesh is a surface triangle mesh data structure and programming model for processing static meshes on the GPU. RXMesh aims at provides a high-performance, generic, and compact data structure that can handle meshes regardless of their quality (e.g., non-manifold). The programming model helps to hide the complexity of the data structure and provides an intuitive access model for different use cases. For more details, please check out our paper and GTC talk:
-*[RXMesh: A GPU Mesh Data Structure](https://escholarship.org/uc/item/8r5848vp)*
+- *[RXMesh: A GPU Mesh Data Structure](https://escholarship.org/uc/item/8r5848vp)*
*[Ahmed H. Mahmoud](https://www.ece.ucdavis.edu/~ahdhn/), [Serban D. Porumbescu](https://web.cs.ucdavis.edu/~porumbes/), and [John D. Owens](https://www.ece.ucdavis.edu/~jowens/)*
*[ACM Transaction on Graphics](https://dl.acm.org/doi/abs/10.1145/3450626.3459748) (Proceedings of SIGGRAPH 2021)*
-This repository provides 1) source code to reproduce the results presented in the paper (git tag [`v0.1.0`](https://github.com/owensgroup/RXMesh/tree/v0.1.0)) and 2) ongoing development of RXMesh. For 1), all input models used in the paper can be found [here](https://ucdavis365-my.sharepoint.com/:f:/g/personal/ahmahmoud_ucdavis_edu/En-vEpIdSGBHqvCIa-MVXRQBg5g7GfM3P3RwZBHL4Hby3w?e=2EVnJd). Models were collected from [Thingi10K](https://ten-thousand-models.appspot.com/) and [Smithsonian 3D](https://3d.si.edu/explore) repository.
+- *[RXMesh: A High-performance Mesh Data Structure and Programming Model on the GPU [S41051]](https://www.nvidia.com/gtc/session-catalog/?tab.scheduledorondemand=1583520458947001NJiE&search=rxmesh#/session/1633891051385001Q9SE)—NVIDIA GTC 2022*
-## **A Quick Glance**
-RXMesh is a CUDA/C++ header-only library. All unit tests are under `tests/` folder. This includes the unit test for some basic functionalities along with the unit test for the query operations. All applications are under `apps/` folder.
+This repository provides 1) source code to reproduce the results presented in the paper (git tag [`v0.1.0`](https://github.com/owensgroup/RXMesh/tree/v0.1.0)) and 2) ongoing development of RXMesh. For 1), all input models used in the paper can be found [here](https://ucdavis365-my.sharepoint.com/:f:/g/personal/ahmahmoud_ucdavis_edu/En-vEpIdSGBHqvCIa-MVXRQBg5g7GfM3P3RwZBHL4Hby3w?e=2EVnJd). Models were collected from [Thingi10K](https://ten-thousand-models.appspot.com/) and [Smithsonian 3D](https://3d.si.edu/explore) repository.
## **Compilation**
The code can be compiled on Ubuntu (GCC 9) and Windows (Visual Studio 2019) providing that CUDA (>=11.1.0) is installed. To run the executable(s), an NVIDIA GPU should be installed on the machine.
@@ -36,8 +47,234 @@ All the dependencies are installed automatically! To compile the code:
```
Depending on the system, this will generate either a `.sln` project on Windows or a `make` file for a Linux system.
+## **Organization**
+RXMesh is a CUDA/C++ header-only library. All unit tests are under the `tests/` folder. This includes the unit test for some basic functionalities along with the unit test for the query operations. All applications are under the `apps/` folder.
+
+## **Programming Model**
+The goal of defining a programming model is to make it easy to write applications using RXMesh without getting into the nuances of the data structure. Applications written using RXMesh are composed of one or more of the high-level building blocks defined under [**Computation**](#computation). To use these building blocks, the user would have to interact with data structures specific to RXMesh discussed under [**Structures**](#structures). Finally, RXMesh integrates [Polyscope](https://polyscope.run) as a mesh [**Viewer**](#viewer) which the user can use to render their final results or for debugging purposes.
+
+### **Structures**
+- **Attributes** are the metadata (geometry information) attached to vertices, edges, or faces. Allocation of the attributes is per-patch basis and managed internally by RXMesh. The allocation could be done on the host, device, or both. Allocating attributes on the host is only beneficial for I/O operations or initializing attributes and then eventually moving them to the device.
+ - Example: allocation
+ ```c++
+ RXMeshStatic rx("input.obj");
+ auto vertex_color =
+ rx.add_vertex_attribute("vColor", //Unique name
+ 3, //Number of attribute per vertex
+ DEVICE, //Allocation place
+ SoA); //Memory layout (SoA vs. AoS)
+
+ ```
+ - Example: reading from `std::vector`
+ ```c++
+ RXMeshStatic rx("input.obj");
+ std::vector> face_color_vector;
+ //....
+
+ auto face_color =
+ rx.add_face_attribute(face_color_vector,//Input attribute where number of attributes per face is inferred
+ "fColor", //Unique name
+ SoA); //Memory layout (SoA vs. AoS)
+ ```
+ - Example: move, reset, and copy
+ ```c++
+ //By default, attributes are allocated on both host and device
+ auto edge_attr = rx.add_edge_attribute("eAttr", 1);
+ //Initialize edge_attr on the host
+ // .....
+
+ //Move attributes from host to device
+ edge_attr.move(HOST, DEVICE);
+
+ //Reset all entries to zero
+ edge_attr.reset(0, DEVICE);
+
+ auto edge_attr_1 = rx.add_edge_attribute("eAttr1", 1);
+
+ //Copy from another attribute.
+ //Here, what is on the host sde of edge_attr will be copied into the device side of edge_attr_1
+ edge_attr_1.copy_from(edge_attr, HOST, DEVICE);
+ ```
+
+- **Handles** are the unique identifiers for vertices, edges, and faces. They are usually internally populated by RXMesh (by concatenating the patch ID and mesh element index within the patch). Handles can be used to access attributes, `for_each` operations, and query operations.
+
+ - Example: Setting vertex attribute using vertex handle
+ ```c++
+ auto vertex_color = ...
+ VertexHandle vh;
+ //...
+
+ vertex_color(vh, 0) = 0.9;
+ vertex_color(vh, 1) = 0.5;
+ vertex_color(vh, 2) = 0.6;
+ ```
+
+- **Iterators** are used during query operations to iterate over the output of the query operation. The type of iterator defines the type of mesh element iterated on e.g., `VertexIterator` iterates over vertices which is the output of `VV`, `EV`, or `FV` query operations. Since query operations are only supported on the device, iterators can be only used inside the kernel. Iterators are usually populated internally.
+
+ - Example: Iterating over faces
+ ```c++
+ FaceIterator f_iter;
+ //...
+
+ for (uint32_t f = 0; f < f_iter.size(); ++f) {
+ FaceHandle fh = f_iter[f];
+ //do something with fh ....
+ }
+ ```
+
+
+### **Computation**
+- **`for_each`** runs a computation over all vertices, edges, or faces _without_ requiring information from neighbor mesh elements. The computation run on each mesh element is defined as a lambda function that takes a handle as an input. The lambda function could run either on the host, device, or both. On the host, we parallelize the computation using OpenMP. Care must be taken for lambda function on the device since it needs to be annotated using `__device__` and it can only capture by value. More about lambda function in CUDA can be found [here](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda)
+ - Example: using `for_each` to initialize attributes
+ ```cpp
+ RXMeshStatic rx("input.obj");
+ auto vertex_pos = rx.get_input_vertex_coordinates(); //vertex position
+ auto vertex_color = rx.add_vertex_attribute("vColor", 3, DEVICE); //vertex color
+
+ //This function will be executed on the device
+ rx.for_each_vertex(
+ DEVICE,
+ [vertex_color, vertex_pos] __device__(const VertexHandle vh) {
+ vertex_color(vh, 0) = 0.9;
+ vertex_color(vh, 1) = vertex_pos(vh, 1);
+ vertex_color(vh, 2) = 0.9;
+ });
+ ```
+
+- **Queries** operations supported by RXMesh with description are listed below
+
+ | Query | Description |
+ |-------|:-------------------------------------------|
+ | `VV` | For vertex V, return its adjacent vertices |
+ | `VE` | For vertex V, return its incident edges |
+ | `VF` | For vertex V, return its incident faces |
+ | `EV` | For edge E, return its incident vertices |
+ | `EF` | For edge E, return its incident faces |
+ | `FV` | For face F, return its incident vertices |
+ | `FE` | For face F, return its incident edges |
+ | `FF` | For face F, return its adjacent faces |
+
+ Queries are only supported on the device. RXMesh API for queries takes a lambda function along with the type of query. The lambda function defines the computation that will be run on the query output.
+
+ - Example: [vertex normal computation](./apps/VertexNormal/vertex_normal_kernel.cuh)
+ ```cpp
+ template
+ __global__ void vertex_normal (Context context){
+ auto compute_vn = [&](FaceHandle face_id, VertexIterator& fv) {
+ //This thread is assigned to face_id
+
+ // get the face's three vertices coordinates
+ Vector<3, T> c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2));
+ Vector<3, T> c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2));
+ Vector<3, T> c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2));
+
+ //compute face normal
+ Vector<3, T> n = cross(c1 - c0, c2 - c0);
+
+ // add the face's normal to its vertices
+ for (uint32_t v = 0; v < 3; ++v) // for every vertex in this face
+ for (uint32_t i = 0; i < 3; ++i) // for the vertex 3 coordinates
+ atomicAdd(&normals(fv[v], i), n[i]);
+ };
+
+ //Query dispatcher must be called by all threads in the block.
+ //Dispatcher will first perform the query, store the results in shared memory, then
+ //run the user-defined computation i.e., compute_vn
+ query_block_dispatcher(context, compute_vn);
+ }
+ ```
+ To save computation, `query_block_dispatcher` could be run on a subset of the input mesh element i.e., _active set_. The user can define the active set using a lambda function that returns true if the input mesh element is in the active set.
+
+ - Example: defining active set
+ ```cpp
+ template
+ __global__ void active_set_query (Context context){
+ auto active_set = [&](FaceHandle face_id) -> bool{
+ // ....
+ };
+
+ auto computation = [&](FaceHandle face_id, VertexIterator& fv) {
+ // ....
+ };
+
+ query_block_dispatcher(context, computation, active_set);
+ }
+ ```
+
+- **Reduction** operations apply a binary associative operation on the input attributes. RXMesh provides dot products between two attributes (of the same type), L2 norm of an input attribute, and user-defined reduction operation on an input attribute. For user-defined reduction operation, the user needs to pass a binary reduction functor with member `__device__ T operator()(const T &a, const T &b)` or use on of [CUB's thread operators](https://github.com/NVIDIA/cub/blob/main/cub/thread/thread_operators.cuh) e.g., `cub::Max()`. Reduction operations require allocation of temporary buffers which we abstract away using `ReduceHandle`.
+
+ - Example: dot product, L2 norm, user-defined reduction
+ ```cpp
+ RXMeshStatic rx("input.obj");
+ auto vertex_attr1 = rx.add_vertex_attribute("v_attr1", 3, DEVICE);
+ auto vertex_attr2 = rx.add_vertex_attribute("v_attr2", 3, DEVICE);
+
+ // Populate vertex_attr1 and vertex_attr2
+ //....
+
+ //Reduction handle
+ ReduceHandle reduce(v1_attr);
+
+ //Dot product between two attributes. Results are returned on the host
+ float dot_product = reduce.dot(v1_attr, v2_attr);
+
+ cudaStream_t stream;
+ //init stream
+ //...
+
+ //Reduction operation could be performed on specific attribute and using specific stream
+ float l2_norm = reduce.norm2(v1_attr, //input attribute
+ 1, //attribute ID. If not specified, reduction is run on all attributes
+ stream); //stream used for reduction.
+
+
+ //User-defined reduction operation
+ float l2_norm = reduce.reduce(v1_attr, //input attribute
+ cub::Max(), //binary reduction functor
+ std::numeric_limits::lowest()); //initial value
+ ```
+
+### **Viewer**
+Starting v0.2.1, RXMesh integrates [Polyscope](https://polyscope.run) as a mesh viewer. To use it, make sure to turn on the CMake parameter `USE_POLYSCOPE` i.e.,
+
+```
+> cd build
+> cmake -DUSE_POLYSCOPE=True ../
+```
+By default, the parameter is set to True on Windows and False on Linux machines. RXMesh implements the necessary functionalities to pass attributes to Polyscope—thanks to its [data adaptors](https://polyscope.run/data_adaptors/). However, this needs attributes to be moved to the host first before passing it to Polyscope. For more information about Polyscope's different visualization options, please checkout Polyscope's [Surface Mesh documentation](https://polyscope.run/structures/surface_mesh/basics/).
+
+ - Example: [render vertex color](./tests/Polyscope_test/test_polyscope.cu)
+ ```cpp
+ //initialize polyscope
+ polyscope::init();
+
+ RXMeshStatic rx("dragon.obj");
+
+ //vertex color attribute
+ auto vertex_color = rx.add_vertex_attribute("vColor", 3);
+
+ //Populate vertex color on the device
+ //....
+
+ //Move vertex color to the host
+ vertex_color.move(DEVICE, HOST);
+
+ //polyscope instance associated with rx
+ auto polyscope_mesh = rx.get_polyscope_mesh();
+
+ //pass vertex color to polyscope
+ polyscope_mesh->addVertexColorQuantity("vColor", vertex_color);
+
+ //render
+ polyscope::show();
+ ```
+
+
+
+
+
## **Replicability**
-This repo was awarded the [replicability stamp](http://www.replicabilitystamp.org#https-github-com-owensgroup-rxmesh) by the Graphics Replicability Stamp Initiative (GRSI).
+This repo was awarded the [replicability stamp](http://www.replicabilitystamp.org#https-github-com-owensgroup-rxmesh) by the Graphics Replicability Stamp Initiative (GRSI) :tada:
The scripts used to generate the data shown in the paper can be found under
* [Figure 6](https://github.com/owensgroup/RXMesh/blob/main/tests/RXMesh_test/benchmark.sh)
@@ -46,9 +283,7 @@ The scripts used to generate the data shown in the paper can be found under
* [Figure 8 (c)](https://github.com/owensgroup/RXMesh/blob/main/apps/Filtering/benchmark.sh)
* [Figure 8 (d)](https://github.com/owensgroup/RXMesh/blob/main/apps/VertexNormal/benchmark.sh)
-Each script should be run from the script's containing directory after compiling the code in `build/` directory. The only input parameter needed is the path to the input OBJ files. The resulting JSON files will be written to `output/` directory.
-
-
+Each script should be run from the script's containing directory after compiling the code in the `build/` directory. The only input parameter needed is the path to the input OBJ files. The resulting JSON files will be written to the `output/` directory.
## **Bibtex**
```
diff --git a/apps/Filtering/filtering.cu b/apps/Filtering/filtering.cu
index 3a85ee2a..b045b573 100644
--- a/apps/Filtering/filtering.cu
+++ b/apps/Filtering/filtering.cu
@@ -33,27 +33,22 @@ TEST(App, Filtering)
cuda_query(Arg.device_id);
- // Load mesh
- std::vector> Faces;
- std::vector> Verts;
- ASSERT_TRUE(import_obj(Arg.obj_file_name, Verts, Faces));
-
-
TriMesh input_mesh;
ASSERT_TRUE(OpenMesh::IO::read_mesh(input_mesh, Arg.obj_file_name));
- ASSERT_EQ(input_mesh.n_vertices(), Verts.size());
-
// OpenMesh Impl
- std::vector> ground_truth(Verts);
+ std::vector> ground_truth(input_mesh.n_vertices());
+ for (auto& g : ground_truth) {
+ g.resize(3);
+ }
size_t max_neighbour_size = 0;
- filtering_openmesh(
+ filtering_openmesh(
omp_get_max_threads(), input_mesh, ground_truth, max_neighbour_size);
// RXMesh Impl
- filtering_rxmesh(Faces, Verts, ground_truth, max_neighbour_size);
-
+ filtering_rxmesh(
+ Arg.obj_file_name, ground_truth, max_neighbour_size);
}
int main(int argc, char** argv)
diff --git a/apps/Filtering/filtering_rxmesh.cuh b/apps/Filtering/filtering_rxmesh.cuh
index ff528279..ac670c76 100644
--- a/apps/Filtering/filtering_rxmesh.cuh
+++ b/apps/Filtering/filtering_rxmesh.cuh
@@ -11,10 +11,9 @@
* filtering_rxmesh()
*/
template
-void filtering_rxmesh(std::vector>& Faces,
- const std::vector>& Verts,
- const std::vector>& ground_truth,
- const size_t max_neighbour_size)
+void filtering_rxmesh(const std::string file_path,
+ const std::vector>& ground_truth,
+ const size_t max_neighbour_size)
{
using namespace rxmesh;
@@ -25,7 +24,7 @@ void filtering_rxmesh(std::vector>& Faces,
"greater than maxVVSize. Should increase maxVVSize to "
<< max_neighbour_size << " to avoid illegal memory access";
- RXMeshStatic rxmesh(Faces, false);
+ RXMeshStatic rxmesh(file_path, false);
// Report
Report report("Filtering_RXMesh");
@@ -38,7 +37,7 @@ void filtering_rxmesh(std::vector>& Faces,
// input coords
- auto coords = rxmesh.add_vertex_attribute(Verts, "coords");
+ auto coords = rxmesh.get_input_vertex_coordinates();
// Vertex normals (only on device)
auto vertex_normal = rxmesh.add_vertex_attribute("vn", 3, DEVICE);
@@ -53,15 +52,16 @@ void filtering_rxmesh(std::vector>& Faces,
// vertex normal launch box
constexpr uint32_t vn_block_threads = 256;
LaunchBox vn_launch_box;
- rxmesh.prepare_launch_box(rxmesh::Op::FV,
- vn_launch_box,
- (void*)compute_vertex_normal);
+ rxmesh.prepare_launch_box(
+ {rxmesh::Op::FV},
+ vn_launch_box,
+ (void*)compute_vertex_normal);
// filter launch box
- constexpr uint32_t filter_block_threads = 512;
+ constexpr uint32_t filter_block_threads = 256;
LaunchBox filter_launch_box;
rxmesh.prepare_launch_box(
- rxmesh::Op::VV,
+ {rxmesh::Op::VV},
filter_launch_box,
(void*)bilateral_filtering);
diff --git a/apps/Geodesic/geodesic.cu b/apps/Geodesic/geodesic.cu
index c11aa7af..8b18b68e 100644
--- a/apps/Geodesic/geodesic.cu
+++ b/apps/Geodesic/geodesic.cu
@@ -36,13 +36,7 @@ TEST(App, Geodesic)
// Select device
cuda_query(Arg.device_id);
-
- // Load mesh
- std::vector> Verts;
- std::vector> Faces;
- ASSERT_TRUE(import_obj(Arg.obj_file_name, Verts, Faces));
-
- RXMeshStatic rxmesh(Faces, false);
+ RXMeshStatic rxmesh(Arg.obj_file_name, false);
ASSERT_TRUE(rxmesh.is_closed())
<< "Geodesic only works on watertight/closed manifold mesh without "
"boundaries";
@@ -55,8 +49,8 @@ TEST(App, Geodesic)
std::vector h_seeds(Arg.num_seeds);
std::random_device dev;
std::mt19937 rng(dev());
- std::uniform_int_distribution dist(0,
- Verts.size());
+ std::uniform_int_distribution dist(
+ 0, rxmesh.get_num_vertices());
for (auto& s : h_seeds) {
s = dist(rng);
// s = 0;
@@ -68,15 +62,13 @@ TEST(App, Geodesic)
// sorted_index and limit. We keep it for RXMesh because it is
// used to quickly determine whether or not a vertex is within
// the "update band".
- std::vector toplesets(Verts.size(), 1u);
+ std::vector toplesets(rxmesh.get_num_vertices(), 1u);
std::vector sorted_index;
std::vector limits;
- geodesic_ptp_openmesh(
- Faces, Verts, h_seeds, sorted_index, limits, toplesets);
+ geodesic_ptp_openmesh(h_seeds, sorted_index, limits, toplesets);
// RXMesh Impl
- geodesic_rxmesh(
- rxmesh, Faces, Verts, h_seeds, sorted_index, limits, toplesets);
+ geodesic_rxmesh(rxmesh, h_seeds, sorted_index, limits, toplesets);
}
int main(int argc, char** argv)
diff --git a/apps/Geodesic/geodesic_ptp_openmesh.h b/apps/Geodesic/geodesic_ptp_openmesh.h
index e8ad2569..b8f45697 100644
--- a/apps/Geodesic/geodesic_ptp_openmesh.h
+++ b/apps/Geodesic/geodesic_ptp_openmesh.h
@@ -256,12 +256,10 @@ inline float toplesets_propagation(TriMesh& mesh,
}
template
-void geodesic_ptp_openmesh(const std::vector>& Faces,
- const std::vector>& Verts,
- const std::vector& h_seeds,
- std::vector& sorted_index,
- std::vector& limits,
- std::vector& toplesets)
+void geodesic_ptp_openmesh(const std::vector& h_seeds,
+ std::vector& sorted_index,
+ std::vector& limits,
+ std::vector& toplesets)
{
TriMesh input_mesh;
ASSERT_TRUE(OpenMesh::IO::read_mesh(input_mesh, Arg.obj_file_name));
@@ -275,9 +273,6 @@ void geodesic_ptp_openmesh(const std::vector>& Faces,
std::string method = "OpenMeshSingleCore";
report.add_member("method", method);
- ASSERT_TRUE(Faces.size() == input_mesh.n_faces());
- ASSERT_TRUE(Verts.size() == input_mesh.n_vertices());
-
std::vector geo_distance(input_mesh.n_vertices(),
std::numeric_limits::infinity());
@@ -303,13 +298,6 @@ void geodesic_ptp_openmesh(const std::vector>& Faces,
input_mesh, h_seeds, limits, sorted_index, geo_distance, iter);
RXMESH_TRACE("geodesic_ptp_openmesh() took {} (ms)", processing_time);
- // export_attribute_VTK("geo_openmesh.vtk",
- // Faces,
- // Verts,
- // false,
- // geo_distance.data(),
- // geo_distance.data());
-
// Finalize report
report.add_member("num_iter_taken", iter);
rxmesh::TestData td;
diff --git a/apps/Geodesic/geodesic_ptp_rxmesh.h b/apps/Geodesic/geodesic_ptp_rxmesh.h
index 11a2c819..a05fc9a5 100644
--- a/apps/Geodesic/geodesic_ptp_rxmesh.h
+++ b/apps/Geodesic/geodesic_ptp_rxmesh.h
@@ -7,10 +7,8 @@
constexpr float EPS = 10e-6;
template
-inline void geodesic_rxmesh(rxmesh::RXMeshStatic& rxmesh,
- const std::vector>& Faces,
- const std::vector>& Verts,
- const std::vector& h_seeds,
+inline void geodesic_rxmesh(rxmesh::RXMeshStatic& rxmesh,
+ const std::vector& h_seeds,
const std::vector& h_sorted_index,
const std::vector& h_limits,
const std::vector& toplesets)
@@ -28,7 +26,7 @@ inline void geodesic_rxmesh(rxmesh::RXMeshStatic& rxmesh,
report.add_member("method", std::string("RXMesh"));
// input coords
- auto input_coord = rxmesh.add_vertex_attribute(Verts, "coord");
+ auto input_coord = rxmesh.get_input_vertex_coordinates();
// toplesets
auto d_toplesets = rxmesh.add_vertex_attribute(toplesets, "topleset");
@@ -36,7 +34,7 @@ inline void geodesic_rxmesh(rxmesh::RXMeshStatic& rxmesh,
// RXMesh launch box
LaunchBox launch_box;
- rxmesh.prepare_launch_box(rxmesh::Op::VV,
+ rxmesh.prepare_launch_box({rxmesh::Op::VV},
launch_box,
(void*)relax_ptp_rxmesh,
true);
@@ -134,8 +132,7 @@ inline void geodesic_rxmesh(rxmesh::RXMeshStatic& rxmesh,
// uint32_t v_id = rxmesh.map_to_global(vh);
// geo[v_id] = (*rxmesh_geo)(vh);
//});
- // export_attribute_VTK(
- // "geo_rxmesh.vtk", Faces, Verts, false, geo.data(), geo.data());
+
GPU_FREE(d_error);
diff --git a/apps/MCF/mcf.cu b/apps/MCF/mcf.cu
index abfec068..34f2d8b7 100644
--- a/apps/MCF/mcf.cu
+++ b/apps/MCF/mcf.cu
@@ -38,26 +38,21 @@ TEST(App, MCF)
// Select device
cuda_query(Arg.device_id);
-
- // Load mesh
- std::vector> Verts;
- std::vector> Faces;
-
- ASSERT_TRUE(import_obj(Arg.obj_file_name, Verts, Faces));
-
-
- RXMeshStatic rxmesh(Faces, false);
+ RXMeshStatic rxmesh(Arg.obj_file_name, false);
TriMesh input_mesh;
ASSERT_TRUE(OpenMesh::IO::read_mesh(input_mesh, Arg.obj_file_name));
// OpenMesh Impl
- std::vector> ground_truth(Verts);
+ std::vector> ground_truth(rxmesh.get_num_vertices());
+ for (auto& g : ground_truth) {
+ g.resize(3);
+ }
mcf_openmesh(omp_get_max_threads(), input_mesh, ground_truth);
// RXMesh Impl
- mcf_rxmesh(rxmesh, Verts, ground_truth);
+ mcf_rxmesh(rxmesh, ground_truth);
}
int main(int argc, char** argv)
diff --git a/apps/MCF/mcf_rxmesh.h b/apps/MCF/mcf_rxmesh.h
index e3681c8b..6dcaf831 100644
--- a/apps/MCF/mcf_rxmesh.h
+++ b/apps/MCF/mcf_rxmesh.h
@@ -48,7 +48,6 @@ void init_PR(rxmesh::RXMeshStatic& rxmesh,
template
void mcf_rxmesh(rxmesh::RXMeshStatic& rxmesh,
- const std::vector>& Verts,
const std::vector>& ground_truth)
{
using namespace rxmesh;
@@ -71,8 +70,7 @@ void mcf_rxmesh(rxmesh::RXMeshStatic& rxmesh,
<< "mcf_rxmesh only takes watertight/closed mesh without boundaries";
// Different attributes used throughout the application
- auto input_coord =
- rxmesh.add_vertex_attribute(Verts, "coord", rxmesh::LOCATION_ALL);
+ auto input_coord = rxmesh.get_input_vertex_coordinates();
// S in CG
auto S =
@@ -95,18 +93,19 @@ void mcf_rxmesh(rxmesh::RXMeshStatic& rxmesh,
B->reset(0.0, rxmesh::DEVICE);
// X in CG (the output)
- auto X = rxmesh.add_vertex_attribute(Verts, "X", rxmesh::LOCATION_ALL);
+ auto X = rxmesh.add_vertex_attribute("X", 3, rxmesh::LOCATION_ALL);
+ X->copy_from(*input_coord, rxmesh::DEVICE, rxmesh::DEVICE);
ReduceHandle reduce_handle(*X);
// RXMesh launch box
LaunchBox launch_box_init_B;
LaunchBox launch_box_matvec;
- rxmesh.prepare_launch_box(rxmesh::Op::VV,
+ rxmesh.prepare_launch_box({rxmesh::Op::VV},
launch_box_init_B,
(void*)init_B,
true);
- rxmesh.prepare_launch_box(rxmesh::Op::VV,
+ rxmesh.prepare_launch_box({rxmesh::Op::VV},
launch_box_matvec,
(void*)rxmesh_matvec,
true);
diff --git a/apps/VertexNormal/vertex_normal.cu b/apps/VertexNormal/vertex_normal.cu
index 3fe2da5b..d8b57391 100644
--- a/apps/VertexNormal/vertex_normal.cu
+++ b/apps/VertexNormal/vertex_normal.cu
@@ -50,8 +50,9 @@ void vertex_normal_rxmesh(rxmesh::RXMeshStatic& rxmesh,
// launch box
LaunchBox launch_box;
- rxmesh.prepare_launch_box(
- rxmesh::Op::FV, launch_box, (void*)compute_vertex_normal);
+ rxmesh.prepare_launch_box({rxmesh::Op::FV},
+ launch_box,
+ (void*)compute_vertex_normal);
TestData td;
diff --git a/assets/polyscope_dragon.PNG b/assets/polyscope_dragon.PNG
new file mode 100644
index 0000000000000000000000000000000000000000..a973ca06271492a74ca963afb6694c7bd33111d2
GIT binary patch
literal 1234302
zcmb@tWmH@1w>?Z-id&$#2QBVWtU!R^h2k#hp}4yhcL@|LPN_g~cZw#s6?X_yw7A1R
z=Wy=*-TQvJAKnic8GC1so&Bu6WUjg96ZJ+-0T+t`3k3xQR}l=-L_tA4ML~Jmg@J}V
zlSE?j8u|6aRZ~F*rEHk$H}U}0N?KJK1*I|p`_}w9@)*+ztnZ3~g4gx;=LyX5I|K#g
z;fo?jTFc91ul1c)rv6RaUBKPLSI%x{&tO9^82smTqM5WZu~G~%+ppG-Z0??(o)>D1
z-?Ho89yYie<%llt87Hy4?#=nF*PKv0{JfE=vHGx(pK7=1{Qlvm84no_T1dx#f87i6
ziGS7jS%gWmR#Irx18ePe_|
z68_jXFN{s%@;r<3?+s!Ww`&^oWxXoPjWhec1iJnX^mE*30d@C7R{Z~3B|tZ9X~@@c
zeYV=F{Ut)rb;SJB7lpc2#2K~#tUMbQU#-AnyZq`SCynj7mn%JN*=LK(htpP|ITvZD
zTw34`yq&V0aW7!t-zSmEtb$`1B^tpx;T!PiLAeqRfZ~$Pix{!h#1YJeqku@3CEA*(
zac;C2Z}#+GmXB+Gzv)@O@JXrJG`i8%&}zEzbHmv%_26tf4K|)bXNU5)qm?6utVZt)
zouv%1Z|>=%mge4q$I6`t^>A@Q=2PTC)^im#AS{ECSrH%M0)+!rqlQzY`r}5BqN>Uz
zSE6tq+=zBly06#60_K6}J+in=_j%T7Fhc8}4i>>92RO-lZr((4d?T=^{vHYd^
z=S
zvb#IP#5El1c6?7X>~&AeF^dtmY9~SSmldM!SbkjHS)l
zNd3+w*ZRqHN%}7fDQs+I<)8~+2+d=V9F4pNkcSZJ>z1V0_rCoNYhY|ur1eZrHJ3$@
zFc$sV{<;(YJSWC!c^}G~wpec7Fby^+oo37ls&A#O0r;rw{2Rfy2{Pr@)BPk*M+6|b
zjjn?|S);qnf3fB4S~9q?(Sx{EEFD-eHl78*P=zq(wZHZ5yy1FSE$`_@Tf
zoC%8L>u_gzD7%D*dRikciqbQyY@cUWCvYrD$5U2X(=2&kqWBKeek)f+rG|P!XaT+u
zwVw>SOWg~1lup~ne!&YQhVY}kFgEM>+&v(XRp{@TE{7Is=xxlR{n
zDUvh_w4t>ZE;HaAyQJda;NZ3(@Mdh+-h@7;u|TP0$iz|pDfVkDKdw8xx9_Uy)8s7Q
zAzo2^qzX|Y4(aN*O!QH99=iJrJ6~b_9wSqtZfvL3&l!MsNflO{GHUr>&agZDz<3H)
zVal-{q`q;bmpPovuIG7wr-HM@6Qfga?0v&rwo?yHu~pQj0=(nT+@QfMT4B(P#}GcA
z(b|UkBd#Jj_^H(68Dbnha?!|!t$djmep6StaqEX%y3-9;8-Z78_1TcDi0=uO566|G
z;PTFT^Plq&U-?RKS#qRI*S*R0F8$L4trtnKi1L;wOWl!t-;)4BAet`KqPZV3eI5h<
z%{|G|RjZfM#Se113{?~0Y)Xa_eG2Rea5ggGepDW2);;JOfd=}ZclUM{BS~_9m@L!6IxrTXsPz5
z^*fGz|MWrHBxFcQJW2ULmv_;XS!;zNo%cj!4$Tv65c@Wd|1
zHmpZb5IY?qh4gzInURGy(}8kK3E(pLd2SIa1HFCY4R=mKdp(H{Bc<9tY_`5^G+IP<
zz;>6Q&$FO0IQ}-^?x2(e+c|x6d$L-@^@PuUNYE~`I#8mNW3Rd?O45Pzlr1&Ev9bZZ
zk;KI~{0rEKIf@g_GCqh5hzQit
zI==AnPQ4;<^5W(S{co?v$Cl7+oNDsfuiu;wrRjsRTRE>z3C7OXrPnPbvDLGv=Thiq
z2(p~QZE6dB;=w`_9>$sgVV{Ag2JQ+v^%DN;Ng55%t21!V66@Vb(j5eVT}6Nn^#TDC
zw}7q)=oUi)cl%v|@l-_t*KFYBG5M=QTBskS1Gg$YLi$p5wl+zUSvvwD+UE=sN}A?3
zqQxHA{j?1HPZ2kbks%3@t0GS9E8t~up_h!goMSy{i$UpCf_+w*e5RgqxGuI4$YWZI
zv0)}Vj!I|X`*VKRobGAHu&98#d@5nPtyHg`Bbg5)
zFjpx4gEOL92k1~3W`G$r`wq;++hXY7rS7q$_#*`2MjuKxQo}an0u+CdAKSwIP2$uJ
zWe2PuQ}QBeNXwKQ$mM8MQF*vs0500ET1x#Q`djcn=F-W1>F`J}ZdebXH59M=oT=+M
zgJnNv{3e>|mT6{S-by7S4`(&Y$*pcX_GlkA@xt&$SLS*gU|(T}R{aAJ3{(6cr$Yka
zmhjtBnZt`vtjmTJ|w?;1DH0r^(%
z8<}F9YnB`-fKLs}YX*?X7}ga^gS^?T;AJz~II-=tyhEojK&tkm_qAQ*LNsFv1?k7M
zqCtk)?{mRMj?}5XWuR0~p_WjS??k`MaRcqKp$dGNQj|t`q#jgav4zJ6VWX7u0URlQ
z8L9eYlK}#y!|To>h@}5|S?u?M$VZ0U75Te`=}(9r-tVPjNsy4FwzE9egxrJ$L5gQz
zWdFn?6e*z!E?^fbHfiTMPYkdnl5hraKrTS~%8wrgDhsE3A2ajOt0N_y
z?aGRZmK^iPtSByS+{zEQ
zJQQG=_;RTCU7zYlaI;yk4WXv?Z}72FC*z2s$aVn6*22gH$d|n`zJ93$&R>|BC#FS3
zOeKoOB_E?-Srxo
zHG15|JbUzyTMm_fjKlY6923d>pw$^7Cltr40UwHqV
zpUW5NlIwVtbdz83%R{Xqhm|A3Zsq=xqv=z*=V-F9&yXVz7VXel%II$~rw$+UJ}bGJ
zBESeu!~oL&dVt*C5X3y=0oqe>K^A@{WL)FwUqf~$I8NIpzXe_s+&z1&EK=x1KFFM3
zu7d0zc*tVYl(q7n!k((FibHV(Fxp`-A0?r`E|3a+QRBG_4(+F)J>v;rF3DMQY79
zMUq!)?Z*u5ALj*9GA73lnt7o!@(=Woe|prhr~hk7fmGg30y^j=!u{e&a!-_bnJuVN
zycN3CWL|2R2aL=3rCvS{FOD#m`=gYF--O~FNoBCB;&!wXYT%Dh8BnHfs(KoBWs@ewtJj**q@%MKQYc6mn_f^xg`AVJF-tq$(F#y>f&
zwe3%R!yu$^_~F?!4a2?y{Qz4`so%ig*-W1O?|$0c`(5W_0pa!A^Q#c>VltjQclznf
zcG7lqgJ1Q`)%n#Pl}GrYTp=iYQqlHEZf)8t=aRVUm%a$G8`?$ifu)iYk!I-x!2UC+
zRM$SHhxgoPPqu-lC2TNrI2T#zBj~Ba>0o^5QK?apwXtEG%Xm>OjBbdO*Iml~V=nBG
z;mtokUn!A4plUMy-X_t&%K!x}bj!pl*>!*DHjj)_?aoVI-6orV&J@5Tvh9kY34M_v
zogqT>hj|LMOxk}GrTcjgy>WKPSRYa1w}VI*{pZ6V)1QKq_TLMgzgHCfXvV-K(sJ%YbL*w8SpNsdVa?I|jIu&yjQne4oSVzYe>D6Dps5}K72
zYp*Ew=Uu1pwO@>H6#jedrFm)3qS0nDH}YZv;=fL|I!1bF8N?MbX76$oh%l_P
zdzdVDCbSh(!THrHVK`H~v_%1$za`q*n?Xg6X#F6R2NV?&;nlAfoD@;u;5Qm%o6m>4
zCV$#CP`Kfyg_N6_!X;I3Gosg4EKt#(%a*DsaO)HEg_3XEEaLX)^uTNT<(gQjqY4T)
ze>@tY<9{20Wt1O-xMT8C662D2TR{L#9b|6COM=eP^*yo@{Px?b*|<5|G0tq78;XRi
zk7kAyJa#1mT%%+8B#@WNj%o}s3kqA(zo`CmcHX15q{eoNhK(jW2$;A!hPQY>ZA;Wy
z6Lq2@=*EaZLce2Vm%jSktoqVt?!1tl{*n3edh)-@!9$#QrXRuwVeU98eM^A)
zDIugSFPK}H2fuO0C`-+!I&6?3<9kl1Xoo8@e%YmUp#(`Etb_H_InvQ+9cejEjlf?_
zKGsuFdQny;wPQKR2y5wU(Rc7Yxu*&Q>Gx|awNbUsX^(!a;|^hf0bGKDhcYjO0ZeZ#
zT52%offp7fZcmTj^&s)~7=Im>n8o=j(WLGN!b`>HB~h~=Lwi^>rPJh1lg1{rWy7D+
zynttKod)y~9RFLIe-5M7)qG2XXq@uiFYrBxq7ZYOXV8Hxe6}52Yb=;;3H=8+^DsB6RShy=tVM#8?Y*XV(
zhBqZjVpL+veRcwdV=K`A5_@4QF$J7wr;NE$9#tI0JW4&n*a(&B(J=qegK~FhE}x;0
zWRV~Anhq5Ddfe|zYq$@R7WdtUo^AUTpI@GN2FfNV*tWd$SzQsF~89XvEqFTQKJl73LA~tCSAZ^@}oVi
z{T-o{spRF{%Ja)%8#{lCCst?c$oqDa{V{hRHHVq}Wv>gt-|ss{m+c)}xcn3zcSf6F
zp7ju23As3(J~8awP6)-6Kmvu!yU=-gvm{?4ieMP!yCqL)Df56o!0*8pP~zM$OhqDS
z(QtQ8p}!Uakbfa?BoP`Gx>}lNsD*%PJHWyPM(w}EzazjvjeG$bf6iwEa*sTVmkZN<
z{^?SczzNl9Q-ooFk8)yujzaB+DQzn;lPND|LEh$7U+b!$de}hxneIuGr
z%82z<=8*1zXAN|e{z4!Mt2Kr`jfd5XXW@$Gdv%#?$f8JW4JD>+8?KF=jfvLQY}JnV
zKj6sVq&Y07QEFWOzCNAObRu0WV&h_~oN8d|qQZSJ`=qI7p`oSZ5oNX~VC}Nuc+jF{
zL{O0imuY4wY>pa|w~!>!z3gA}%`^1#6j$+!RaG=hEKuv^Q28!zXmU<$GvaQUp?jV}*B`t|UNqUqHOyNi%6_h)xVuQEt}dMm*mq
zDQK6IFcNi;&;M2l4ohSQRnQe8D@7nl&uTURKriQLFCk>(xu($@p*15hatU}_Q*Eui
z#Jey08Q%Q0HO6`FC46*ebnFMJ(A3m6VhYp6Q*>c*5wJ@eLldx(&YBa8+m;zy!2RY0
zIBbdc9#e&MC{!FAzuv*Z1N1`8y5jm?xf6#r@2s^gIFle-!b8$nV8nL1_WX$823|8CcJ=GEq~_o9V^k%`VD;#^jURA9wxpq7?*Vz#ej
zl5tcX^*}`r0|HqIVV3Iok`5?kEKA+OQOm0`Ywz52(+YHqBN`P?czGd&RWen1ZkiTN
zi#SyEuYx9Ytlp!E;>a}Sv7G<8FnboraE*1{p}nG{TU}cHy(`!Q`}?3T{pt0@e)x_!
zX?)DRUj60|NQ>PI0Pl9mevtRvyM|O}{0Hk)_wcMK6M1MIJG0HVb)xjxn{!!gc-M`U$k_Ar^PFtUSJR7?&`@?%@vSg&(<
zXrG0J@E8h^7U>j`QHqnC`=M(^7mu3pmxYb8*S?hF;U?1-np|-bsaVs8DnHIT#&G^m
zW>XW+)Z}nCj50`Q)1zn!?r&+8jv%0-iiQSR@&Fwe+cf(9jcmNl
zPx}jsCRRp$`~1m}ZoilQ>91dTYGR2x%NbbWvTg87{7*HmiRJsDmUGo2-&*6%=PLgk
zoy(=@3JUq!b*Wy=H>}^|A=~QEUZoJC)NiAvdiYVzsJq$I>ox1Yh3JF1<0&HmBn+b=
z6VF~Qytq8}`<>-=e~a=SwTP8C#p1)z>)1+4*|4QIH|8t9FJPE!E3o2|f4Zsn&;D;l
zYEz&;XciAj`zCyVF41w6H-Uj4|8wp`95_1vnuxnlZR339P)k2-+QGh_6E!73c~Wxc2rYMlbwdRi|dnqsd>
zGFGxsu3)anby!EnuFBRh15gCUh7Ron8)TERVK6y923|O(%Jy9xuW;K}UQfJ%!x2-_
z)Y3!q-E^@n_n1a!Z~ey(G=S2%?nq?Ehw#V3RKBhg(=fVmxVQ*K%(K*dUo-N}9?m?z
z1c>jR8i%rh`A6?GqH>^qcC2>W4}$m9p#qJR`A
zLOAFswaj{oY^=hF)g#?v+-6+tR~i1;Bf%Uky>W8L*{A(Wlh1*pS=@qB`7&=~p3iOc
z7Pl15HEOM#P2fp?@I<-xowQ7Y(JXkhPsk&nMjq{GQjSt~s}5ge5+Jl&C{?9Z8VH@x
zka%=uo*(i9O)?gw>j#DoekD`KFiivnz)h(T?Spgfjt^+KgzlOtfDH(?E?x)F-+q
zV)+s!p;M=Kj{~+#?*E&ax)v4IUAG+bPT;Vdyk}q>9ZUHYJ3HK&sH{sP@fdI9V%55H
zIuV7;Y$uY6w>W}|t-L=y>9}r;q_afuCJQLQ*@b$J)opR3l<%alvY#8r%Ba75V*$C*
z(3*?e;hAGS-Ql?Y8Dxd;8IO)?6mX)=caPyJsO+q8f{HIlyCyNTbqPA*Oz~c}{@^K?
zJ60@FrD@~Il3XiOcin+|8V{GG*U|gfC$(QYL-8?f$?W{jHf4e~HMlzb^W#`Lwb_A8{
zh}`%!OwRe-rAE{(K>M$YJI%qkfdOPSK;G(TH`EujVxM-?fbgM)cJM*atUmKP
z0hi8h2zK%y22iU+Z^mg02a7l3W|lwf3q&wk{MA;QMsNkLg$m%sWbX=SBAuf<{>tR!
z3|)FP?~ZKH%}4`tImJuw(iK^W|)gql!`dmU7)&iyoA8AAeyW!a`zq<6n$
zD`Z-qR8Pi?eRt@jI|FMK-4s+ep}S1VfwisZ6|>_{PmLhoNIM4xu%ng~-=AiEkfl*@
zn!hX451M{eYvN7@q?(i2B(M5ZPjWwXPqy|7^&6G?m>l0igG2f+bHbL!rBUbO(e8}s
zq84ki4%#4#0N!pRUaCtxBtKH{s|n}1P8b0e$cf;UP`lV8`%>!tQpRz3lUT&(USKF+
zSlw(Q-&0Gd^&(k9rFGiXqU1!|qj0rpf0>T@O-Op3emOY&w`kDGQfr@rkx5+8w$CHP@>y~5$aBdf)YgoUe3F|a&GsB^FK
z!4|%?^-l3uHVofvIx6J<`(7Jno6JAmsHk}4M#pz#^DsJzNS@sgRzKUy=PSd?Eo!r)
zKc(pwm(OU5sGRipa>&MZ0>-FgcPyWvTOU=fve%Abj53D4EmUXA+K=4sbP3XkeDY^K2<;xAkItGsIqcPf*dg?X5=Ik{dTQj^(UVC<_OC<
zWe|aS(4yOuyWPPFSb$aQR2yXNe=^H_g97@E}l)?Y-IOzt`Nf
zqJ98?c+yq6e8`bpr+I><*Zf_Meun4b9cQRlgax%@h5TUqU;|`P6wjf)BIW_%DXA
z8oU$Qj84TKn1K=>7w8=;4c>5l+6jaZtgjf#!=4*l)%+rLL+$n2cH
z{sp6zGvZ=pGutq|<*z4yIp{@8u;8!XccF`1ZaO1wgWr4WXVf|HrqFJlLpf%eC@yFg
z6#8bGgl@1DCRe5w>nklTm<{aRZac~CL@GWHIS0x=Lab`1d5V_Lf!v6#Zf)1ft3M8}
zrg|)GhG-tthX)9KKz8RjLgRaqNx78D6Y0sGn-&K~6E$rSrIrU&gvVFD*QYPWVQmu%N0)t;I5|>dd!i4@Xv5jBdzu^l+lBv~ZQ`36n(iU?gMP_Jv*L@|JiA$W
z)~FDOmv-wdqh%e5hHSL4zo9RCT_HkNWKxvAHyiqy24)CvwWkrItBme;9!UP{
zT@F3b^W@dv2T`x4U$HhSF**?v3`Bj7F%LcNN*v*S?I_KO7Kvu@Aw_cPnTD-##A|7~
zg|HH(Skv_{n5`$%T=g?T{&yB^uWhVi9V+ie;KdS#0UF!xw5ACW3AE-rb8`0Q
zafXsHO3BeLW#{*9s{f7S=8HoOYn!H0FSkRIga=MUVI5{*$v
z@;uTyezKJ87mLr7yK31-dZ}_uOasm-Kne^_!R4K$HRyUtZ4}uPzH~x-cwLMP6}m1c
z_zXF#)sfuRIr1$DY>K%3627ZnVfIFU8gRI|9Z@5nRV-SeXOlt^H6zw{{>I@h&Z19y
zg1{I3Kea}zMGl+^Lt4V~$}4g;7N9o)A6>Js7WWsj^#%+R
zs7lj0fYyPlvW>VKhl`84;$R+BZxwnhW+Hwc;<~Pc`{B&8X@U*CXxw2p
zs1#;%z#&{C;_2`Nrbv{AEs5ryBt|mnbr?pyO|^P|Q#!{+o=t@mY@N}+D$j}d`3(CK
zJ8uPqq+A-#;|QGP$ma8(Op*w$oQMqRUS9)9NTLQSj2IQZeD>4|{HUP)*;VW1@#j}E
zlxG&cMjt8qzl*>>Yk^gVWDIPCYKmSV3Q2(cg7`@o@fFpClZLIM4_|*+dAe0|$&osf8vSW}|IsHr>sU0br%m|{GuP)e?tU>WMID<^
zL%E-ebJSB)QxzmLKUG)aUrQ4dpt>+(CM7b#Wh(R1hr?i7$bNWtqzHy`@ET{Op4dmU
z(?}=3?{5OM_i#m>$Oj~pI5nD7>A$_6lZ>`rDpp^4gP=Enf>80iRSI#S3iI&?T}8@U
z-&=SGtrZ~80Iv+=_c}J}%id~$AYB`o)8FCrBg&h-%>g%@E#7Cq)IGg~m$oD|4S2HP
z`L$#g_3t%2cDsV_SQq6@CMx1;7AQ7H8Oqg>I-GD&Y~mL4v`Re7v|PM#vt#DbLe!5<
zjYfgMUAA?ha{70&J!omb)2=W)+rL$W4JCz~GqUNI6wEDQGU@Y}iOr
z{oSJGoqwz612x-6_`MpgzthbUb;(jQF$1akmuZXdJdC*j7p@Hc_?q;jO25l(K#DX(
z>dL4@tRlP4b!6($w8>Rssm(KOi3q6Wt40%Cds9Eg%W@d85C|meYy5iESZPPi=dlT*
z!JM@^$eQ{)k$*co%BH*?-xAi%o_>Y!mF~42JFj&UFI#MqK1^8^lb8Hl7<_WN;p%1=
zg+di9i2A^7o`rL%eoOF+46ts<
zFMeVl2o0?H`bJXs!Jly}d5bsY>MO4>7dQK3K_7W*JBhv}LP}p51q!ic5I89ELR{cb
zKNB;z-g7|`Mzkjw>=wHDJZ%`FijC=M4Je9=iir)1Z{mc3zy86JGNi$$RTAp
zBnLv@ODx`O>V-4tC;$*ILKScOz*`WqFd9WR7UHnful-|(fH%;fSLxF>v&zO^yauWn%!lgn-{7
zXBdJ%fN3>xAv0F+pN;F*+MldL=2@;a#ZKqLoW
zDI*`TPaOMS6VpPajEsCqk@&}wk?BC0FgQLi_dHiJy!BiBn=83^ro$LGnq{oBpX5_o6yc521!VqoH
zE&6h~1I$_V?qqW*sCu}K0`L?bpb52}QSVKORIkd7{H2gUn`w{DEEA>mjXd-PCKQaI
zNnDO%1diCVnMZXqj3>aAj?-JX06u{=8&d>tqpT=gVV{)$CihAcScr2GSGXy;Qly&i
z6c4BYW#RUq-bL>%E@DpvFTe-ETa+q5C0o1;rwT!sB*H+S>Kq!GG!KTwE)24)E{)Qc
zbTxw%P(_MX7!y>(XaKti-y_4hkfHx>AkArneRZi2G}7F=N{kIGYdI6oJ~pOU`u}8R
zJ6KA%!hl(*n>j5Ibi-hlW@~g*Y9hh->r%a>N2VJ2^&{R=ZVujHQY*8}!*dcD$8RHC
z4F=5?Qd=K1cI;Q|L4~z1wDGjJH)n7|o`+~y=#R^6PkNG}SjF=r)k`abt~PCRx(UOp
zU&u%$$@^BV!HF#9;9pLAN&7IO-{_Z$#Zq*~a_9Yc(p-cW~T>vBY
z^7XciJriN=#-chGU69`_Ju${_)WhR0BSovHjinl(n!%Z=z_kh`=8m?y62zqP#SVJs|;GNObs
zfOMcr{4PUZVMv{e=AGOUX`@h=ZE_eE+?EIvY?3?JuAl>CvM&ThC!T9|sRO~0e#Lea
zQ)_}akJMB=J?Kr+l%*!Qf=Nh=NqL{_ckv&Y<^hw2z|%-KfkpY^rAR+PYE$4O+?R3E
z9rPi|nQxDQenk45$Xqd$rcvrEXHVgAf*pgUN=^t|V%$K?@{VxM!(JqQVSdVW8_Fi7
z$q5c%N%4Vb>rZIwNAUn}aIw=`?D7>PDNEnw9=a3G*>wp@ucLgz;;~ATUVse()yDFLrSXE>9X$IC+fw(s0J4M6GSDl
zTgeqY&Y)Wk1o{otrc?fh-o`m`tKsS}T7Bv@>ucUTz~#4gNyT8Sz`)FkhrLe1z#@qr
zrYNVE+Cb$=c%XrdBjl!m1t%YDtL)*=?Y0s!J^AT}%|BUB%dU(RzKRgNMOFQqruQEt@)#8-5&tW(7k=U4&I!)N7xkm~
zQ_Ll(M2Ezzyi|_cfp((4jm2`fXUn553EiV~y
zuxZ2ER-cWsw@_qvHPdq!pjZP87(tBEfP5K9Xe?GMa`+m09nm>8MIvp|0X|idVHp|z
zenIt`v{mn*k#JFBfTz<(pk^v#oBd!~X0!CvI10lV_0mYPSWL^(uleY`-E^e*fH-ob
zKNe!@ptA{OeXjJvCFc7m!|$%&F)vK8>^p`KR{6OVQpOvpk5jwDtuOTpgc+|Hajjz*
zQ*0rWSBiy+`Ap*ika$rS9H=Y^m0H9#R_6(z<26#@CTn1)7?$(f`)kbeTb+Cb2nG~S
z)s}QzNv!G(T887`DcRe6h0SMT&PPKe<;t28?x-n=h-;kvxWZ2fWjjBX)k4XNmKa=<
z$q^qLAE!Xy+`$`^&}c~IGd^adoKA2)T3H4eQb}Lu!GhF)AEmG~u+s58yj7x@)^x(3
zF}(3vpz1ahR@gK`w*yb7EZ!Z&lo&6tBjX=uQ-XXf$^+70MAif(Et}SjGSIKh(wzj9
z8DDSxO{1ypuJ@58k1$$W+am!!Xenuzx3qwPK3BWTD+8(bGCr@zkO)cV4~5%l-dVGq
z!mV!l4X5LXgZ@_g%4F;@XqSU-&$6$U`;{irDbv!3{~79k=3G5rof%xl5;hdZaiR=O
zpq7I9e}iow3L(q7O{@x4R5IqbKZ>_j_8^+RlS}T^)dNLjVHf;HnZC+LWO86ab|4l!
z+$>+Xq_V6Ju2!<8KLw{ArFOyQ(9fUx;%xKeg(2Ua%`7o*7Fp;~C*3u*eSKq7Jbc<+
zUcvZf(=MCFm))5ACUdfU+iuiFx507FPQpcll7Yqer^ciReNws^u@_5Af$<_A!3A__
z-xCuPgHJ>4+>}`59AxNywr$>g-;HMTc2r3laTw-Mv4v)&&jI-pH*#Yz9x!lDDMK>?
zsw+kv0XosT<{@4Tbupf1$=^Ea^&^B`T*O|*0AJ!^5LO=gqY*p
zSVgx5_Y|XWG6l3rE>lQjoyL5!lyW>U2NxfY-eq-?Vfxn3dWz4AJ?H>g*e4?(ZPim-WETH0f>aZ7
z1mMUPdKt%GYu0YD;96>c%~u$i1yVVhYM?1SeeI1)?>Y~ad!f=3ZygXb${!jlMQ+$-
zwJ{yKw9Pc}J&&B#vMQ@;j60vEPq9xgs{!?}x7|@^+n?rvsb#SxPG%%A
zZa9-*mn5QWtxN>rnqKxnnM8F4zW7UiXJ~CCpR*B?GB2Vgj>UnkVXYB{RM}YA-Jtzx
z6bHF}UhQ;sXtO!#j$Y?A;U{htXyug?rd;0G5x!Sx`(+jHfGKX@$|J{v`)w|_$M4?a
zy`3igg-+)(H#l&mTo%>X&vksl^;5E~ihNVNd}fY9GLn6)sBqHSy|z?6&{^p74b`&8@%+mAC
z33fqq@TpIA*-($oBph;A*=KTnv(t#GyHH+Gt}ErX<}RPAN^%}~Zk7vvcnD=Che%+g
zq^4kOg)?-#!X#hA89d9z4y=?CIJkc3brB$&E~^Y1oG=|n8gh_e5Vxj8K8rB?Jk2Sd5xOED
z#)4$33tO2Hr$fa_<{~V{w2dd@AxJ-}JMcoA+vh7fT%gg@m(+$SDBC1*dT<{xaoFbrLak%PcxDBOF!iH0yG
zj&BkCs%@T-BQh5W1R3(ESvlT%)E4^*=J7zE11ra$6GDgidR2`#-77Y=|3W5MA@bFmh`T+iI`boRUsa$}LF`o1YS&ha`s(QS7t4wqOWr&SXJr!jmnvm3qHe1wfF
znC5=i{NbP&DMH1&<~2CR@ihk2QtlnEq)N@a^Vb?BQIr*YzEOSDP#DEph+Gf07&@^(
zLh0JlWvxOq2!Av~Uc(ysGsdD#$9+OKqh&eCHqJK0BX3@Uo-f^rct&`IE$Ash^Cx
z<4a#}A{TxnO`1j6;j=-sXUJNJ`3tXRwcs1M1*dk!w};vZ)>*%uwwgbWHC4uq?wU_L3*kVsbZBnehW`ftIPJFJ
z?O*f(1``EA_Z2>+c!#@H?V$#}jvNl=lF!%eJq*IgLA3L!=N>F*f{wpbR8}!{4L*Wv
z{6F6?B?vWf6(GUVD2c>2QGUQ5HAy?Pe59060T2oSKn{cp#pTvXx;rc`DCFB7&188`
zOeC6Ca|f+c4bfr%28nYQna#Pu{j?M5L0~A!TTS=B2!g*NlRjhN*~dW*`LZ-PktX0B
z+tVl%hNK?4A~4(?>U|G)rzjgU`g;ywK&ze;1E;PsdtPM5XenJ>8tKGKC5znL3x0d-INgfL?03aW~w?wZo#k~V){W!
z#H)k~07$$W;yXZ|idtPtzK}7=tqCH0X*;7XGZ)@x6@syzk8NUhv3h~H|9pFeF
zA;0s!nPgup1G)EXr`s}Y*CV)B2ZT^nk`CTGmGs{K2yQqrD=~vYTA%qeYOhtC@PeAz
zH4a5k91Rz{h{pI)MbTHCZ<{X#?nQNa%H>>SuzDHed-g-cheUeYHZQ{wRXIG%xJGBbGRPI2
zHH2nP0EouFX;jwJ`9F7=h|yqW@5P$6K|=R#-Ks-tp@F)5-S@z6A!-Gb25Oi-pF26BgvaG5Ma
zYI`MB_54sc{xrI!5$tEiZtOMQ*b
z*8<}nht4?O8l2M}IJmr8()USk?5Ke$gQfDe7*MH)U9tzLt0Me{msq=0_8b~5<9ZWd
z&tdwP!E)CWucz8>;z)+*tR0;8t|WP+`|ye*c_ggy@JD0f)?KaTe|Uv4!W31U?t_4t
zovLT|2lr%$JWufD5{^_(B~^vL9f*(w5-#^VL~N?vijigTU|B0#>t#2jmbu(0-fm$y
zWAh?+@m0g$@jhJ)eqXx;>3HF#=bP8xh6b}!Un?HM{Z#9*Ywl>9Ah-~xH@a|7@te6z
zAZ?axK}$qsKK(o~C2O=4+&fz!g4~j~Y(UuDhvTa$KOROb&&$
z+oSx2@z=W6jjPDq>Wgy&nr`M{A&!mGpIRL_I7IQFFm1u?x7@wVcmZ_*6%(dJ30<-T
z+!#d~8AD@0NjCB^4d=*cO?p8=?B7MAa3#*yc5c8bh-f;8ZKw0Bbqo6qob?1u1JBWP
zzEE4SWIIziVf|yghb3!{<(2nMPQCj0U~7D_K+MZX(s*%7eI-Ft
zjeN};(+vTh&?~=H#A6NY=XFY#Q8)aP>L`Q=X5)aV{Dr;5vs&aG0P88sj@9gCcHE}hrH>$$q4
z&K=jhBKXsp)8*z(i58Lwz6NHSv`sebxKjS`7Y1r4hkgDeJcN|WJDxWF@Q#-J6tw;2
zhirW9L0ESU@q`>*iKxz>O)uGgC~%gK{3`W|wU$+`RE{&;jV!p5y4S+uMiE@4#W2N3
zV&j12KTP{6_rbg$U9V{N!*J0%aFkjXq{q$ZjZv5Dm1je7Ls(BV@pdt5hzhepMhTVj
zJ9cMn52D=uHjujhb0Ks`r0QPf#C&C|fy@o}M3&^UNDYzkwrqZ6@B1r}7nh{wVP3ds
zm$>p(?PxkcKmkaDl{nR+#Y|l1NXL@flY8n!=;!$RiR%Q)b7%FQjy(BTE*SX0nX1gDlbd+OdMh2#^vFA8{hT;+Z+&ZFnP8+S{W7!Yj
zl;H4r{gYTTie=4^{`HGchz1)~Y1V1Tqd(1$-XnVHaHFWnUv6vmr}77^G9Su
zwq(pse)^XSqns^|^jN5CH2bB@y+m7Dm)%lYV7O&8KUsNtaO}2ScOM~-p3sYwe8Cx#
z7Rn#8Q|2acpbD0CXT>OVn|co2SXLQI`3c=!59@{b7zTdw&&f$*$+9I3cigbjQ4#}l
z=3Y#S%A|TM-v(djbl(P5yx&AD(J=9mE0IIzoXrxL=-$yxh<;A>Xje6Qr$iTEAfX3m
zScUB@+N)=DcDq8MCXFk}+}1Yciz9~|7Rz(dgGTTQV_v>-h`#ysD@g4;{eDCqa#yb+
zmEc0dK?Pv8QyHNuXVSt7n9>pAmi+iQb~T~+zeuv~W15@Mf!SYX?3+pPK?%-7*JyaN
zG#yws+GOjaAIj7g^W1ba&A%y*(eX18$-R_gRe!1Q@bR83S)7DF$;v|i*d^33qHAP6
zp+Dz}3)^}4261FMXmR*M4;e$71;k{?5zT#?4!P2>5iehhVGE=fdK?T
zy1PNT8|m&4l5^vOzR$bA{T=S}Ih*@mYhAx<&0x!b>`Boz*cwJ+8s6h$RDv5PkJc*uPr=bceshd)I)Ff(}aEnkd@wci9^g{YSd3k?VzZJ+S6T?5Q<(PtaM;u${6urcHcF)fb_ed;_ehr70e
z>Dn>K1k`0u)FSno-T~h+4F4w@aBrra0po_ON@2d2oUG{q^Pq~UA7;VcgtGl46rdAD
zjDr*qF$~GKDb(X(S38;A{$gl3>
z-b1*Kif)K42(A0tj7Dlqit8V(_GbabAlG{{a;O%l{}U7X%^xsSg{QURlNpk4?w>^Q
zK2R;b_y2(Y#l4LfM4LE_zTqjf=}&PP__3h=RI22``uB_YG|+P-9@;H#kg&(#oN(F(
zAG!xnqdgW(doDR%W&VXmJEE4Meo{4GL+i{GnY5J*i>-tsLYSEeqG$>ZTLilgcj|T(
zm_!AUESr&t=zD}qaup=xRHIjoM)#)xkt~D>1RYa@eryVCp;-54VwbfS#*^EPp}-e;
zp`Bo&R%y>{-+*J8ehz6XtCLAeDH1q{ir4rj^}`pH*=xO1twLefx)>o6RkUO)nYPA@$aV#cEw
zA)XBWKJaE2Q^&x&LbfqEx}@d*l{V!6|AKXiheeqRlw?`fvAdJrLHY)h3j1{mB~9IL
zu`9r@52!RFpF`v{7YaHF+Pvi;0&XXfOf_pDTi{u+Y@I6At%ni){NLr~VIG2W&v8#U4
zMal+Zou_&YwZAhj=UjV~FDSZ19q_Z#vlXZn!)PyQ9xJmSc2H4K`%zcaT5`QgCKM>V
zC7k{do;?#j=?0lYZw_Rw*0k5P9d;dM!F?c>K?0#@iW!Z?xVCM)LiuHchZ+^(@wTj!)P3+FDp1r+o_h_ks$&&1>A7*NieHr5Bvz+A3_<@X`Lh2tM1zU?
z3=XUoy!apFMJ7%$Nag<>&jJgQi_V`WsL8J#IFOp
zB)_xbF>$;r?rWf$HJ)8OX{5ASo8HcP;OWAZk*`lYJ3EsV{ClE%+(ZPxw>gbsah;EM
z4b5#b$^N(tLiL4s&WinShfx@%!OCMTu%!;x&WKa+hreKnV)t$Fk3lPt-IVbji>5i|
zWMc4FxrC7(2BEbROqCYN2#3!O>AIP!+42zCT<}j2f(T{_WE$7qyneBXBqTG0i`l-taXRFDAaIsOzNo6_byBf5JXjQ1a}_?Z3Z
zR57z>0Us(8=Rl;Oh}#0$8jQto0IL~hd2vU@MjyC6K=Vix
z3!CKocDIY+LG>jk-Oxc{9kDWlpqzbr6Lp_NQ8u^W^7~f9gBj~J=F+c5jAC(l@B$l=
ztM>Wg)LP73ri~iMZ>$6@){+pngOUhcicSiE{yWzQe>PWcV!EI~hDEg``{I;Lb2Ce}i^XenZbC=-FdhWDdn(s;MsaXBf#1d0FTCC)Q^ckBLS$5xrn7Li(Af+3#Vxyvh@
zt$C#9T9a&(BF{EkC89aowjxc5T-+tYj~F0c1$+`aR=D=4-?`}AC2nJ@{aO?_8Q9Y)iQ
z`(x=E+L4QRLx&Qs7V`1F8h$#N*G$+ben`P`&nh!{=^dKn9_p
zU268h9yPUuWf89vf1@=O$c>aDG%B^_$M$f6gZs-W92L^5XJ@;>0&W6;FXC(C%~)ZZ
zTJdC4d~6T&n5ei98hq;#(7%s?%M>jy7{tKU;Wh_5L=R
zOwhyhFBDf7d);EBMOaXD+{TAsZSk&O!tlH=(*OBy<&SEt3K{2akuGkS-REW!Q&AGC
zIx1TayJ58bG{AYCGN&;*;$NRnl0LoL&UMhS-)Y$rVK3a3r%v~CY-;$Y)>so6%)*6Ptf#&-NpML4zra&c0Edg(W
z=z?(2<089#ZMRVl{e4w=q>A=Fkw3A6dcZpATh8;-Q|%;%aZ5VgIqA>9M5oFEnP{
zWoPAsb94BgbK6ZnUkXgRWhaQVOU$$+2M}hFa${h7B`#Vt=PQ0ux}R*!#W4{zxpF&<
z2%LJ*M-|Vgj_qy#b(|?Pm0}TY8m^Zj57BptnFQ0N)rg6nyckc0=vdDp9Oif+UP5Sg
z<2J|Lp_qb^ZDeT?TypK?GQ@)ltxEE&f>>Wy2~ck#!_nIn4SY-&3SFq(v=I724AbG|
zQQ83kkhb;$m&AwvU1q%hXPK@42~(%Si4Y@l>#Xv{2pe$mav2nf!phndLNoW43`Y*9
zrIQL}dj?*J29ym{mE*5h)q4`w$MsqNOvZ@espqyc`vM6zQapOJtbI`&2>}VzChSqX
zXI|E8AT`~;yi9f}wh!>Rt?N%NBRfRt>pZ&k%ZdP%EP2ojZN*@z^7x^IveULJ!Kvg_
zq`QvhUPN@z_fzqQ{=9k8$hrVs(ynUMQ1S}1s+C%TNSuhi9r0Am2kJ4Lfk#|aBs!1%
zkHBiY6E{wN1(dZtoRGU*x@oe=;}s@g0+1y7H9AV-HO^T>m+okzlkVArrD%^+520mY
zh)uQ37e{eNcS?E^%v^4qFKhO$%WA=ajKU{Ap^CjroXxXiwK#O*bX%tTALzOhvn)h$
zILO3&YqTEG*(qM-^>P
z@wGH4U`CyyBOYF2lydwGk@WxHzs&Fd`^%~m45`B4jch!AY8CZ4F^~Q!Gi!9>Z~KG2
znSR3S^h;vUc+K^|gF&}MKL%8VTw#tk!(>U*hCkaoqMdo9IckA
zawoITu}|5|Lwx7W<#4`|s3L82mT)VdbJ-U@Q^0+sLD5e+BHHphk^O_2%3cbT`%`K;
zoeEx*XR2p3;n=ovzTm8yVJNE_p;)c>;iISQ?*xKGNQDmbf`3gpU46MX;grAOko<{B
z(Qr3LS%o!l+LhgD-tq5M5J{KD0{0d@SjN7LMOv&?kCinvj}~jt6?H3|CL;}RG1QkGqm&)N4-s)#Nth^Ja4
zj?K<~LDUZ+D!whY<)^fJXl-|wDl`IFuS`K}ji_(C?#*K-2y)ycpgO}w%T6qWetav_
z;i}9LcgDm_EvdMWNRi|K+?i(>^?h_MWYh@hjh}Yv7crBRvUnOt@>HkF7{XJu1)wh+G
z7_MTff(Nq&ikoXrZ9v
zfWDp}KR~_XaOPPOUNZFU+tUaJH_3||gaVR$oD=5cWR=`4ZxQlh?s0V*UJECB{^L~U
zAoPNV&yT1#Jvgj%c7mJTH%l-3ukI5fI8C0`C^gY2)u765mnv3kjAiQ;%P
zg$RI1i7npSAs}`DyhWvmPfvsnU8y9y;5L5uX%F?t6}DG#_40#!OjusRy2r;B>qTM@
zt2z%Vs~|R0hzR;8QKIYycwg8RIhplnWjrd0INWQ_5
zJEtQt*{y{6k1BFPDRNtYXvz6^y9nb{sRZ>PoO=!W1_-(!l1zsTPogMx#;T&wWzN4m
zU#?_*>&%?q!RD~5T(~l#7=`$Z?Ob&3f@tM1pR&!$lt<|juQaJ8gp@a2AD~H0wGR?J
z$kBoeCGC+~Js6g;9CUPYP+%q4*z?JeF9S2fH!759PNt^PANM!
zrIk^GdK-4~q2R%#>J63k_?|Q&VUk%Fvf=kNIqEf2&C2`Du8Taiz3jcl3vi8=qjb<)
zv{aez4K$WWDkvs$<_f$`hMru;@dLNSVAo93@INHG-Gt^PFC>QN-61ddoovTE$m<@oVhUFU9H~?GTh4oU=wyIZO64G$cz#eV
z^GFAQo3GegpN=>@Vtyw!b@1w2E#%JSqJQZP?adm^kDPs+)eQ9{@Q^Ly20ahH(nYJ?
z>i3jNgksyT%+pTO{pVds(2^mrvfy{|Rzb$-TLH$C}O~JamaQRb)L$z2+6*U+q
z_{ep64zkDLs;9WWNIpa@6JeN8Fv@)Ow)sOJ?sceNR~AZ&f3KklW*bQ;a3@qNlqY%w
z!Win=?VsAodt
z*ArrMJskegu4?!oLPBMO_xR@Os<;}VA{oy(hVyT0)Er06ez6#8tq*{IMcp#rzYidJ
zVz^;RDpc2sgC#S^V`PVBCHusCP@>=nzI(#wBYBaapgd~q_hh;6nlGO0V_GK;4n2Ry
zpCu6jkt4K8)oSi^Z8P?!2FEqWV!)KYmO8cq9;2#zk#Y5LRZf;Mmai;5489nA^n9xH
zoV-Av;7Q_v7NoR!N*vA6+wD61_w~Y^)4=O(A
znE7Kn3(_`ZqX1DxEXZNl$lH{NmKBe;RVBwC(ajke>3X`JzN(gU6@HY
zrWrECr;ubd2sKxpJ1h;3=iS?C&=ArVjA*wY)g*%Uv5@2|bp%EFW+L+x
zYIxZYW<1Z=Y?lfOt;VSeSUM@>M~)PB2Au!(A%lrtbPL<9zmZMmiv9%RagmY8M@!8p
z+1c4w`;l1WlT`oqjr)K5MwHce{tg1)u(sv(%r$4)i>a@51)Noq+v(?7RN@%rgf=M@
zX^A%DNv>~v@mDncXG%WBple~huELa6+*L>^k)ex$sJ~^6wdg7`<|mlbbWgZEVu5Jd
zx@5-({4~G62Xw_T4(*1oJ5}WErN~Zh#ryrYlW5VRhZ$b}M5CsfmYWL#a@~a49UeAc
zFN|zye$#5`v<>5=l-4PStF088?KbREI8-NT4U77RTSQ`#`jbV0if@Y<61?@gAv|
zu7~Wxh2dWIW(c$tL$qNalQ%BWKT`pM#ZL?;E*jm^?144Mv@EhlxE&t*lwaeEye+w0
zwF|Fosh!t6_=R?nVv`aTdqarpJ`5m3G)W(m2=b7719);eR1U_EV;jF|`If!eDU%+W
zH611AE8MG??UdOkf^ZG;!lNr8-54B4r6u6!P~TqfkNkj{
z_4OIKDFXrR4=}+)F7{^04oiCXN3dFQwAQ1R=7I9=JK4elSk0uth)IG|s8h>eL)J=)
zVyE1*kG>ugF^qN_{V*q_L|tt>Z6IoYh5d{!84PozJ;OX4JH<(ml3}~HEMg{zjO5C{(~jT06MY`HYJ{@D0W`$6*Z`kA2c@6i(9(MtRDQSZj5
zvNA?7zbn+cyF165($dnc(FAhMdn-#z)0gL`#U5Yx&pZEwbM@SXe}8TsAIJ4fs@$i)
zcJ}opfQf9Y?
zw(8dC$D9_;(oJVbnXT4lc
z3;uv;1z%TDPiD0qvMV!FX3L|`OQyr1^UsCIvJ}DXN}_U`!yVhB``1S^=KcBPzDCXU
z8fL8!`I{}4Eb2ByoWVR=t{zlAIva~(ON`M8#?};HCQF-SP)5wN>vad2J0Py_
z4v3mK3E0%h0f4mU=JIgW+icKH8&M*NMV+W
z6lrI>}r7
z)ecA;>qjIDhWq$NQB)w=O(+>Yk-vM4V#S%?9mY)mn68{$N$nwIVIh0{=94WnPkJu}
zGZRzVWTRIun|PNh^l-bAgoX-TvhURu26c3r+5VX$$(i!vI{-nC4@$YdRllGFe95Y=
ze>Iw6;N!MSON`161B!0McNXiVKs-X~h*qIfec9y?6bn?VF;4yOx&NGLm1V|+Ws
zI)u99HCL!K>wI_sXc!_Xgi3K3iWL~~Wc%5cN-$|E+XmM!Vim%AUe?DG1|aNCuUJJi
z0V-x(J;G%ImGvo74@aiSinoJ;lhDtAvYqY+4}pclDDZ>Nx#VqUZ9tmotK733>rs(d
z5zP&*MYL%&AfcMvnJ#wU1!Rm%k4w>@VVY1b>0lLcGWIdrWa{vcM`@05yRHZ4Cb|=dFo8$MZz2+n6SHI^cd>pj3iL~pD6P(t%NZSdCh|K4!v|!)9@jC(9
zySx1;;Db!Ee8XyqhZI>Mw_S1pZ=`m^MCqdGjs2^OjP-PVvIpDpH;zDB{q@($$*A(M
z_8=LhR7*07)I3DGYaIIi2-u&Nc`yAR?94^xGm?I>MYIeX*b=DddMZdfekjz`O(@8cIfdW5=*=HKJ8K`JniW
zRh7ZNnA`)(uaX?Q*d_EBh(k(6;s648cHIe?vKGasawXk1-@TR*8_AZ>9BMEPqHN5l
z5zOx6Y^QP>kY=EJJC(bEk0meI!EARpMPMzTcQ-sv6DNf>nhYW%o}-P6J8~)V7fzE8
zB-DZ-Q=s;fglp^4^?N~qf<7zpa4}lPH-u*8O7JU`-XsXZYMn9
zfaM4&1%_U_s?K^T=PEyx_9ucYjtr0^NSVy8+(BUKop_0{n(bLmJWUkt
z(xf$yQ`fhAOP+lY9h7~)*^0)ItxBh#S8
z>~V9a(DbH|qeJV))5X)J{O82&K455yEIO(?P+KKmi)EYb>|j!iolrGgrG(!1{C?Zy
z42d*g9KB5%BHngRfH?QO(8EHD8W{XO$4ovDc@qHE_VvJP8Ihub+U&WX&!`
z=9^4u@It|4Efd{fDG8MX8S!A<7RoJ~OGCIejk#=MrCDvHc`IO{-B
z=1HbnNA|2|tUFc7gdZb1I)pm9JBL#KAcB6zw)Nn-?hm_%a!nI&t-4fnurB+39{L&~
zIZfSWXdr}2NlPpThL8Zy0uHx35nZ&JS^33}=;En#W;{n7Hx<7YzGpant!h+ue$9{Z
zh4jkPe&wfnp?uZ(qcB`_bG+bGzF^CcDylG~ct-TV2X`rPiMgF)46YiXfk=N&z#weW
zVBADbXNT~-7MC7bfQRmVDF!e1VEiQ5{Rot8$JP`~dn{Xez*19?;fq#8&w8XTOI8eIx;pT_T`F>gHZYqm>PGeUH+}VpX?!+Y(tGl9
zh2G#xI1xY35y!twBuD&sYS^=9ZfB8@LP%HuUERdSsMVc82TfljIuvWKl4!H%B30s}
z>i|7*hE9BENnc*rIGCh{EVd&{D1F@Dyuh5f%xaoHRxA}`Nic7DVDG31{!0@Qg+pjFKX)rmCbt
zo59LWPlV1wj5F(kq;MfkN<7gdRFEw%croS
zZ6}g!VJaKm`c02LtD0}Ck>}YOvuw>o!g^3u+0i2ICj3ab#DPWF8-=7vxmXeWr+FD3
zBXR1T`8jULTD0G_%a4kO;;1h4+#jVdjgK@Mxlx*)M5*(n3hL#5Om2UDU+YCNKi-*{
z>%*gtmK?F!uD`lVt#e{SOEwRv}t)xDT*MypfzD|H>>D-EwdJEbzXCd*So+
zuKB!fdEb9v2u7~Z=XgA7ghcwF4G`F0qFD6
z+J}F!DZESvMvr@O3L))`2&i1LWT?1GKDRTG}=`V=$TXO7GVaf~`&XX^!S_plZgP9jxeww2Pqqhvc{by>I
znL&uVi`!y%k*h7p)FvfzjpJEp0*tK+_QIagxUxWhmGu@@L}a-YH}FUc9HEuVO1Sju
zJj5L8-?i*wU|`yLUsg5DBiS1mu3aT;g6J2WCIRyM_vGPvP#WbvBWs58r2j3+iQdV#
zucp4u;vVP_LQZk{S#L6S@oU0>!h-cGI=b6^A6vk2S(fcwL1Ba{CBWC$mzkNlfH0`!
zKB*?z|8!m4<772kaj>b0ciCFWc)Eq-BcM6o05M(}QbS%o<7qEI9rS%6mlxI12{9uazbFH!pETln>K{oS(ouM%
z_@k-j_^QxXto&z1tElhUkhiiQE$Eq3iB?if72P*P=_m&vRJj(F)R?7IxP#bs=5#
z`}X}oZ9yNUno`q5e{c84(c}T{tnIq6+Hk7TAk_Y$*V2b1m~LwzB|3w#NJZoo>30VL
zYwwZxBczVhCxIY(t&e@lC`;6)b~hp(C)5c0$oZ3CMZ#BcC
zles?c9w=shPruK|zq(73F!^&S`+&x2xjvjZDQ-KS6Skgx?ZHE$XZhLt152s(z>8UfDs1y5IiVdQyXHi2nL5rjB5Nia2iRZ$*MC
zQJLx~nOOa}&_PaS)}zR&2yr!I-MFqnX%c-Qq6DNl?q+~nB1{C{?~NWAdmUMp^sM~M
z8rntL!mqxh^B+{P%NI8BnP~;ywYFQT1ByIeR}(h5E*2}!`
zKq&(kVp-17@>DKYRX2h_@E^byB-8nM|SZ`J#M`iS8C6ZZ
zGB$<6iiL9ZsRA%z>3gY?>)s@(rlzJUi;aS>Z)YRx>DkthZi6XY?VhOL6?5$*7~W+n
zN0@}i9MN&s>bQ-a{iPQuk6dIC^5Bf5s*oJ=6M0to>1X_{l~B^GE{J_2r5J$xyLG5w
z&Pz7U3;TBqfO;Lr8yd-b%wG?huGvn`bI61_p?hRg@`$@j07l6%%D?5Bzl}Wm;?(Pn
zVJ*`mBo8q_4|brx1EdhEBvV^NP4+8|N0XmoL+&NoJA>N1-HrOb>K@FPUFD9ngwUrzqKBp=3;!E92iST{#uB+5bp7aP#CAL{Z1jjR
z9_BP$aFlY821NXp0@X5&BBqvpF^``__8h6e##tRy;CP!a>@F@u59dhl?g@6irOkqL
z2!46XxCVKr?|v-Q2omnv{VUu4h@QTuQ!Lge==bYqA!{GC{ly=(kCLcQ7g0qL!HCGe
zl16798$C>Ho_u$OF)+G|-kWj#AXu>Q84!BQkricyMBeaalSYNWaJco+`;7Yq9FX(E
z(L03_ME8DdH+%Xf6oEPRFn+8m>8$fNN|q}`
zW`@;Wt(*e2l$Yw)c-|t(7t&4^5sx$vnId{n3_5=#i=27=HccL@M1K-G%@1*kpk6rL
z-DZl$UNl|mPQvJ!z99GY)Kz&wJ7JLs?~#m#c2)d!6d7>TA3k`s6gLkwuYv7@k&cp%
z@WS!JgQMkt-pwNq63;r%Gxw_@hp{D9(JD1FH8btq1Si2!-3($5)!BTrK~z~ZLSR~g
z0$#=8yYG=JKP$N`4N0|o8jGMr?FRAo6ZS&xIciR7P6U2CE6N|p$FRPv{4L{s=a>>d
ziLF=BQZHMeOe3C&D$bOMk1a%HOMBHO>Lg!cQSvNOWV_->AxXvGqn~~zG9bMheb{p-
z2Sk3(AJL+TShs9`z0>(syRG>dh{zbTc3~)1c7S>7iwnYI#ncNwaxs5s?ni(+n)X<0t^ZzLJ#>$D5a*5FKtRjK
zMH{DW*x>@V((YW&TX!AwFW^4e?R&}S=H|9cVFTO!npSzW(&Nkbl!ihdwH-I|L(brIYlRk1c47C9Dp`+Z6(p-s{-yzcCr2t3M0Wz?g{GuR9l
zVa&kXb{c=}C7SG%A92+|!5(7}*3#iH%z~#@T~&-%FHl8&^EGjk?@;iAW^l@V0Lv5T
zsTjGqy@*R9Em#R)AP@*-xS9OX7A&0yaF-N-&T%phU0?qfSb0U-JeXE0!K$aVMbTNp
znUtc2f@i`>qHOtFXxbM(^JDymu^VQArQ;&tDHaGTnLVK6fM#l^h+
z%*;oz+k@&JEJskbANCk1W-eQ{@fJ{YV
zp*onUN&U8Wh*udna@j?Acd>XSth}IVg-{46Yqdtax>CWw+SLY6M|sYJ2f-;A(!hFps~
zMSDNRamV5lH(eRrHIIxM{v5FtPX==gkY#xb;x%_W@aC6byo;Eb3zlqdLDK9Xw(Bmc
zZYzH#ODCzd?YA2745+@v2~zLRvh^mc3jFOd@c`#{r)}o28%_kw+z6ynk{JG+I-yI@
z3BR?KqAE99psE+$IqUAK;)oiLCBRk+TG3D|oxZFEOf+!&I3SCg7j4MNPG$_18k0Qg_?7&wF
zc)~H7Gr|`++=?&yvIhxHlE`p5uXYV~heCFwsxJvLXV04cLW-m$gEIA5uMCCj7Ee3+
zB4$>8bzqDnj^JL798Asy8~cRKM*SpS4{+8*t6GWfVdyDsS<7r0O(oid84#%^z0Lyt6Wk%@Hs#3&=AxpZP;1
zP!G}-RrCE=Cn+(-8q&7ZW#Bt=O8Z})
zEP>6PtKQ_PZIfPPTuG9U@MCJl>IbmCbOH8%U*=#Oo>;{6-Sq8>D`msj01T8;rbf9g
zH0AGl2TbJQ=6W{tuvFSKD7ar3)qbB~acNoFo-iV!Cqwf$7APK)9LU|aay$g-wmNJ6
zE1wd`MA0suRNS^d^E0U68Og-iQe7CY}?n$CVGy$tCH|t|N#m>d287^n$@rx?8te+XiQF6I40iN>d)zv&g
zc8F#4--f@J0A<7$ED=wT-}{H1mK(f1yN`w4ZSa*WdKK?TBbC2)yJ|cox6K=CZww$^
z8?=<(cw7)tyzBh{{%YgOeMfM^NIf>Q+o+&2TjZ@COu*nN?F9kY=E;%aa66JTZa5iy
zx#HkK(;xI#P+tgP7yXr#8pLLokpS{J8ABPu_yf&OL1AVFz1zo%mp5IQS=$R5(L@xM
z4z{ayEK`rtwW7;7lwH=ce=slc60?5)(KO79DbaR%JNAZ?%3gAiY5DYA_W>=_{wJ-;
z3aJsnT7yg@N