From cd96133e9e45df588c53042099a0b8b37b6c75e1 Mon Sep 17 00:00:00 2001 From: ljcc0930 Date: Thu, 8 Jul 2021 20:47:32 +0800 Subject: [PATCH 01/42] [lang] Support ti.fields with shape after materialized (#2503) * worked * finally worked qwq * Auto Format * Update python/taichi/snode/fields_builder.py * worked qwqwq * passed qwqwqwqwqqwqwqwq * u shall not pass * Auto Format Co-authored-by: Taichi Gardener Co-authored-by: Ye Kuang --- python/taichi/lang/__init__.py | 4 ++- python/taichi/lang/impl.py | 17 ++++++++----- python/taichi/lang/kernel_impl.py | 3 +-- python/taichi/snode/fields_builder.py | 35 +++++++++++++++++++++++--- tests/python/test_fields_builder.py | 25 +++++++++++++----- tests/python/test_tensor_reflection.py | 2 +- 6 files changed, 66 insertions(+), 20 deletions(-) diff --git a/python/taichi/lang/__init__.py b/python/taichi/lang/__init__.py index 97678d9c0..28aa9ff2e 100644 --- a/python/taichi/lang/__init__.py +++ b/python/taichi/lang/__init__.py @@ -235,6 +235,7 @@ def init(arch=None, ti.trace('Materializing runtime...') impl.get_runtime().prog.materialize_runtime() + FieldsBuilder.clear_finalized_fbs() impl._root_fb = FieldsBuilder() @@ -404,7 +405,8 @@ def visit(node): from taichi.lang.meta import clear_gradients clear_gradients(places) - visit(ti.root) + for root_fb in FieldsBuilder.finalized_fbs(): + visit(root_fb) def benchmark(func, repeat=300, args=()): diff --git a/python/taichi/lang/impl.py b/python/taichi/lang/impl.py index e3d010be0..c74eba62a 100644 --- a/python/taichi/lang/impl.py +++ b/python/taichi/lang/impl.py @@ -11,6 +11,7 @@ from taichi.lang.util import (cook_dtype, is_taichi_class, python_scope, taichi_scope) from taichi.misc.util import deprecated, get_traceback, warning +from taichi.snode.fields_builder import FieldsBuilder import taichi as ti @@ -237,16 +238,20 @@ def create_program(self): if self.prog is None: self.prog = _ti_core.Program() + def materialize_root_fb(self, first): + if (not root.finalized and not root.empty) or first: + root.finalize() + + if root.finalized: + global _root_fb + _root_fb = FieldsBuilder() + def materialize(self): + self.materialize_root_fb(not self.materialized) + if self.materialized: return - print('[Taichi] materializing...') - self.create_program() - - if not root.finalized: - root.finalize() - self.materialized = True not_placed = [] for var in self.global_vars: diff --git a/python/taichi/lang/kernel_impl.py b/python/taichi/lang/kernel_impl.py index eb2a07c38..2f648e4ec 100644 --- a/python/taichi/lang/kernel_impl.py +++ b/python/taichi/lang/kernel_impl.py @@ -351,8 +351,7 @@ def materialize(self, key=None, args=None, arg_features=None): _taichi_skip_traceback = 1 if key is None: key = (self.func, 0) - if not self.runtime.materialized: - self.runtime.materialize() + self.runtime.materialize() if key in self.compiled_functions: return grad_suffix = "" diff --git a/python/taichi/snode/fields_builder.py b/python/taichi/snode/fields_builder.py index 41d688ae8..0c526b9ef 100644 --- a/python/taichi/snode/fields_builder.py +++ b/python/taichi/snode/fields_builder.py @@ -5,6 +5,7 @@ from taichi.core.util import ti_core as _ti_core from taichi.lang import impl, snode from taichi.lang.exception import InvalidOperationError +from taichi.misc.util import warning _snode_registry = _ti_core.SNodeRegistry() @@ -33,10 +34,21 @@ class FieldsBuilder: # +-- pointer +-- dense +-- place(y) fb.finalize() """ + _finalized_fbs = [] + def __init__(self): self._ptr = _snode_registry.create_root() self._root = snode.SNode(self._ptr) self._finalized = False + self._empty = True + + @classmethod + def finalized_fbs(cls): + return cls._finalized_fbs + + @classmethod + def clear_finalized_fbs(cls): + cls._finalized_fbs = [] @property def ptr(self): @@ -46,16 +58,26 @@ def ptr(self): def root(self): return self._root + @property + def empty(self): + return self._empty + + @property + def finalized(self): + return self._finalized + def dense(self, indices: Union[Sequence[_Axis], _Axis], dimensions: Union[Sequence[int], int]): """Same as :func:`taichi.SNode.dense`""" self._check_not_finalized() + self._empty = False return self._root.dense(indices, dimensions) def pointer(self, indices: Union[Sequence[_Axis], _Axis], dimensions: Union[Sequence[int], int]): """Same as :func:`taichi.SNode.pointer`""" self._check_not_finalized() + self._empty = False return self._root.pointer(indices, dimensions) def hash(self, indices, dimensions): @@ -67,23 +89,27 @@ def dynamic(self, chunk_size: Optional[int] = None): """Same as :func:`taichi.SNode.dynamic`""" self._check_not_finalized() + self._empty = False return self._root.dynamic(index, dimension, chunk_size) def bitmasked(self, indices: Union[Sequence[_Axis], _Axis], dimensions: Union[Sequence[int], int]): """Same as :func:`taichi.SNode.bitmasked`""" self._check_not_finalized() + self._empty = False return self._root.bitmasked(indices, dimensions) def bit_struct(self, num_bits: int): """Same as :func:`taichi.SNode.bit_struct`""" self._check_not_finalized() + self._empty = False return self._root.bit_struct(num_bits) def bit_array(self, indices: Union[Sequence[_Axis], _Axis], dimensions: Union[Sequence[int], int], num_bits: int): """Same as :func:`taichi.SNode.bit_array`""" self._check_not_finalized() + self._empty = False return self._root.bit_array(indices, dimensions, num_bits) def place(self, @@ -92,25 +118,26 @@ def place(self, shared_exponent: bool = False): """Same as :func:`taichi.SNode.place`""" self._check_not_finalized() + self._empty = False self._root.place(*args, offset=offset, shared_exponent=shared_exponent) def lazy_grad(self): """Same as :func:`taichi.SNode.lazy_grad`""" # TODO: This complicates the implementation. Figure out why we need this self._check_not_finalized() + self._empty = False self._root.lazy_grad() def finalize(self): """Constructs the SNodeTree and finalizes this builder.""" self._check_not_finalized() + if self._empty: + warning("Finalizing an empty FieldsBuilder!") _ti_core.finalize_snode_tree(_snode_registry, self._ptr, impl.get_runtime().prog) self._finalized = True + self._finalized_fbs.append(self) def _check_not_finalized(self): if self._finalized: raise InvalidOperationError('FieldsBuilder finalized') - - @property - def finalized(self): - return self._finalized diff --git a/tests/python/test_fields_builder.py b/tests/python/test_fields_builder.py index f315b4692..04f2c351b 100644 --- a/tests/python/test_fields_builder.py +++ b/tests/python/test_fields_builder.py @@ -13,18 +13,31 @@ def test_fields_with_shape(): def func(): for i in range(n): x[i] = i - for i in range(n): - assert x[i] == i + func() + + for i in range(n): + assert x[i] == i + + y = ti.field(ti.f32, [n]) + + @ti.kernel + def func2(): for i in range(n): - x[i] = i * 2 + y[i] = i * 2 for i in range(n): - assert x[i] == i * 2 + x[i] = i * 3 + + func2() + + for i in range(n): + assert x[i] == i * 3 + assert y[i] == i * 2 func() - with pytest.raises(InvalidOperationError, match='FieldsBuilder finalized'): - y = ti.field(ti.f32, [n]) + for i in range(n): + assert x[i] == i @ti.test(arch=[ti.cpu, ti.cuda]) diff --git a/tests/python/test_tensor_reflection.py b/tests/python/test_tensor_reflection.py index 5bac5cdbb..406dfc22d 100644 --- a/tests/python/test_tensor_reflection.py +++ b/tests/python/test_tensor_reflection.py @@ -59,7 +59,7 @@ def test_unordered(): assert val.snode in blk3.get_children() assert blk3 in blk2.get_children() assert blk2 in blk1.get_children() - assert blk1 in ti.root.get_children() + assert blk1 in ti.FieldsBuilder.finalized_fbs()[0].root.get_children() expected_str = f'ti.root => dense {[n]} => dense {[n, m]}' \ f' => dense {[n, m, p]} => place {[n, m, p]}' From d30f64468cf3a57fe0909d0043281ca9c5fd7833 Mon Sep 17 00:00:00 2001 From: Jiasheng Zhang Date: Fri, 9 Jul 2021 20:55:18 +0800 Subject: [PATCH 02/42] [misc] Added palette in taichi GUI for circles (#2504) Co-authored-by: Taichi Gardener --- examples/simulation/mpm128.py | 6 +++-- examples/simulation/mpm99.py | 6 +++-- python/taichi/misc/gui.py | 41 ++++++++++++++++++++++++++++++++++- 3 files changed, 48 insertions(+), 5 deletions(-) diff --git a/examples/simulation/mpm128.py b/examples/simulation/mpm128.py index a32bf420e..8dffdd7ff 100644 --- a/examples/simulation/mpm128.py +++ b/examples/simulation/mpm128.py @@ -148,7 +148,9 @@ def reset(): attractor_strength[None] = -1 for s in range(int(2e-3 // dt)): substep() - colors = np.array([0x068587, 0xED553B, 0xEEEEF0], dtype=np.uint32) - gui.circles(x.to_numpy(), radius=1.5, color=colors[material.to_numpy()]) + gui.circles(x.to_numpy(), + radius=1.5, + palette=[0x068587, 0xED553B, 0xEEEEF0], + palette_indices=material) gui.show( ) # Change to gui.show(f'{frame:06d}.png') to write images to disk diff --git a/examples/simulation/mpm99.py b/examples/simulation/mpm99.py index 9a9d8fe59..cf71c3877 100644 --- a/examples/simulation/mpm99.py +++ b/examples/simulation/mpm99.py @@ -122,7 +122,9 @@ def initialize(): while not gui.get_event(ti.GUI.ESCAPE, ti.GUI.EXIT): for s in range(int(2e-3 // dt)): substep() - colors = np.array([0x068587, 0xED553B, 0xEEEEF0], dtype=np.uint32) - gui.circles(x.to_numpy(), radius=1.5, color=colors[material.to_numpy()]) + gui.circles(x.to_numpy(), + radius=1.5, + palette=[0x068587, 0xED553B, 0xEEEEF0], + palette_indices=material) gui.show( ) # Change to gui.show(f'{frame:06d}.png') to write images to disk diff --git a/python/taichi/misc/gui.py b/python/taichi/misc/gui.py index 073f8c280..deb69e9b1 100644 --- a/python/taichi/misc/gui.py +++ b/python/taichi/misc/gui.py @@ -210,7 +210,12 @@ def set_image(self, img): def circle(self, pos, color=0xFFFFFF, radius=1): self.canvas.circle_single(pos[0], pos[1], color, radius) - def circles(self, pos, color=0xFFFFFF, radius=1): + def circles(self, + pos, + radius=1, + color=0xFFFFFF, + palette=None, + palette_indices=None): n = pos.shape[0] if len(pos.shape) == 3: assert pos.shape[2] == 1 @@ -235,6 +240,40 @@ def circles(self, pos, color=0xFFFFFF, radius=1): raise ValueError( 'Color must be an ndarray or int (e.g., 0x956333)') + if palette is not None: + assert palette_indices is not None, 'palette must be used together with palette_indices' + + from taichi.lang.expr import Expr + + if isinstance(palette_indices, Expr): + ind_int = palette_indices.to_numpy().astype(np.uint32) + elif isinstance(palette_indices, list) or isinstance( + palette_indices, np.ndarray): + ind_int = np.array(palette_indices).astype(np.uint32) + else: + try: + ind_int = np.array(palette_indices) + except: + raise TypeError( + 'palette_indices must be a type that can be converted to numpy.ndarray' + ) + + assert issubclass( + ind_int.dtype.type, + np.integer), 'palette_indices must be an integer array' + assert ind_int.shape == ( + n, + ), 'palette_indices must be in 1-d shape with shape (num_particles, )' + assert min( + ind_int + ) >= 0, 'the min of palette_indices must not be less than zero' + assert max(ind_int) < len( + palette + ), 'the max of palette_indices must not exceed the length of palette' + color_array = np.array(palette, dtype=np.uint32)[ind_int] + color_array = np.ascontiguousarray(color_array) + color_array = color_array.ctypes.data + if isinstance(radius, np.ndarray): assert radius.shape == (n, ) radius = np.ascontiguousarray(radius.astype(np.float32)) From 3297c6c8ec9775f34c3cd72df370eeef7a84161c Mon Sep 17 00:00:00 2001 From: ljcc0930 Date: Fri, 9 Jul 2021 20:55:57 +0800 Subject: [PATCH 03/42] [refactor] remove global_program in kernel.cpp (#2508) --- taichi/program/kernel.cpp | 40 +++++++++++++++++++-------------------- 1 file changed, 20 insertions(+), 20 deletions(-) diff --git a/taichi/program/kernel.cpp b/taichi/program/kernel.cpp index a5e3ab019..cf884cbc4 100644 --- a/taichi/program/kernel.cpp +++ b/taichi/program/kernel.cpp @@ -276,25 +276,25 @@ Context &Kernel::LaunchContextBuilder::get_context() { float64 Kernel::get_ret_float(int i) { auto dt = rets[i].dt->get_compute_type(); if (dt->is_primitive(PrimitiveTypeID::f32)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::f64)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i32)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i64)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i8)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i16)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u8)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u16)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u32)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u64)) { - return (float64)get_current_program().fetch_result(i); + return (float64)program->fetch_result(i); } else { TI_NOT_IMPLEMENTED } @@ -303,25 +303,25 @@ float64 Kernel::get_ret_float(int i) { int64 Kernel::get_ret_int(int i) { auto dt = rets[i].dt->get_compute_type(); if (dt->is_primitive(PrimitiveTypeID::i32)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i64)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i8)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::i16)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u8)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u16)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u32)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::u64)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::f32)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else if (dt->is_primitive(PrimitiveTypeID::f64)) { - return (int64)get_current_program().fetch_result(i); + return (int64)program->fetch_result(i); } else { TI_NOT_IMPLEMENTED } From e37bdb5e7c688004e6cd34bdcfc14aefdad65cf1 Mon Sep 17 00:00:00 2001 From: Jiasheng Zhang Date: Mon, 12 Jul 2021 15:44:46 +0800 Subject: [PATCH 04/42] [release] v0.7.26 (#2511) --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 8b494f769..5196b16c2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -8,7 +8,7 @@ project(taichi) SET(TI_VERSION_MAJOR 0) SET(TI_VERSION_MINOR 7) -SET(TI_VERSION_PATCH 25) +SET(TI_VERSION_PATCH 26) set(CMAKE_CXX_STANDARD 17) From 447177d7e3b81817ba91040c374c5a3cd102b834 Mon Sep 17 00:00:00 2001 From: squarefk Date: Mon, 12 Jul 2021 21:36:00 +0800 Subject: [PATCH 05/42] [wasm] Add wasm_set_kernel_parameter_* into runtime.cpp (#2510) * [wasm] Add wasm_set_kernel_parameter_* into runtime.cpp * Auto Format * Update codegen_wasm.cpp * fix nit * use anonymous namespace * Auto Format Co-authored-by: Taichi Gardener Co-authored-by: ljcc0930 --- taichi/backends/wasm/codegen_wasm.cpp | 26 +++++++++++++++++--------- taichi/runtime/llvm/runtime.cpp | 8 ++++++++ 2 files changed, 25 insertions(+), 9 deletions(-) diff --git a/taichi/backends/wasm/codegen_wasm.cpp b/taichi/backends/wasm/codegen_wasm.cpp index 1649f11e8..cf5d68224 100644 --- a/taichi/backends/wasm/codegen_wasm.cpp +++ b/taichi/backends/wasm/codegen_wasm.cpp @@ -13,6 +13,12 @@ namespace taichi { namespace lang { +namespace { +constexpr std::array kPreloadedFuncNames = { + "wasm_materialize", "wasm_set_kernel_parameter_i32", + "wasm_set_kernel_parameter_f32"}; +} + class CodeGenLLVMWASM : public CodeGenLLVM { public: using IRVisitor::visit; @@ -173,14 +179,16 @@ class CodeGenLLVMWASM : public CodeGenLLVM { ir->accept(this); finalize_taichi_kernel_function(); - auto wasm_materialize_name = "wasm_materialize"; - // compile_module_to_executable // only keep the current func TaichiLLVMContext::eliminate_unused_functions( - module.get(), [&](std::string func_name) { - return offloaded_task_name == func_name || - wasm_materialize_name == func_name; + module.get(), [offloaded_task_name](const std::string &func_name) { + for (auto &name : kPreloadedFuncNames) { + if (std::string(name) == func_name) { + return true; + } + } + return func_name == offloaded_task_name; }); tlctx->add_module(std::move(module)); auto kernel_symbol = tlctx->lookup_function_pointer(offloaded_task_name); @@ -199,9 +207,6 @@ FunctionType CodeGenWASM::codegen() { std::unique_ptr CodeGenWASM::modulegen( std::unique_ptr &&module) { - /* - TODO: move wasm_materialize to dump process in AOT. - */ bool init_flag = module == nullptr; std::vector name_list; @@ -211,8 +216,11 @@ std::unique_ptr CodeGenWASM::modulegen( gen->emit_to_module(); gen->finalize_taichi_kernel_function(); + // TODO: move the following functions to dump process in AOT. if (init_flag) { - name_list.emplace_back("wasm_materialize"); + for (auto &name : kPreloadedFuncNames) { + name_list.emplace_back(name); + } } gen->tlctx->jit->global_optimize_module(gen->module.get()); diff --git a/taichi/runtime/llvm/runtime.cpp b/taichi/runtime/llvm/runtime.cpp index 43cb5a1ef..cd79077ca 100644 --- a/taichi/runtime/llvm/runtime.cpp +++ b/taichi/runtime/llvm/runtime.cpp @@ -1751,6 +1751,14 @@ i32 wasm_materialize(Context *context) { (Ptr)((size_t)context->runtime->rand_states + sizeof(RandState)); return (i32)(size_t)context->runtime->roots[0]; } + +void wasm_set_kernel_parameter_i32(Context *context, int index, i32 value) { + *(i32 *)(&context->args[index]) = value; +} + +void wasm_set_kernel_parameter_f32(Context *context, int index, f32 value) { + *(f32 *)(&context->args[index]) = value; +} } #endif From 6235a5243afe5e63650b21b0a8597e2adb9bf9d4 Mon Sep 17 00:00:00 2001 From: ljcc0930 Date: Mon, 12 Jul 2021 21:36:35 +0800 Subject: [PATCH 06/42] [misc] Import template and ext_arr and kernel in python code (#2514) * remove ti.template() and ti.ext_arr() * remove kernel * delete line * Auto Format Co-authored-by: Taichi Gardener --- python/taichi/lang/meta.py | 60 +++++++++++++++---------------- python/taichi/lang/transformer.py | 8 ++--- python/taichi/torch_io.py | 11 +++--- 3 files changed, 40 insertions(+), 39 deletions(-) diff --git a/python/taichi/lang/meta.py b/python/taichi/lang/meta.py index 0a20bb365..73250fd41 100644 --- a/python/taichi/lang/meta.py +++ b/python/taichi/lang/meta.py @@ -1,26 +1,28 @@ from taichi.core import settings from taichi.lang import impl from taichi.lang.expr import Expr +from taichi.lang.kernel_arguments import ext_arr, template +from taichi.lang.kernel_impl import kernel import taichi as ti # A set of helper (meta)functions -@ti.kernel -def fill_tensor(tensor: ti.template(), val: ti.template()): +@kernel +def fill_tensor(tensor: template(), val: template()): for I in ti.grouped(tensor): tensor[I] = val -@ti.kernel -def tensor_to_ext_arr(tensor: ti.template(), arr: ti.ext_arr()): +@kernel +def tensor_to_ext_arr(tensor: template(), arr: ext_arr()): for I in ti.grouped(tensor): arr[I] = tensor[I] -@ti.kernel -def vector_to_fast_image(img: ti.template(), out: ti.ext_arr()): +@kernel +def vector_to_fast_image(img: template(), out: ext_arr()): # FIXME: Why is ``for i, j in img:`` slower than: for i, j in ti.ndrange(*img.shape): r, g, b = 0, 0, 0 @@ -45,8 +47,8 @@ def vector_to_fast_image(img: ti.template(), out: ti.ext_arr()): out[idx] = (b << 16) + (g << 8) + r + alpha -@ti.kernel -def tensor_to_image(tensor: ti.template(), arr: ti.ext_arr()): +@kernel +def tensor_to_image(tensor: template(), arr: ext_arr()): for I in ti.grouped(tensor): t = ti.cast(tensor[I], ti.f32) arr[I, 0] = t @@ -54,8 +56,8 @@ def tensor_to_image(tensor: ti.template(), arr: ti.ext_arr()): arr[I, 2] = t -@ti.kernel -def vector_to_image(mat: ti.template(), arr: ti.ext_arr()): +@kernel +def vector_to_image(mat: template(), arr: ext_arr()): for I in ti.grouped(mat): for p in ti.static(range(mat.n)): arr[I, p] = ti.cast(mat[I][p], ti.f32) @@ -63,21 +65,20 @@ def vector_to_image(mat: ti.template(), arr: ti.ext_arr()): arr[I, 2] = 0 -@ti.kernel -def tensor_to_tensor(tensor: ti.template(), other: ti.template()): +@kernel +def tensor_to_tensor(tensor: template(), other: template()): for I in ti.grouped(tensor): tensor[I] = other[I] -@ti.kernel -def ext_arr_to_tensor(arr: ti.ext_arr(), tensor: ti.template()): +@kernel +def ext_arr_to_tensor(arr: ext_arr(), tensor: template()): for I in ti.grouped(tensor): tensor[I] = arr[I] -@ti.kernel -def matrix_to_ext_arr(mat: ti.template(), arr: ti.ext_arr(), - as_vector: ti.template()): +@kernel +def matrix_to_ext_arr(mat: template(), arr: ext_arr(), as_vector: template()): for I in ti.grouped(mat): for p in ti.static(range(mat.n)): for q in ti.static(range(mat.m)): @@ -87,9 +88,8 @@ def matrix_to_ext_arr(mat: ti.template(), arr: ti.ext_arr(), arr[I, p, q] = mat[I][p, q] -@ti.kernel -def ext_arr_to_matrix(arr: ti.ext_arr(), mat: ti.template(), - as_vector: ti.template()): +@kernel +def ext_arr_to_matrix(arr: ext_arr(), mat: template(), as_vector: template()): for I in ti.grouped(mat): for p in ti.static(range(mat.n)): for q in ti.static(range(mat.m)): @@ -99,36 +99,36 @@ def ext_arr_to_matrix(arr: ti.ext_arr(), mat: ti.template(), mat[I][p, q] = arr[I, p, q] -@ti.kernel -def clear_gradients(vars: ti.template()): +@kernel +def clear_gradients(vars: template()): for I in ti.grouped(Expr(vars[0])): for s in ti.static(vars): Expr(s)[I] = 0 -@ti.kernel -def clear_loss(l: ti.template()): +@kernel +def clear_loss(l: template()): # Using SNode writers would result in a forced sync, therefore we wrap these # writes into a kernel. l[None] = 0 l.grad[None] = 1 -@ti.kernel -def fill_matrix(mat: ti.template(), vals: ti.template()): +@kernel +def fill_matrix(mat: template(), vals: template()): for I in ti.grouped(mat): for p in ti.static(range(mat.n)): for q in ti.static(range(mat.m)): mat[I][p, q] = vals[p][q] -@ti.kernel -def snode_deactivate(b: ti.template()): +@kernel +def snode_deactivate(b: template()): for I in ti.grouped(b): ti.deactivate(b, I) -@ti.kernel -def snode_deactivate_dynamic(b: ti.template()): +@kernel +def snode_deactivate_dynamic(b: template()): for I in ti.grouped(b.parent()): ti.deactivate(b, I) diff --git a/python/taichi/lang/transformer.py b/python/taichi/lang/transformer.py index a3a52b9ef..aebe0fbe4 100644 --- a/python/taichi/lang/transformer.py +++ b/python/taichi/lang/transformer.py @@ -4,6 +4,7 @@ from taichi.lang import impl from taichi.lang.ast_resolver import ASTResolver from taichi.lang.exception import TaichiSyntaxError +from taichi.lang.kernel_arguments import ext_arr, template from taichi.lang.util import to_taichi_type import taichi as ti @@ -679,9 +680,9 @@ def transform_as_kernel(): for i, arg in enumerate(args.args): # Directly pass in template arguments, # such as class instances ("self"), fields, SNodes, etc. - if isinstance(self.func.argument_annotations[i], ti.template): + if isinstance(self.func.argument_annotations[i], template): continue - if isinstance(self.func.argument_annotations[i], ti.ext_arr): + if isinstance(self.func.argument_annotations[i], ext_arr): arg_init = self.parse_stmt( 'x = ti.lang.kernel_arguments.decl_ext_arr_arg(0, 0)') arg_init.targets[0].id = arg.arg @@ -725,8 +726,7 @@ def transform_as_kernel(): for i, arg in enumerate(args.args): # Directly pass in template arguments, # such as class instances ("self"), fields, SNodes, etc. - if isinstance(self.func.argument_annotations[i], - ti.template): + if isinstance(self.func.argument_annotations[i], template): continue # Create a copy for non-template arguments, # so that they are passed by value. diff --git a/python/taichi/torch_io.py b/python/taichi/torch_io.py index 30794729a..8a800a151 100644 --- a/python/taichi/torch_io.py +++ b/python/taichi/torch_io.py @@ -1,14 +1,15 @@ -import taichi as ti +from taichi.lang.kernel_arguments import ext_arr, template +from taichi.lang.kernel_impl import kernel -@ti.kernel -def from_torch_template(expr: ti.template(), torch_tensor: ti.ext_arr()): +@kernel +def from_torch_template(expr: template(), torch_tensor: ext_arr()): for i in expr: expr[i] = torch_tensor[i] -@ti.kernel -def to_torch_template(expr: ti.template(), torch_tensor: ti.ext_arr()): +@kernel +def to_torch_template(expr: template(), torch_tensor: ext_arr()): for i in expr: torch_tensor[i] = expr[i] From d2a139a48c2bc8053f9f50c451ff24cdff921b12 Mon Sep 17 00:00:00 2001 From: xumingkuan Date: Tue, 13 Jul 2021 09:52:29 +0800 Subject: [PATCH 07/42] [misc] Output "taichi_cpp_tests.exe" to the "bin" folder on Windows (#2515) * [misc] Place "taichi_cpp_tests" in "bin" folder * restrict to windows and add docs * minor fix * Update docs/lang/articles/contribution/writing_cpp_tests.md Co-authored-by: Ye Kuang * Apply review Co-authored-by: Ye Kuang --- cmake/TaichiTests.cmake | 9 +++++++++ docs/lang/articles/contribution/writing_cpp_tests.md | 5 +++++ 2 files changed, 14 insertions(+) diff --git a/cmake/TaichiTests.cmake b/cmake/TaichiTests.cmake index 4404a1cda..d8f594474 100644 --- a/cmake/TaichiTests.cmake +++ b/cmake/TaichiTests.cmake @@ -23,6 +23,15 @@ include_directories( ) add_executable(${TESTS_NAME} ${TAICHI_TESTS_SOURCE}) +if (WIN32) + # Output the executable to bin/ instead of build/Debug/... + set(TESTS_OUTPUT_DIR "${CMAKE_CURRENT_SOURCE_DIR}/bin") + set_target_properties(${TESTS_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${TESTS_OUTPUT_DIR}) + set_target_properties(${TESTS_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY_DEBUG ${TESTS_OUTPUT_DIR}) + set_target_properties(${TESTS_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY_RELEASE ${TESTS_OUTPUT_DIR}) + set_target_properties(${TESTS_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY_MINSIZEREL ${TESTS_OUTPUT_DIR}) + set_target_properties(${TESTS_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY_RELWITHDEBINFO ${TESTS_OUTPUT_DIR}) +endif() target_link_libraries(${TESTS_NAME} taichi_isolated_core) target_link_libraries(${TESTS_NAME} gtest_main) diff --git a/docs/lang/articles/contribution/writing_cpp_tests.md b/docs/lang/articles/contribution/writing_cpp_tests.md index d41a42037..14792216d 100644 --- a/docs/lang/articles/contribution/writing_cpp_tests.md +++ b/docs/lang/articles/contribution/writing_cpp_tests.md @@ -21,6 +21,11 @@ make ./taichi_cpp_tests ``` +:::note +On Windows, `taichi_cpp_tests.exe` will be placed inside the `%TAICHI_REPO_DIR%\bin` directory, +so you can directly run it after adding the directory to `PATH` in [Setting up Taichi for development](dev_install.md#setting-up-taichi-for-development-3). +::: + :::note Consider polishing the CPP test infrastructure: From 286328cb916aef653d5981e020faa127b4d1e488 Mon Sep 17 00:00:00 2001 From: Dunfan Lu Date: Tue, 13 Jul 2021 20:12:15 +0800 Subject: [PATCH 08/42] [Bug] Warp reduction bug fix (#2519) * bug fix * update tests to cover the bug --- taichi/runtime/llvm/runtime.cpp | 26 +++++++++++++------------- tests/python/test_reduction.py | 19 ++++++++++++++----- 2 files changed, 27 insertions(+), 18 deletions(-) diff --git a/taichi/runtime/llvm/runtime.cpp b/taichi/runtime/llvm/runtime.cpp index cd79077ca..2888e7f19 100644 --- a/taichi/runtime/llvm/runtime.cpp +++ b/taichi/runtime/llvm/runtime.cpp @@ -1083,19 +1083,19 @@ i32 op_xor_i32(i32 a, i32 b) { return a ^ b; } -#define DEFINE_REDUCTION(op, dtype) \ - dtype warp_reduce_##op##_##dtype(dtype val) { \ - for (int offset = 16; offset > 0; offset /= 2) \ - val = op_##op##_##dtype( \ - val, cuda_shfl_down_sync_i32(0xFFFFFFFF, val, offset, 31)); \ - return val; \ - } \ - dtype reduce_##op##_##dtype(dtype *result, dtype val) { \ - dtype warp_result = warp_reduce_##op##_##dtype(val); \ - if ((thread_idx() & (warp_size() - 1)) == 0) { \ - atomic_##op##_##dtype(result, warp_result); \ - } \ - return val; \ +#define DEFINE_REDUCTION(op, dtype) \ + dtype warp_reduce_##op##_##dtype(dtype val) { \ + for (int offset = 16; offset > 0; offset /= 2) \ + val = op_##op##_##dtype( \ + val, cuda_shfl_down_sync_##dtype(0xFFFFFFFF, val, offset, 31)); \ + return val; \ + } \ + dtype reduce_##op##_##dtype(dtype *result, dtype val) { \ + dtype warp_result = warp_reduce_##op##_##dtype(val); \ + if ((thread_idx() & (warp_size() - 1)) == 0) { \ + atomic_##op##_##dtype(result, warp_result); \ + } \ + return val; \ } DEFINE_REDUCTION(add, i32); diff --git a/tests/python/test_reduction.py b/tests/python/test_reduction.py index f390cbe6e..0b859665b 100644 --- a/tests/python/test_reduction.py +++ b/tests/python/test_reduction.py @@ -39,10 +39,18 @@ def _test_reduction_single(dtype, criterion, op): a = ti.field(dtype, shape=N) tot = ti.field(dtype, shape=()) - @ti.kernel - def fill(): - for i in a: - a[i] = i + if dtype in [ti.f32, ti.f64]: + + @ti.kernel + def fill(): + for i in a: + a[i] = i + 0.5 + else: + + @ti.kernel + def fill(): + for i in a: + a[i] = i ti_op = ti_ops[op] @@ -62,7 +70,8 @@ def reduce_tmp() -> dtype: reduce() tot2 = reduce_tmp() - ground_truth = np_ops[op](a.to_numpy()) + np_arr = np.append(a.to_numpy(), [0]) + ground_truth = np_ops[op](np_arr) assert criterion(tot[None], ground_truth) assert criterion(tot2, ground_truth) From b08bebfdaec0298548650fb594e5e670ee5f5118 Mon Sep 17 00:00:00 2001 From: Ailing Date: Wed, 14 Jul 2021 03:54:05 -0700 Subject: [PATCH 09/42] [Refactor]Move llvm runtime build for archs to cmake from C++. (#2524) Co-authored-by: Ailing Zhang --- CMakeLists.txt | 25 +++++++++++++++++++++++++ cmake/TaichiCXXFlags.cmake | 3 +++ python/build.py | 1 - taichi/llvm/llvm_context.cpp | 7 ------- taichi/python/export_lang.cpp | 2 -- 5 files changed, 28 insertions(+), 10 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 5196b16c2..878f9b666 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -76,6 +76,31 @@ add_custom_target( ) add_dependencies(${CORE_LIBRARY_NAME} generate_commit_hash) +if (TI_WITH_CUDA) + set(CUDA_ARCH "cuda") +endif() + +find_program(CLANG_EXECUTABLE NAMES clang-7 clang-8 clang-9 clang-10 clang) +if (NOT CLANG_EXECUTABLE) + message(FATAL_ERROR "Cannot find any clang executable.") +endif() + +find_program(LLVM_AS_EXECUTABLE NAMES llvm-as) +if (NOT LLVM_AS_EXECUTABLE) + message(FATAL_ERROR "Cannot find llvm-as executable.") +endif() + +# Build llvm-runtime for host arch and cuda (if available) +foreach(arch IN LISTS HOST_ARCH CUDA_ARCH) + add_custom_target( + "generate_llvm_runtime_${arch}" + COMMAND ${CLANG_EXECUTABLE} -S runtime.cpp -o runtime.ll -fno-exceptions -emit-llvm -std=c++17 -D "ARCH_${arch}" -I ${PROJECT_SOURCE_DIR}; + COMMAND ${LLVM_AS_EXECUTABLE} runtime.ll -o "runtime_${arch}.bc" + WORKING_DIRECTORY "${PROJECT_SOURCE_DIR}/taichi/runtime/llvm" + ) + add_dependencies(${CORE_LIBRARY_NAME} "generate_llvm_runtime_${arch}") +endforeach() + FILE(WRITE ${CMAKE_CURRENT_LIST_DIR}/taichi/common/version.h "#pragma once\n" "#define TI_VERSION_MAJOR \"${TI_VERSION_MAJOR}\"\n" diff --git a/cmake/TaichiCXXFlags.cmake b/cmake/TaichiCXXFlags.cmake index e6f629163..030d58c6a 100644 --- a/cmake/TaichiCXXFlags.cmake +++ b/cmake/TaichiCXXFlags.cmake @@ -44,11 +44,14 @@ if ("${CMAKE_SYSTEM_PROCESSOR}" STREQUAL "x86_64" OR "${CMAKE_SYSTEM_PROCESSOR}" message("Setting -march=nehalem for x86_64 processors") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -march=nehalem -DTI_ARCH_x64") endif() + set(ARCH "x64") elseif ("${CMAKE_SYSTEM_PROCESSOR}" STREQUAL "aarch64" OR "${CMAKE_SYSTEM_PROCESSOR}" STREQUAL "arm64") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DTI_ARCH_ARM") + set(ARCH "arm64") else() message(FATAL_ERROR "Unknown processor type ${CMAKE_SYSTEM_PROCESSOR}") endif() +set(HOST_ARCH ${ARCH} CACHE INTERNAL "Host arch") if (USE_STDCPP) set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -stdlib=libc++") diff --git a/python/build.py b/python/build.py index ffcd10a40..bb78811d6 100644 --- a/python/build.py +++ b/python/build.py @@ -88,7 +88,6 @@ def build(project_name): assert os.path.exists(libdevice_path) shutil.copy(libdevice_path, 'taichi/lib/slim_libdevice.10.bc') - ti.core.compile_runtimes() runtime_dir = ti.core.get_runtime_dir() for f in os.listdir(runtime_dir): if f.startswith('runtime_') and f.endswith('.bc'): diff --git a/taichi/llvm/llvm_context.cpp b/taichi/llvm/llvm_context.cpp index 19c46efc9..9597a2523 100644 --- a/taichi/llvm/llvm_context.cpp +++ b/taichi/llvm/llvm_context.cpp @@ -230,13 +230,6 @@ void compile_runtime_bitcode(Arch arch) { } } -void compile_runtimes() { - compile_runtime_bitcode(host_arch()); -#if defined(TI_WITH_CUDA) - compile_runtime_bitcode(Arch::cuda); -#endif -} - std::string libdevice_path() { std::string folder; if (is_release()) { diff --git a/taichi/python/export_lang.cpp b/taichi/python/export_lang.cpp index 8a3e98f4c..06b28bfb3 100644 --- a/taichi/python/export_lang.cpp +++ b/taichi/python/export_lang.cpp @@ -55,7 +55,6 @@ void expr_assign(const Expr &lhs_, const Expr &rhs, std::string tb) { std::vector> scope_stack; -void compile_runtimes(); std::string libdevice_path(); std::string get_runtime_dir(); @@ -743,7 +742,6 @@ void export_lang(py::module &m) { m.def("test_throw", [] { throw IRModified(); }); m.def("needs_grad", needs_grad); - m.def("compile_runtimes", compile_runtimes); m.def("libdevice_path", libdevice_path); m.def("host_arch", host_arch); From 6cd41cd5fed0a115a0db842e18a18111345b764f Mon Sep 17 00:00:00 2001 From: xumingkuan Date: Wed, 14 Jul 2021 20:05:56 +0800 Subject: [PATCH 10/42] [bug] Set use_unified_memory=True for tests with pointer fields with 32-bit data types (#2521) * [windows] [bug] Temporarily disable tests with pointer fields with 32-bit data types on CUDA on Windows * Fix issue number * Auto Format * Revert "Auto Format" This reverts commit 6cb4d601 * Restore tests with use_unified_memory=True, add a test for this issue Co-authored-by: Taichi Gardener --- tests/python/test_fields_builder.py | 2 +- tests/python/test_loop_unique.py | 8 ++++---- tests/python/test_sparse_basics.py | 18 ++++++++++++++++++ 3 files changed, 23 insertions(+), 5 deletions(-) diff --git a/tests/python/test_fields_builder.py b/tests/python/test_fields_builder.py index 04f2c351b..a4edd92a7 100644 --- a/tests/python/test_fields_builder.py +++ b/tests/python/test_fields_builder.py @@ -85,7 +85,7 @@ def func2(): assert x[i] == i * 3 -@ti.test(arch=[ti.cpu, ti.cuda]) +@ti.test(arch=[ti.cpu, ti.cuda], use_unified_memory=True) def test_fields_builder_pointer(): n = 5 diff --git a/tests/python/test_loop_unique.py b/tests/python/test_loop_unique.py index c4a33d15d..7ee0cce48 100644 --- a/tests/python/test_loop_unique.py +++ b/tests/python/test_loop_unique.py @@ -1,7 +1,7 @@ import taichi as ti -@ti.test(require=ti.extension.sparse) +@ti.test(require=ti.extension.sparse, use_unified_memory=True) def test_loop_unique_simple_1d(): x, y = ti.field(ti.i32), ti.field(ti.i32) @@ -26,7 +26,7 @@ def inc_y(): assert y[i] == expected_result.get(i, 0) -@ti.test(require=ti.extension.sparse) +@ti.test(require=ti.extension.sparse, use_unified_memory=True) def test_loop_unique_binary_op_1d(): x, y = ti.field(ti.i32), ti.field(ti.i32) @@ -51,7 +51,7 @@ def inc_y(): assert y[i] == expected_result.get(i, 0) -@ti.test(require=ti.extension.sparse) +@ti.test(require=ti.extension.sparse, use_unified_memory=True) def test_loop_unique_nested_1d(): x, y = ti.field(ti.i32), ti.field(ti.i32) @@ -77,7 +77,7 @@ def inc_y(): assert y[i] == expected_result.get(i, 0) -@ti.test(require=ti.extension.sparse) +@ti.test(require=ti.extension.sparse, use_unified_memory=True) def test_loop_unique_2d(): x, y, z = ti.field(ti.i32), ti.field(ti.i32), ti.field(ti.i32) diff --git a/tests/python/test_sparse_basics.py b/tests/python/test_sparse_basics.py index 2cece4351..c6b27d73b 100644 --- a/tests/python/test_sparse_basics.py +++ b/tests/python/test_sparse_basics.py @@ -1,3 +1,5 @@ +import pytest + import taichi as ti @@ -75,3 +77,19 @@ def func(): assert s[None] == 5 * n print(x[257 + n * n * 7]) assert s[None] == 5 * n + + +@pytest.mark.skip(reason='https://github.com/taichi-dev/taichi/issues/2520') +@ti.test(require=ti.extension.sparse, use_unified_memory=False) +def test_pointer_direct_place(): + x, y = ti.field(ti.i32), ti.field(ti.i32) + + N = 1 + ti.root.pointer(ti.i, N).place(x) + ti.root.pointer(ti.i, N).place(y) + + @ti.kernel + def foo(): + pass + + foo() From c25c76772471aea540095679de190c56cd1feeb8 Mon Sep 17 00:00:00 2001 From: xumingkuan Date: Thu, 15 Jul 2021 13:03:11 +0800 Subject: [PATCH 11/42] [windows] Use %lld as 64-bit integer output format on Windows (#2525) --- taichi/ir/type_utils.cpp | 10 ++-------- 1 file changed, 2 insertions(+), 8 deletions(-) diff --git a/taichi/ir/type_utils.cpp b/taichi/ir/type_utils.cpp index dffe9ff99..7ba723994 100644 --- a/taichi/ir/type_utils.cpp +++ b/taichi/ir/type_utils.cpp @@ -25,17 +25,11 @@ std::string data_type_format(DataType dt) { } else if (dt->is_primitive(PrimitiveTypeID::u32)) { return "%u"; } else if (dt->is_primitive(PrimitiveTypeID::i64)) { -#if defined(TI_PLATFORM_UNIX) + // Use %lld on Windows. + // Discussion: https://github.com/taichi-dev/taichi/issues/2522 return "%lld"; -#else - return "%I64d"; -#endif } else if (dt->is_primitive(PrimitiveTypeID::u64)) { -#if defined(TI_PLATFORM_UNIX) return "%llu"; -#else - return "%I64u"; -#endif } else if (dt->is_primitive(PrimitiveTypeID::f32)) { return "%f"; } else if (dt->is_primitive(PrimitiveTypeID::f64)) { From e44a3ec2049e7a121df62182ebb0b67eeb757f4e Mon Sep 17 00:00:00 2001 From: "Chengchen(Rex) Wang" <14366016+rexwangcc@users.noreply.github.com> Date: Fri, 16 Jul 2021 10:41:36 -0400 Subject: [PATCH 12/42] [doc] Update links in the readme file. (#2523) * [skip ci] Add a formal security disclosure section. * [skip ci] Update links to docs and zh-Hans docs, also fix other broken links. * [skip ci] Not sure what to do with the hub link, hide for now. * [skip ci] Use formal words. * [skip ci] Also update the PR template. * Loosen the gitignore restrictions for docs folder. * Add back the missing metadata and image for docs. Fixing https://github.com/taichi-dev/docs.taichi.graphics/pull/71. --- .github/pull_request_template.md | 6 +++--- .gitignore | 3 +++ README.md | 12 ++++++++---- docs/lang/api/reference/_category_.json | 4 ++++ docs/lang/articles/advanced/_category_.json | 4 ++++ docs/lang/articles/basic/_category_.json | 4 ++++ .../lang/articles/contribution/_category_.json | 4 ++++ .../contribution/life_of_kernel_lowres.jpg | Bin 0 -> 122978 bytes docs/lang/articles/misc/_category_.json | 4 ++++ 9 files changed, 34 insertions(+), 7 deletions(-) create mode 100644 docs/lang/api/reference/_category_.json create mode 100644 docs/lang/articles/advanced/_category_.json create mode 100644 docs/lang/articles/basic/_category_.json create mode 100644 docs/lang/articles/contribution/_category_.json create mode 100644 docs/lang/articles/contribution/life_of_kernel_lowres.jpg create mode 100644 docs/lang/articles/misc/_category_.json diff --git a/.github/pull_request_template.md b/.github/pull_request_template.md index ccd420f17..ce564e2b0 100644 --- a/.github/pull_request_template.md +++ b/.github/pull_request_template.md @@ -3,12 +3,12 @@ Related issue = # +

Documentation | 简体中文文档 | Contributor Guidelines

[![AppVeyor Status](https://img.shields.io/appveyor/build/yuanming-hu/taichi?logo=AppVeyor&label=AppVeyor)](https://ci.appveyor.com/project/yuanming-hu/taichi/branch/master) @@ -13,9 +13,9 @@ **Taichi** (太极) is a programming language designed for *high-performance computer graphics*. It is deeply embedded in **Python**, and its **just-in-time compiler** offloads compute-intensive tasks to multi-core CPUs and massively parallel GPUs. - + -Advanced features of Taichi include [spatially sparse computing](https://taichi.readthedocs.io/en/latest/sparse.html) and [differentiable programming](https://taichi.readthedocs.io/en/latest/differentiable_programming.html) [[examples]](https://github.com/yuanming-hu/difftaichi). +Advanced features of Taichi include [spatially sparse computing](https://docs.taichi.graphics/docs/lang/articles/advanced/sparse) and [differentiable programming](https://docs.taichi.graphics/docs/lang/articles/advanced/differentiable_programming) [[examples]](https://github.com/yuanming-hu/difftaichi). **Please check out our SIGGRAPH 2020 course on Taichi basics:** [YouTube](https://youtu.be/Y0-76n3aZFA), [Bilibili](https://www.bilibili.com/video/BV1kA411n7jk/), [slides (pdf)](https://yuanming.taichi.graphics/publication/2020-taichi-tutorial/taichi-tutorial.pdf). @@ -67,3 +67,7 @@ If you use Taichi in your research, please cite our papers: - [Taichi Conference](https://github.com/taichi-dev/taichicon): Taichi developer conferences. - [GAMES 201 Lectures](https://github.com/taichi-dev/games201): (Chinese) A hands-on course on building advanced physics engines, based on Taichi. - [More...](misc/links.md) + +## Security + +Please disclose security issues responsibly by contacting contact@taichi.graphics. diff --git a/docs/lang/api/reference/_category_.json b/docs/lang/api/reference/_category_.json new file mode 100644 index 000000000..ec4123c51 --- /dev/null +++ b/docs/lang/api/reference/_category_.json @@ -0,0 +1,4 @@ +{ + "label": "API Reference", + "position": 99 +} diff --git a/docs/lang/articles/advanced/_category_.json b/docs/lang/articles/advanced/_category_.json new file mode 100644 index 000000000..84fc16f93 --- /dev/null +++ b/docs/lang/articles/advanced/_category_.json @@ -0,0 +1,4 @@ +{ + "label": "Advanced Programming", + "position": 3 +} diff --git a/docs/lang/articles/basic/_category_.json b/docs/lang/articles/basic/_category_.json new file mode 100644 index 000000000..60b7ef558 --- /dev/null +++ b/docs/lang/articles/basic/_category_.json @@ -0,0 +1,4 @@ +{ + "label": "Taichi Language Basic Concepts", + "position": 2 +} diff --git a/docs/lang/articles/contribution/_category_.json b/docs/lang/articles/contribution/_category_.json new file mode 100644 index 000000000..74136a441 --- /dev/null +++ b/docs/lang/articles/contribution/_category_.json @@ -0,0 +1,4 @@ +{ + "label": "Contribution Guide", + "position": 5 +} diff --git a/docs/lang/articles/contribution/life_of_kernel_lowres.jpg b/docs/lang/articles/contribution/life_of_kernel_lowres.jpg new file mode 100644 index 0000000000000000000000000000000000000000..0f7dbc3322f942ef9f3bfe307525fdcd781ecd2f GIT binary patch literal 122978 zcmeFYcUaR~voIR0bfgOiN+5J_6Pk2|5J->~nuMa#6Pkda6tU5JLI}MlK&VQux)rGb zfe?C=s#K+mC|urs&i35*KIc8(ckc7u``1~?Gr!C#vnFfJn%|n4wNA!PCIO)U+6(_k zr;wJB_987KGd(RWJ@X|dCT1q4ORS8ikeQK{g^itqosH%4WnSLPmxX{pppfYAL_G%)rFJ%*uS4SlRy~zX-+->-U4W70B6pgg5T-l z`3p4XX#bISnttd0%fARcjq65&JIeYH>g^R3g?1CIXSz1mm zZm`lV6NI_NeeZ`!I12T$kRqn4x@PExmYLu4P%Po~pqzrKrM=GwVR_%YciIjw!rp@% zoeJ7SAkHt7^NEYKqGG!3bF6=H1vq>9IdjU&Y3DR-XIN!T1aI9xtMwf0%??Doz~r4w z02t4n7O|dT1!w}^KKlog|9go=Q{;PR%2hvESN#Ai5<3)cIx;RU>w;xp{vp(DXvTBB|>CqO)n5TZ=;;KS1cY8};8*y~j z1KmJL^Ro{oydtiyHGqo{2t>JQZt@bik z#%dpT$5R$0l-e_7ykqW~r{kPbS#I8HDCSfaWzH~whdB<1EC$d1{_9Md>=nGkQs=&k!9KbaPIsqTXw8cEqI}s< znm>o|zKEP%y~3gdSB$%ympj~x#LT4-%53615>-X5zXFCMY-`1>SXo|QY#D3t_TL~)n2Gmk=S`-YN=+kbTE%)P z1@ekew*dkPls!*z>A5GxF%zvpRJh3|ISeO=vOIM`-KwiT$53WGy(m% zwg6*pq5;vyCQ}Y3Ngk*1p;?X|0nfKLieD;C1ir^EIBxR>szgD@(-YxRS&Mxl+O>M^ zjHrgZiYIK?4`am-#;tpG6j*Oy;c=?dOF$FAt}7k>vd!ZNWHGg1<{S6H`Wm$#YLPQgGokWzc?UHQL2a?vsUHd>wG$}D=g@)cE z_j>L1%LI%R$;IA!8kFt}&N&oShn~xt9=5%+uY{}V%|ryHgSP{gM<{d!aBN{pSyEcr z0b5Ye5l>3#p8&wcIj|!<*hR$XNyDI()bx`jHojXfcr+q5vOcx5v#C7Fhc}9?_5viO zUtdZNl)_a_R>R_?dJitO&0`f)gB}#svZ8+_fu0To6kOz+us1)iUL;_--syLDmuCGQ zfacpyZwJb&%Uc>ej<${$!{Smr;1&?=ohU{ATiB#C+K5z>MWa7z8eW1o#4k9CTt3um zK<^Xm^X1o5w-z_hvUh|K)bVS(QUh75aL%5)MX#H!m;Li?a0ewrE`RL8wKwvZ>VBBL z@YgcW9;~U9$>YJ4cFs?6=bow~miA23W~q#Fj*?p)H#{NeTcG_oXEDBeQ;hm8OA z>++)tGE%NzH_7@Ld5h}$B*h6FmIqscQ-~N$VlNB#e7fS}um?r;((qsTpr_P}Rrx(} zDyoRNmMGS2m~_g7i8KRjTU^4>(a~X_cjl^6)oYIrPrr>79T|Y`?(tpwfAsb*(364x zj*8x9R`TmDI|Tkp#;naPUh!N@v)a})Eaf4tVU;@t7peEaX9WGVF{{vwYi@AdA&d&lLu)y@<8f6(>5K`|Zo@w= z9K@x%c$i3!Y^sI;QD1cf>Kv;EOC1+Hm$4SfJd1=<>7}uLjXxQ5>U{@gCeY{_agc~? zTAmYz=us9KJLye7lg#yWPr9uYK{EaFExUMFxSoygJZW^=1SEca`N@m*1QlN;GXc-V z(jhwCR~TA?IDeeGN~rL$3r^+dz$3Htu-=&aMU1%!D(-};DmvSMX zJKSSaWRFK9ln}F}Z)fFS57>XeOqB3JKWs1;#^ZY9si8$RhPj#K6Tk%jD_#LPLUMRA zMdQ6gP9F&KIhMz;Oh1ns<(1RNvmg-FVOUrD^}GA8yweiSwKVPR?jy-R)n90egZZL( zRt!H^->3Hr)R&{YP!P}=SBnfO_?G^{=Srx+wy8cWC0XRA=+^5_mZg>tuG<9ZLk4Yw z+IN{k)LJ%#)l!Q1en*@CY~6IyyXd!q?iKd7W(aM4U#sdO6Xyx@6bqYzF0Md2a^16+ zGw|Xzq3f(~o^SBIxS}jiiigDyt)b;QCbG0FQ1e$!ixXsS3Rtu;-ZUr=RKHVpqk_g< z)Wj5taFfr7o&-_gzH(sSW2nQBKk=*QOk9uuY+@sAI6-j7+RaVKil+$p0P8|c-8Q_E zSMSvvP>)E?^G$^(<_p0v#q~Uw_Y4j!6CB2C1BYe93*h>WQ%};_GH3lDWH&I(vHU`- z_0dgDhI-ERH5IK~DoDoQ7Qc-cO4J4#)RieGpLg+X;&AG_+W7r8Q@$wcLh6@h1W0jx zY!Nn4FjWobA}%cmr$&9+Q*JsIy6ImSd@WU}ge9Z1cw_;kAP86R$sUGTY}qYG6zQb7 z=#cP9NoVZo7^Gsm>nQ~3t3=4lUjWHJ3O^KoaDI4kP(&&eJN|^omROVHB%*zA$r7u^kKUyI*Rril^?C0#0 z$g&DZiZPip&GYcVFeyLdN+x-n#E~Be=w_H%ecV zH{mhkcK~3ATmBJQ<>_+A@td2uzAC{HX@`AZN%5jT2WO>LW;6JH9G-dJyRt3d@>%Ik z;SG9Q&$rNeL5Vz@bgEc;1AgVJ$$=QmnOOP(E5|M#Jm_!A;=Z6=40lfm>h>TA|3r!M z!{7J_Hcnp-@S5Qa^guL)b?GQD3$oHd@l#41VXVzXAuvi%jMUC-^Mf^;v7r06>zAM@ zw_h6EFEn$_9x~`O^94@Lsl0$~83kanb;oJOtG9SZTDv_49gB%IS)T6_vocno^CPf6 z7{XV5IY}+tr~Xcu7lmkn7O9G2;?wP`NvFbkV|?^eaF0h>O4!67wVyF_O1ErFb+Yp8 ztH1Kr!#$8U9dgR-F& zU>pQub!qlu!XNj~ErSLkkLVLC$78o*j^*qSkVGRdu+J6Yq>5|gEgO?`#MV>uwUJ$iKE0z$f-610n9q0W zrs$C5QykzP2HE*S@?zxZc~$tWuR=5>wo3eLt8J3e1-DY8Fpn8l5STPxE`!y98mTCl z1=Kn=BIlLyKqkDP+!?0DE{-{B+yjdV_Vvdz217 z+yRv9Rxxd9qz$0FFHrTit(jo=YUFQRGX-Zl?pD5ZMJ!%R_-O#rnV|WC{-`;&+lY?h z1eatepll+Gp`DLp!kTJ$TL@NH-91K5mBFm+bYS?8xDoIi>s7pW>{o z^TA!Hx?nt+j*1MtXGqjF!`q?U)KQMLHzHtYi&j0MS#plcD<3KLV$HLxE8-}qUaFDK z6VUfAQ@9oDAy;VWE?{?KJkYqTc(Bki)4LFK!@aI<(lpZtRv8PMv5xffxUExn9p)ow zEa3MUSZAb0tg55H5H(zT7_VN}pMzh7s!<>PO%d)fyRS|F+r0g9TC}U&^hd(Ij+jHU zQ&l-|G5;puVrpXnRwrO6dD@ik1dz0SVd+`&W%*#$GIO+%h}w*z6@A4tH?`QIG)eoH(vy~t#M`+MR@}ZhdHImIGm?Mg$>V{YTFCqO zE^MeE2#dhVyt{<|q&{s_Hu0q%G3i|*@3`W9T4-2BFLt8W>t2*t04Z?__A%XNps4V ze%Z4hB(rT79vNoO_ql_Hp+yQgoQ#4{d^c!pc=&2;Z{VLBCH7LaNILysj@&EKR}ez# z*;z%0wOczIUuljeM&nuow9F>XRVCtku}2a@+T~)0bREMj)$3OU_Zj^vCby@n46j0% zn|fS4=)&qKSyB^&r53T<$B9(}Ht*5986m$2(!_MD%pceEy|yf-!q}Zkf0oJ@w>V0) zP`ZSOlWXgRkWkE_@yioHF1@-aTRD9@?~``|sA~PfvW#RAEsjAp9K;bOBn_9NLLfS6 zsq^FyNdE78I!!`g#I&kYO^6m-Xr)J`VQ&=L@>kk6hDB*iv9B}o(v=!gMLxfPa1zN3 z)Bw$&Xc+&a=V4J`GLfEQb)cKtIv2tRYpXS%^v{icsc<&lyk6rJzaytG~3m%(k|;23B&t!4O91|i?l5>dEvL)JxR##6zQx6qs@Z%CDh@eaPW z*nrGLa%byc6TF>{-4B-e$(onT8Q!DCy>Yv|FB{h#df}!u;8|jA9g)!g!5J!c_Py_& zB#$XJp{f;J38cMlvb|v+lO<1dIsruSBv{!yO#>%i30vHkbW*YBO<;2Bi4f-6E9wYE zn>_Bi$ZaTNRY#~bb<|i~#je$$7I=?%ZWwM zz%aeCT&#}}*-YEYsdzG=RvEK@JL>l4hRXttdxBt;yOM|`DIfxCD8GgV38GA#lR>a4 zM@Kc^m@vH1)alkIB0E*_xgj1Ad$c^S{)Nr<&)GC5ga0iP?t-Gg z<22;f1N&dYn6<4XMzh+Al8S4VqBrx6l6g9r?I-f7B+o5{_c*+sg82|=6vNx$km}w)cICseP$RuHE|;SL6lLsQ#uzU zXpcEu)wIMPl){qVENL(H~Qxf20cq&tnQ{IUJlHz$yBhngL7;4_Z0(RvF>{IJ z_y;KC*(}gXELVpd+RqV(9YM?2K?At5joD^g&bd4`%iKhjUeDUsOdOF-mmAy?sjn51 zE=lvN*687A3vl{*Vdt3{LwJn8kxwNo?R*4jO-mq&za!C2G|fa9H;=GNb7}B=+!;o> zu4q+05T$k$K3-GD;fTd{^N_?q0`o}5e#4*ywb(mMVJzg zRUQm%!s=`jfnSj619J(_)DsHD_p|DL$*9I(BaVE=*MLXzC&X*S#k7)()M89NPtcea zy-U!4&XsW~PByP<%Ac)3fk%-7t6?oZ$ev%~PXzFa$O+s*+4Wb2GknH&?Q}2*+)2|u z0mvs*wWu@2VnqAfRx^1WU7X=hys_azbwxH{ai(%B!E-JRShW@O;MXI=x6Gd+-1}0IF;*tcR)w2g90*xEK*>eHH5N^oF3 zYFJ5dlbTj>xPCxd&NUYPxt^NYy15g;7H#`Z14Nxaerx(Evx{`ToXFMKwX{d7T2PlC z7Fu*JNbXREx|fPNXYz*_&44%7N+Y)sIM5<5j>roJ0+(dh(#yg8=nR=8V<<&<1mE^$ zQC(766tmCRsCRVXY}XCbTX|EXv4Nl`^foA|)++cF#olQxMPi5HF;Qab`Y*v(Yum~K zRiAfQ`68IF2dh@pn-c%n?fNz&`pMj0(UTRW4x3cOn&ckQgSe75(&){VuA&h_f)O5c zJ`#;6#d+R+x4YjF9@41(+K%mKdB{{RH!I((S-0+wuKNM^bJPItU zL~Mt}*w85t9Zj=wn-5TWz>lhQFjl^pnF8c+TH007n8dngm-X9I2uSl~q0e?`_6b10OK@Wd2<|$-Y$xtqxOQ)^&BMZ`VnQ}ELx191 z>PCmOSG$R%tKK*l8ga>~-NdCir9&V`K=AEVNa@7wH-<6%%E7;ncOBm%{0&Dg|kMMN6VM%u7ygQE-V*X0UnX;K>>H^z%CS%f%4uiVC4_ ztxHBi2wVXa0aVG$- zHxiLC>XTJ8mG*sX$)0wUFT1Wwyd2IB>>g~mg1+0QFF{;2BDS|*c9rplv?6E2-JcrD zKQ}N5H|Dd5LDhG@WWQxcH@#;&>SUGhE-nR*7xGdwzxUPQ@WIjRJd5^gp-K$3ZU`pX zA7*3B2nRVHgm8^({Qiq^Oy`4s1p!!}PSFE;+Ec6UqP+L#z&~Y%Wh)QgUNctC9L~bX zH=8SkJzTa~1Jh)u1uH?b$H^-X~ zKl7E>ixaDepxbp>Fi|??luH-*xV_^`Z|)kZHUx&hUCdmF|@VO#-@4tgDv7n-?N zQheS=uJ_pLYnS+i@`aO0TE;ahuwDujMTG{{FZF2Kzp5n9lzT#|ig(X{8?h48t5k7= zxURFh>VG5*V#;fLxZcR7h~7IFI0&~48o*+Kuwi>vRv)MLtX=BfkwwW7J|Uh-duRbm z^S2M~JR)di7=~!$NPQY3eYIb(lgvp`))GLqZj{lN^{)acO=z&^ZEz|1#JT`eO+Rwpfww9 zwW*O%kSE5oUua3D+xe#8RG$y8j!z>i6teRg=|{JoZvB8#dDwvYvbCsH=gm#OsaP_& zve{4^1BTS`nQmS6urU{O9`}TkUzJP9*xAk#8Tdy!0zcKF)%O%S6I94?0kC_ zxSg^>^L8FlGsewx0@!3p`Is4z@du>qH8?KEk;ocY`hL$__F>=q=nTI?L!EbE3F^97 z-0Fara=JrZcx`%}D;SpT?1B%7tl*d!H5a|)np89>+8!OjuV2~9xjYGzjxa{zYgoDV z0K10QvNO0AHIMnOuUe=zzbqb&GCQ-OvijCpwd8a+zIzN*0v!po3h*(*5`QZ3pu!Pu z-tY%CjyU1i{t|)5@0%@FiT3|Kgw}rVy-$a55(TQ7mqU!a0M6uc9J>z5yttR*SvX47 z<9#c4>jY3|bGk|OV_#8IoO)U~T(Un5<^(cEZGM~g>zJ>p$Kmo)Ow*&POfb%`D}6Og zB&l$sheDuD11*L>$-^hr8&^C~>j`VXEeAi~eg&Xe)^OcpoiU^Clgdv3eFvSsEK+rq zWkRK*7+!bPN%CO}@9K}$A%>_WR5_4dvVWD-C@hgiN0@rP zL|6r*=QUt)*S6uda`K*9WZ}x)tM_(jGuBIlM}>M|tM`W$e<*XNa}sryrqqo)F3i>7 ztrL5#=uwNH4k{wEi;_>(B~Qo}%x|)o=aj;lr1d(b)1sA%h9-`aS*}{F>HU1z5bOY4 zK=;4y9zB9vb08jnX$FE(O+Xvu<|Pd(r5>d|`23#T^}b zC!r9cYkVZ)a|H~l%SYbs>JQoGO^_#@0M-NPo_@Xl1PL8S*`~BESvw<&qjQm4l-eee zBN`*^oTTjvt*Jx7#P)cIqTTrN1)Jqi!bqvCm~8Rxio#iKtj**6s`?r(JT@yxzEo;Q z?U`Rddevuxr~b#6I$TZw&+V6iE&XQc!@SZV3ieqN*={~^vEK?UKE9D8S;Xa=W?ONE zn4oFqnhE6m)xmw`C9xC0;hMmL5yDPOK(=E10N1o+=w`78Gia4QNUUiK?&>xqK^z?u zF7MygA$65WqvS~jd2o1)kdQnkjR&yBzg}*c{62r3y8I-tGxznpWJ#Org>iw%s`>J% zMmkk)q>$!MeL>cSa)u7SkaFeHKtdhiHg=oKcor;QK3ooIAI6Mg$ZS~KGF<86PNL80?2nq=-FnO2-i^tH}B*R-ecV!?eakuBjeje0b5L3 zD-=HFw?cnyeOZd{92{SsVjRxX@xP_@2*p;!S*y}B)T*q4b9Nm`k|#-!t1&rZMejOG za3h!6xE^*Zm~&;t`iUEmf-|a6`oJy65MvH~eU|c4P`BX#k2kRw&UQRFtemEZmS3t& zDv@34wYDgOHn5xk9>aqBu0@o%cUCk$Og3zu^in~D2b8NCt$I>ic^YMN#3IaaJUpH>E6pI}`ecC4c%k->P{uE# z?ji(PCu1Qp%3aU-oPU-RHJvL=uqKy6C>0WtY8JBpXG`s0VOLKm{~ZhNUy=F$Gyi1h z?X!-Jnay6suaBH1sxsPy&+^p=HUBZGKAmcz{~K%l8>{;-%KmEiyXOBs_Rlo`6YMw5 zzsi1N|J=iWQl|b-T>p!9|4ik-z)CsW%h90XexQuW1dnXthp&q?hZ;oCIh0z0>~{~F z%~;41tBT0OO~A^=DWX{M2nzm2R(+-tWT{%fRGKvy_NltzY*DRm(;pM3?ruAV|0(*P z(fB_m8_NHXD9saUEda2(__qT2Z zC~=&{wxTYym8YI<*NkUnPRt@yKJ)vy6g9J1we3(d+`9***a42U4>`~@IM!U6k##Nm zD2IP5Sk;}x#m@Tl69D(O1-YN5xJ|jte1~Jq&m6Ms)$Yq3MYK^-CxBX=PuTSR&FoRX-ROOT-BuE5Ovn9t?H}FFK9!td)HG zEJU4fl2CCmUtNyBFj*pB6Fvoo+Xfl_S`LbAhAKau>wmlo^k_wgXsGUdKMg#jT+P~) zEZ+=pmE{R z6()g*$D`){$|##}3$o%cL~5X5si9+BiGHjLNm5sEQ>|lLR1N&z9()wn@y6cXJl-33 zyOX;5XkGPEQ|%$`BV9voi?@QzpB7SVy%)AHRBoJKSM1?MdtGBY6UGkx_;QuEu(oKq z)!6i?*KO0c()gmW%YB77C(IZ4;~trhA%Z;)0qlIqI!A9PF+wk3l|yNXx|o{A+t|n> zD6*ch*|TN`<(I;?rY|+HfChndCPAWL*~%d)K8=Li2K>?Tx~Xs&i?Seo+HV2aHL?~1 zEAH)o8uL1_%h*+(?3_9+aOw~ZrA&L*QOlj7t9|WU1v8GPzR4%g-GYBwy|Xkcj&I9g%- z>ahm4WxZAWS1d_?mRQT9%vNY|?+KL)yy_5W>QcIqOt(vmSuZwUo+L~t^}*%~3CRp% zHvfPdhhmLGiJ&eVAu^^5V#U;USaA7R^Vg%(V1Tp3=Q_iu?QTC;w^%=#%eU`gCa4JM zAJg#2$m%JRwUWj z@LPhZ_9NGpMy{tfX!})77tA|kPXId`kMCOtZ;l(@J#PNvuZ4Z_-xu~=itX!vTi5{y zm+u+Yd-d=ihnwvwFvQd$^wqEH*f~=Y5^L*piZpD0<$V$7zZIYhhdkaM^O34E0m`)# zm-%7RwM8|A;@DW9v=K-@dArDrw3jM9jPl`2I={QY_k>=^5Nfqpu}H4cZ3(AuX7-yj zjBkc7t00h1-ut^l*uMlpJxFEUdO`zLOG03N^ZOP8;=}46MJR4FbptzJT(>uBB4Agy z`ONgZ#kY`QG~x3f?yFeE_rREP8-3uW4fS7n8sg;^Ru~BrmcC<{2b8py9D5>{@FS3U zRP3lAy*`o;p%_*K&6hD<}zP*RQTK z#igsus3Qzek#Y5=siE}}_ImsAg6FVsvw)?B#GKSeT0U5(3q?pv--a|Kp*BEc*ic

9|sh(7;Vjy0u5SYIvcZo>yyU1FNz!vn052ZI@r+YO0 zzwFkbD@S=FCnT`eBEkkh_54Z0%*29zsfMV-h=-!mN!PL$Kz%kf8;_HD&sETgk9Xbg zlo#xsRLrs}5=B(4Vr`#mv4yyLVT?7E>-Si~)nBN5fxhfxmrzBdmvku3XxZG#IoA*G z!rk5HFSuHy(pV6*#(iMaz$L-G$1m4+`A1Fcvz6klK<$9xKEByIZ=0U{>06qv^yc-` zheywS0My4Qrq7o0UQg_gR}N@CR=CdmX#X{|dA{qc$s9o`%Vw~3FXpId>f&#@+}D_N z!xLxxe4z&;*Ee$>Tv=U?`&xE)TH}+=GX(+O|LNxc&7c`3_wnlFsQ+UQ&R&1X?bPN^ zDXEg-$&~$4WsOR9pY(#8y+p*!V_&E*jdLdN(QjW4-%)*vux#0?EH&=USLiU}G6q~p z{Ocw1H;B8-tpOAeR&uc#=IESbo%9R8qRs2a4i9Kk67XMr_@fP(Sj|% z8%bLQlC>==+!gv_~Oucq>L9)~(9_(17VUVt<{aJ8nX7#{&=P9aWM9ga%*1_M5u#SCn({&zKfiuy>;fs=)@x%{{(9l- z+IUl{(XFmeKC%fXBeekLi|u5}+F z-#qjOn*T(f!{o+{I2}^|12}URgW4hd+l#gF4WGoRNnmD8llJAAHLms zHhZ5-tkOMelXK_tzrn=a=Vjl$BnOng-ltrs{7RR1@|M&ovDg3y#xk9Pl~pN(AtU&PJLu-AJm!RD7h9t3l!x5ypCOvt)%Fj>AWoxX7`o(42I> zMKhL<9>25~0_v{2*?)7hr4QtP9Alc8`MvSf2`vYoaau5Rl=uAoN{@@Ac1-e{f4NJ2 z*e~y}N>QM@+;r-(DX3Hzmcdk3mf4t$yrsuSUb z;nI0uuz3Rb;7~H&qakT!zKVkzfE_Bx>FC?`W-gQPIB`UO{5)GWBz5}Sad)y=R2Yu? z%K$ym_*1!HB-u}@K;b&9p`Jq*Yph$mE!8jGb{_FTb(7q%BEL4lBs<~c)RzGkM+cb( zh|^jG`S$L5u(`BLokr|knqG6p*nP>jmN6C9w{@iGfC{>$u~TDs++vV6?5l~h6_@Tf z->ZyS?qBeURu8hj)MW}!DG7HaqIQE%Dd*GLAfc-%=_+v}a=Uo=JC|)WD`}J;GKlF$ zB=M@)^ux2pw}u^x12c@aN2S%qM*5r5^8BfHhgXq8sxV5tyUUD+jZM2tQ2c z!#zh=liWVZZWU~Rr&@T9Uc0z?8JIL<@ZL#**Qz{CmIJNGQL|2U&iYY9st=G`no@LE zZgjF8!Kth_Z0PuCWog}UP^Z&36vr(U_czdB8$AQni3Xo*y@%ykT=D*qTAAwl{YIQs za_daCEvUB7I0|3(cHAC_>1mZ7(4=@54>xtRDacD8X7j(DPriHt_*HPZSK~wUd6~i7 zK}Se&`L298N8MaQp#jOH)9=b3Ji{a&ELtW7P@#Kbt6#o8=fh&_C*-}xo1 zwFg%h-sfhuxd_~f&3ej?#&59A#EGOwT$k83#toU13e&vNB0CS`;8tA!xTo-S_U!}u z2XuaRCxAwG_#QAPFrTr>)3l}O}y!1 zz&U!(xoOh53;0^a)l9!yUYNd|JV>~zPq)a&cl{YDzF-gTZiOLjgKDP6!=1P!@Q}51 zgquaEOK}%ANlsQa*WP@r{W#6(TU zkf}O1)?A(7uK1$wBKsgUpX9xnw z1fQ0hVxNEyrOZC4*3A}EJ$^8V+$b;m*Q1Z#F%5hemswm!|M^}(<^uHjV>Qtjl8h%XzHSywI9Q?f(O+vLL1-h zt?yyV#WXaMm+)nW7wtNln&j&iY}DAFb0n+aIww4QL^4gS)+q6|>UpB&Gu+et;mI-K z1-D$ydUi|TJvgxHo z1=XVq?4PV+!rJQw!*B(#u9$3EB!NLrT@L$AlCL*1LJI7XF!>A;eqwTD>nweP&&LYs zrayGE?V+SOzY~D_=o-t7nyK6D!zuhg!xX2|U*)w<GUd$|E;byZJH&; zqi*N55t<-mkt{{khbt>9Z9ApfaZ@F`?D={(e!=W^I^L9PExyUK)4nw0;oS=|J9-vj`ZfpTgwxF_zTQ&At4RMu_;C>Tf$6ks@Ng?kgE3xf$FBOS~ZoBAQ= z14>g~iUc{Y-a+u26oJHf(E4FHrW2f(e|owYLROFv3F482MKb=ewB+dQiMI%* z`mHJB&sM1_b}2}mBv+Pqx`!vRA0`vf@{|X4a3M zG(OUuDz6{UXyaCj<3nc0}Go}H9^&C335EFmTL z8Nl>(w0W4-s`|&dd?9_RZkSw@L3l z0*VEc@10^?mKX#QN1Xt2T-4tDv9&h+gFYohKD;EuFA4I(LqZX3nDii3?v_BgmD5<4 zoBu~qk)tY4I>tP+i6t=mYJgO6FO~?Ns}&5(J|TM$VfbiA(jwy zr-tbar7!H`^YMBfRRe^TC9?F7YB|Pbi5VsJMn7wg3E@p9wqioPcx#d&`uqMsu*;K$fe2Jr^sz57Qyg&|%Gr=11|}T~6C!`!RGge> zyeqzY1Rccig8YjoDUTg{`xsxn;|gRY24*6Y+vn3**}^aN^W5*LwdcN&`f)Mq%cm=%`Yi&!?~<4CtKLn|nK+Rzn0;Lz5J20a4oM&N;OWujSD@QcQ6joM z4hoS9=so7+DlxOwmJhd*-U*Rm&1r9RgwstEcv#WMq2a5J=LzRLe>9gS2SbYPXR-O# z6o2Fk-pRVW+ifCUok);h{>7uv<@CMt?SpBO;uEbf8}SE{Np}k>BnCEAg4R-whe@hR ziS@PO;zOk-Y%w7>73bnAM@e;kKm(tO`Sh?$SkH2Cg;E8QbZ(URY#IroB^jCF7MQsQ z%hWE`Z!X~F+R#w_kh%<=H=cWNgNS=bNY3`;!JOWfjE(!MVmtVJYtw2w+u8EseWPlA z$Ryhpv+z`tH#14S)rK_~C(N^%>4r5W(6fzpC#byHQVnrPh&)&ty;A4wR6JNbkc80d zm)h=iEbEkU@o12-qt*7S1=)y8uSicMxD;pfoh#=-3vc$`M%^~5|0ermSfx|}JQGS@ zUws(I8prbWHbNNx!b&o(W#?+nZRBw43u@Dk#R{ju;+?_N6ww7H38kFxZ6Hr|rdq7t zflJE0^M9UGT!?sJec8|GO}>o+XzE-lwn+L=(qEM*Hk2XC)20|DrbL^B4Z5mi#QOE< zg-({U4GxxIdbh+Q+FBOxsx<}eqc`$`NlOg9P5krjvw#CO+tydxPuP|v^KA|a8}hTx zyW|T>w=$S|{E;+RWg(h~vB@G2P7D^%z54sKcczZKEh8Lm4+!zqg*M8ZmXQytd z|8QCupF08Ysf>@D?x~JUuW7o2*PePF<~>n}D?_I`jT}$!X#H^GyR)E<``cGBS_}B( zqtbZh*7a{Ogu|R4qG7>$pkvmDIj8Xu$hXYo2+Ny*ORoU%t$jbOEawj?2B|+qhfGi+ z>n;|QI0>hp6tp`fz!?$%C;=qZdRr_Z{rhq}n(95HTfuTW63pWIpJZ)3L>5O!oL$P1 zRqoaO>c7ekT-%|qr0aS}rdE(>h!GXdOzO1tvnf%Kq81<$Fyh;nm?AGULf&k+g}bPd zs*NJV`JCtVoW?$u7f3@E^swQ5DDZ5$@07;W8)-$C{z7##R3olqQE-T5hmY^ZMT z0Vc)qqWi^+SB1$E`*~ArZVD3xKKES3d@#zpNmh09anK~YWGs+?)GC|iLw+}nD%04DJ3r`I+5Z%kTH(Xl}%Rh}LakFJ#*P~Sss(K*v3 z4a5xds0AmfB_wMrn>Q3cgx|N2s2-0qG7+A{5>J&L8R3ugRY;EsVDp_iE28rA+w-=7 z$@gh<6N)YDu32^cNp}6}$lw4blglKh3MA^M1l+k?xNctJ9K}!1#7&fKrp<&f)X`rf zdr|k5ipS2_YW%~y1x6~h0&&6I(nThRia%=?&*08lmdt~4bkcX}!*ulrTbdgVG;VtaUSBd|~+B1kC}kf$hlTFb}nERr*1R14twn>Wm+AKay$6 z7a76_ku99YY0_zk?d@gi872lY>K+ByT6PbTK~p%cKrIrxJ%B$9QmMLp#pQ!|ogY|D zg|TqFjBLeAsq1rtIDy_zDx#&c0|)Uxe?Aj871b*G!>5D|8O1LGDat`w-{MUG3L$!9 z9{-}`o3LRWQQ+Rwa5W~%X%P6KWu4w3McUi_y~jSEL^>VJW;(Ee81w-ZI}$Z zPin|@T6k0P=}p2({m_bzq#4y=$@k+0!&dp`5(YZOdRjK92tQGIF)E}yz!lg>KjSh+ zlSxxDC=2XvC~?s(7|y_`#^q%(3!!pjhpS8@;g%TI9FBhf<6Eb(RsdayreDj`j|;zq zc#f_wIyD|vi4F_jlRYY5IDT|><#PY{Qz7a8K8ywl_}b9@6EgtGpK6Qz84A47ZQWcj zyNfL+kdyTa8?Nlq;*S0mus-|q?Z!vu0Kdmu@7JXQ zfD5*KZ*QpGZGXKrAfo!+==5a$#n$H`e?{0M)T*dSnqrnWAAg!c*ZPfs1^{w{me z8+oX&DOM|3%ZHt-{{uXm<%#xKA~tLj^HgwcntUXuG#*g}qBu1;ZGW5BVEI?Wk)oe{ zvWSk(d%uf!E8wR6Ja5Wh658Th(-8z#seRW~y};aoFyR3LCbz4L;?&S}s4w_`G56ka zO>NuSFm_Q<0i_p#gieqyU0NUsMM9HM6oEkKAVs=L?~sJhdkqkf4g#X|DiC^8X(Cmr zDhhr<_jd1d?m54E@ArG}ec$}C2y2d6)|fMEjydKtp7|w}8`)ycJe0PBurO=z4I$3{jtiFGn_jps~ zz3j31!GiqE1>J7AsSY_Jc8WZYk~)G2fnu%dC0A#>Ybnm1F@ZhOb6K6SmlcD~PQ0^c zJv`HtxV2DQ5lv)pOi)4r3PiTJ*jc^c$`;CotUP2pmXRp%nKz#Ap^_jrbRM!a3C&sFRrjmw z*GSQ9A{Z^H6Q*QkySh?I{0kTv7D$&EKUy`8hfgM^##DWi0j`=vX*GR{U07JSFUC3U z@zeQ+xN^{9jX;Vf%#Fj*W_QwAS4+6nF-3(i&e+dG#9@(!%iWX0ta;85Em4hY5whH7 zHGSW$=+|GwU5tDb_W%e9Z?yxe;+2V1PenG;(EA4Ij6!qcMw}gis#wsLRhE)?=@u$ zKe$k1Y$KW@RIGzg*h=Z<9x!j#4Mt6!^y&%L#O=1L%s=VkCz2C z1)QAb!i<&xWFF@9aFHV1XatEZW?xU-{R#voQBqpb8$6gEgwHAtm6VQ5!fOnLr!`KUm65hKF6Th6T0DRA?u;Qwe;f@mw^_;E7hRobh=K z*ze$b85M#>EZxVgPGkFNidNkik_>MQ1l>}OX9%5!>U+9tvA+@H*T_Ri^ng{5_X^_Hb?CZw{2rtTki!i~Q0WR1L(MLI%@jHpPcB;Q653TEycHA7-Vu&b);F@` zl^MYi)i~rPSJ+P!&uiMvExnBA9(zx?5`W9t8|4Mez*CM996lEpzCDzu5Ww z0uw}_@}0U=JQrL&6;4EHvn+4`qpMkTCiuc!tShf&SKFSA7hkOu&B`m>7qptTG^}|s z^-!k%Ty|CNfC&3QZ7c)4v@rC0!%L?eptud>JvgD)^g8IX+F?L*$IB`h_)@7yb<^Le=D zA@cgcp$L5{XwB2nv;5G81TRtX6(1zKO4>#~%Rq8A4V1DkWSppX6HiulH;FY;77y-+ zY6i!*0g;spmWh3jhq{x}cVyZe9OoFJJ6gM15p>ZS-#JjGsH3p0yKZ+&Z<9%JlK@Sb z`Af|_ZqO2F5P=;}0Z-MVG3GL-WO~?@Vn2@-CFZ?Lu{Rd0a|b~Wo!E;^}+wN zIYQv&>pHakWu>X=x;R1m^uh7U)$=h$*WaCZKLgKuojvE404WlCmEZ1W-(?Fo$WsnF z5<%vsFRar|lqHMbmq9uVq$x;6yM~0SHJ-X1Yd4AQbjhy(j4xTLx0^3XQozbqV~elp)HY zU)qJdr{K<=VTb1qRRM|!=aNb;ngBTt&rP1Nv z&IDhqIY+pkH3ia{DHx-(KW`%wRHM9_gUHmoKHXJe6Z0q3&T65lsE!%Ln5`%f zteiIeRrNJsF3)WZIZ{iwWpzh|*0zO7CrTKJ)JU(K7LX?jYA+f;ogf2znYNy(b>-y{ z-@Qy%SK7Fao~n8i1yKOJkhC6<{T{&es3QfQd??&x#oT(^NMol$U91A0Fw9ZtP?e?M)-RzyToMTIE@Q9nRH9iIH{_#sw%5AA&AhFey~v`;{?gv zB_lr)cdH~E5%%*{$h5X932Badj@81a_FF-<`yZ2Ohwn~x{2dw~GO4^vLMC$1{TTd- zOe87X z#g1q*p!PsoFdweT8yn9+B=r<>o{7K0=^NP%QHCHb^0i&Yr%JDTYaP6N3@il<@yZx` zhUu+RUMx7Z{*I@+hn?}@^)uat2p-hfY_8fku$l+#JS)aiD%H<~=NvKWeU{UU)m=^W zxLtw};-7rK6PwAFtBi&!sM$>lkRuFoG;=gjav(#x4`t*V-84DwK?+|VW_CEu7IgVq z&k1Q5<1O37(Zjv%aP`_R5jR>!?WE`5H53}?DQfUCRv9qPw@7jp7{VGPtmt3`t)wD>E^k_JF{0adz% zQ)ZVNk98FeYwr$$&JKOa&9j_m&E=QQdDR4XD*(Y3rB=XAEMW@e=t0~z0AtbAhW>t^ zs?Rk?=Z38Lkl9Ujoxw%i)s{=!Fk2UK6o-h&cXQT9A7KSLyg2i!E`pM`%>s}3Ros`z zQpu{+Se5#-IlVo+v^eIWkDj`dZV=%_7o4r`gUBjCfq zo}qA(9{*V3%5nOP=yCeYz=7~$*XjJVIDPNrlNZ&U++>++;$RxL$Fw7E6+2ZZ0C(Oisp~Qe8fH4YJirs12i$*@yL_x8&BW zBpMP7GcbN;g75vcnQit#Gr^GvPS3Fg^OF`!1 zwfHP$aFBpX4szOmiWOvT6VEKXAbTB&Qrf`h@knEpD@2oUo*&Kyxi;ai=rv#vE} zv7!%hiIZ?99%cU_nhhdWi0P2Muo@v#MZer7fbH1>Aus5z(nRX7l$M`W1S*Gv!EI=; zRpaNsDZ!t{v8>te_ET zUiER5@Ox>R%*krr3?6$)X2?J|5{=?YR_m7@h#j~|wSFserdE!&KNIdb5^k|_MS^FA ziz6^^eo{~$N3`w>;~qXZXMWX=VTt34eJQO_Sa@}+PoB5~!_^(F=kXP|{Ye`*`RPfn z1i!f#N-eisAuguX#iLS8hT)^^m#i3#`~=b`tCL*@oYRp)caxxbuo%=_#mHAs9-ppF z&PN^_+W@WwCTCdR8#m%3qOw(w=YHrpQ`l zG+B-=HnW=!NrQ3lN72;9s?2%GE{ajgD0#Ydb2rWFi1y0-;pAJNEDO_On7Trz!LI}0R(0V398UO}z7^ZxwE4t+ zkbRm+Ri}-AA5c&VQno`=R{{W7D2HZR$8AwL>(>GWVBsb8-9`T+r2sPxhSt_Kv4h-SMm_Xup&`{rc}x2O+m`bhYv4h8Sm2RN(YzAp%h);DmP9j70EKP zLo2K_(SkI-mGo1icxc(QP4%k9Lagqvy)BT+gu_a+zpa?Rnld*ABS>?bjjmQ7KeMkv zv*?#V6AyLQmQc6VZV=PJ*Q#UlfamS=BU4r{ipb7pWpvi+7oVMmg^TgH&?82C^0UL6 zmcJ=kc$_XIZtULS8nY3einMTTe?BkT#@`QxX#fG3@pR~D`euQa)%~)$CGq9z2UB~b zcZ3s%;zVImmN0oow;D7)?{vU?A*sVKosDM`q8C+yU0XI_VF|^JV#fNu%cw7oYCe5< zEtTkn!bQdLd}*-F4>I=a4c|L>N7Y3`Osf|Vd#i0xI&XZUOiLzo|=4%;i@pnz_>DFsnYphapW92N;qjj8UA)vV9yuCd<# zR`T5#XIZB`7R?)rwu{`QcNAFBy3u}bh+3~*`hu8cw(DZ+lh#CMX>5+~sK7-n>(M0o z3lDCN$@JvO;UA8^H@=YTuP?-(q@ zWk5NhA{qjxyD=c1#WiROJVwND86Dw7p7w+6EZ~w|rSan)4v^|;oo@g%IT+65C<0h& zRzPiE92~Jt^i(Y~(d~U<@kcQNq+-^Mb(RRc_E4FT$^jtFT~LUh5FL(4HEgi+!TvNvk?MO3JA_7AS((tLw3dhDB6*~!WMu~IT&uhYEgw$O z5-X`mzASQJ|X#UFXL9;dz)GB&VU}X`Tlj|Z(=kt8voB_yixIKnEHWooPfb@ zJf}3w6T(9+B5coD2`Xz7^FE6SP1xnZJTS$Wagv-+YF@f|3{jBrBxS!J zVnI2>|F30C{+vhJ_kaC+-~V!aVyZtWFy-g7OyBHjNw`0}dUdJa#wNcp1f37hMY{q9 zEeO7i9hO3+7@_=avQuc%%22XXuGZ!Is^G9tQp;v7?sE$37{8hG0AZ$C>Ug#Ktl`&c zHG=l=`tI~?GR51+L;1vs7n|btOtw+ZYSD}NW^Il-0`x%W6xbQvHQL}@Cgevki{{vl zW8JzQxl%5rCPyXe@TqS%Dx?J*^@ct@uKeEoXP!g!kwW{X&NtxqpnF?2|CtO4V-SSC z5*WqUC?d;Y9ut^$b$~q!Hc(w*crFUNh} z>exv9qBhjfDnb-zC1;R$g!BTp`$x`}M>jpjLvfH@BE$4dih6Fvaz*}0?@#_LCdE*X=8i%h0V z@)h^G1i$>ZTJvNFq7CyL%!~7H`(4@KRt^XkU&=$vz{N)C|2bb#WM)kzdHE0ZKl}4p z`4%C>bY3s)5s!OB+yB4X|4o*=lptBuv>?OLw?aas{jblDM?f<7w${@Z+@!yN0%p@*k?5~L~SY<#K1 z?L9)OZA^dKhkmpPkCl1a&BGyzc1<@-czrTED6cKgTASwrJ6spj7a8M8Pk*IDeZr}m zH;JuqMc)GBhEkx%?euELzM61d)#Dt_A1z>{?JA6Vb04v%$3#6~w9Y2YD>UfpTjeD$Vv; zwS4h3X7(|hcjhad{TJ;e3HLG|yWwk0nqbMZkQye8EH=dsmild&8O8Obq(~>)?g!bY zMe!3+1;*E))slKPcKg8FcGvKzWy|LxXXxKm*wU6t?(kal3noBgzA z?#waXq4^$RAmasXJc`7M-j<2A=GZ7^XV?@_3(;wL-dE>HBJLyHTobIyMqE?$)bwJB)P;4f(hx~@J}`&&V#0K;NUbyvI`NJ7|$9g zh39{PUi#8F66Au8@Zc^QEEWPWh%(Ha5{Z(;Hcql2Foktp`bO1&9mtoiv$uc#4vsIU zoByOgk33A+7M}S&Nxh1`Nq0E5^@B`0W;fDmy*UoI{JtwgH|}nMxD%O^z~w#=!B6ps zoj!uu)9&Zi+{N7EJ69X;uJt1Qmwu2rIz8+y?u;0T&TPynk^Y4XU&t^_xJ^d&@^{7m z@r!G*{K%-FXxtu*%G64Jm3KL$zL!2K6JuAQ6 z{TpLVtPGn|*NmQ~xw92kil>Om?SwhiIPk>V0B%++DeQyIGu0n`YLPFI-E3JoC@Q(Z zF?-UyGnn$w?4$eTCqfIY=Z~I!Za5~AToE+`zu7Y{9u(iu;xa!;Ajn%NhfcE7g#tmYKP8Nfpz9Mx1a%8e>HupL~Z17eRTWl2~W?}ujmDy_WrGnSA(1~kIY!dSU%q`_@7K^WxQ9r^?W8ry;P&Y$VHT< zfBa7GW^zBk!WgdWiu3+vrbR4Hq)@tO6Mtbe1d1QE02zeXc{ z*Qb0u0NSNb?UPgN3g9^!wzAlGp!`0FvJmGVp40c4EL{(;CpmllPENb-RIW6>Xt#^J zD-8p6W4)%1?hEtsa9F1x>~GSRvs|C!4B>Iz&Xc6UFrSVpcDdHLozg+NnQigzA*dB4W2%@j=R<|M1;^&mqFD-NQki{E<|KRo4iz& zyU}|ONLE$uxMpBpGF;h4wt=_b&?PZ{7B|B9;tA+&A&-4IHwHskB3x7qYZH*%$D(H^ z*-AQFQ6Uzv@*bW75uw}W;>3lj-rW6Yyw{!uq+hkJp3H9oyr6EDM&7AULQR11 zFl{v4@Wmz)=^dQW&u&p{6RB*8E})Ow7oI4w*ZEY#KO+})KVEx(GPEg2#}zP=!$q?K z93!HDfSM%BjQf#3#&lXI-&KFWHH=D9l&i3FE~21gb7)cLY8OiP^Ys^MA~M4MT~KX*4w}Vm+xmRzvlI5`pRC9UK!a7m>y&^;HvR@fI=q{b9v^1gFgKrb` zD5(T^>uF#!C;8SM#U4ayauuPnyU^7*KRE^bJNr20`ofQ+Ys%{hD9r9~5|{q8od6GN zPX=3q8p%sX%NlM%E;0*G6<1%*oOPLtGf0pPh^i4!j_&0a#Y-hccB96;zVv7(r7KH) z@i%LK#zVxHb2bRoYkdmP!9z882XTE7T3^*pm!_%ak4a`0Yp*&F@k;X;KQ$Xcr@)E5 zfHZ6UB@3GYe+%EKIJ)eGS$7IHY&oUS8{zOF6c~m|=qHCbuurkfuV2uRzx-hMV|~ID z?-h`G3&P;);8QzvRUSQ6L{x*q=}VGfC%-57ae9|RaWLpS1C>Uc0J$T(wsS4j`O^{S zEDY8z&k`r{+lyc0tyNUIiSf)RrEm5SC?sSLXq8^7sdz|VK;R4Seie})bI&sY?jX7= zY}S4Y3v_q}(_%i~8)>*`Qpt#!xa0aQSg+f1+_cuRf5c^ilQEtdTy1NrA%$Y+Ao)H@QJhoZuc;2i(`i=+`#h+}A2yvfl5?8sM{*-ha1x9ApB=I*?QON#U82FXkR30e+2fc z;@L@IL7q7OtE32S9pQ>hdisQ;VMol-sf7m;Qg(byrYrJQ_xf%%5&f;36C;MdrhVAC zuaa1_Sd|)Vg`H;)HBk&;vX@~wIA^{g!Z}C z2JH^jl>j%!=3!Yzvri5=9tT!Ac031qq_r~SF#E`?eG~Id<$GJeyIJtZ-2w1~4`bo% zQPPL$3x4tci>;}jS1F9>7*0~RL*q%ixYasZqRX{*F#s3R87zQ))&jYgn0IF~U3lT8 zD<6oXGNUvn4H_i*v^vUw{>Z6{?cz4^d;k25cIJg+(`8Y7-~XW_V*h-;KXqcxbI?=# zpU99VOA=|~FJ)f-rGCJlUDEoe;&gV4ntv&?^_TX;&i&aVe|F=wL*BBjpIv+Nm-Zhh z{MjSFyD>-}`hTH0vlGJdY;ixxPO!=$SA{4uW_+{-EaglFm_nK-O$4N06<$;}=wdlrG38GL+>*d;-kI6sC z_GknTidSmL$O0e#`sw__~hRBa&J8S&1he*&}Gew9)k# zVEKLng$}jX+Zj`9e~nYN4EdaU&2^1~N6BYZ??*Q?QLkAkA6ji_T_DSNmtiEuwI*-# z=*~II<};ds(WNlRTijl5ZGt)0uN$F{MqGnJJ_0Gu3+kL$e-h@Po1P>`}Pn8dIX44|}UBU>%-NKF^}tJOe(Mei7R!?`J$0J{?<(+2(Q z?^C4Ey^gO+bhLP?>ID_+%qG0K`$xfveon4y93eDl+IIa2qzr`R-@!`V%tF+218~8i$$R7y&}B{?5s)y zcZW^!LtB?)5czCSjb;|InE2o8ANctzga{Ywx+Kch;6Iilz!WT9>AI4X5 z(L3i?$bMzZvoK}9?$vRFh_7y&zWSi2WZc!B`H7a|HM<-EuMFpPOe&jVd3RDGkD}AlJM! z`VsOZm)n5C_N7mOJ5q$uK~U1=y1<*m>v%`{xEry7jX8#epHb{@6G>d#`g36Sq7jCO zv2Rzfs>Ot7nDB%v{B=c6>7pM&;se;O$9Fn1!qXfWia_U!YiA(@I7#1C&mm>Rg#1L9sTnzQ~G zt|rkPN!8z?;_(l(-AP=Fd+nFu)B{vlbXQNsM_Z3=$422yq)P`->a8MIMLOendn zaG*NYcnk9(?Q`r3tI?-- z{WFD>`5#k$=4reY{||Y}j!XL^PyBI(|1!@%)HC{DXo3EBQa%5^wEvLuyVZYxo}X>| zqt!o^@h{r;?{)Bp0!ev{H09!U8~xjP#r>Jut@r57yX6Ff7+U<#T{3qIw=+x3W}R}C zpvPpuUo#J8N69Dr$=CkoLLdFpe}(hUm#>RU{e@h%I4k#OBtFLtsghrF784Id@i_hy z4*q#ck8wWo6y0H`a+GnDy}a1@5_ObTN1z44)?O-dL=~6uP0T@Lov!B!|15FMQtwIR zdg1G9$wiX|!WO6=ekFkF=+tyxfg8=ewtJvX;0&EV+;0+|=9i--aX-ZyS0(k-6l?jV zM)C|8?g$``#vDJ(W{2;56-b*gu)yR#XbZ!o-Q$UE4B&0hh|fha6@cUNy{TUqr#;usbBeJftUx2V6?|rN2zf0_oi# z8Qb{qMe>~JiwnIeqSSUJgYgVTpgbP!xUJhZArXmAY0FMR2UH~LyR#i{sf<0h=+Y=! z+(mJNJC_UM?CVG~tKmM;Gn7`O@ooij*1p@UHdd7uP#0Jd5ygekrrW)6HFyv=4PiUa z4!w%*3XO24q;w}U;~h^cSHugc7V8>C+`(c&GH&R87%W0O)-sRNXbi)a#^T?h95+u$ z$yjg_mgk}{*r8qe<*`=0`rBiDll3;IXroHKqXaM74xwv+(#>P^;mlqObukJA6DS9X zT40pG)If#dQnO%piMfhmTY?7>m15{xU}Qspxy^DCjb#DzOwj9i$s|7>KMCtI>Zgx1 zxjM`#7}`#^L&r!V@3*Hcz+W4l8F_2Shg;;VdY*qsOMMp;+}jLgvyX*3rga+@zSf)$ z5bx&4WY9w)dosbN{7=X~7oswu$^%#5*O^Xm8@o!3r`Jq=&@tXFCN&SYfvXWq6$c4_Xh)UL{mVEVdh{VA9mT{8ehgIPw!35LP@9ZcNC&O6TMbppSK zKqzb!8HwrYJ7VsVV3&G--ZbB7@!fDHPpTrj4(I3LqnPqK=-4pc4m%a!pz3P)IRA0t zn6rMCdh@;V`8#^rKIrtbb}t}CuWz;7gX+Q{Y*ledYlg)xY6GB`h+7%@_&0EV@1n?L zJw0-G(nNg~B$={?SPGZb*g|K^Jl1&(qi;Ky?ba?+tymc{4sjGO06+_G&1)*;jv?rj zfbFA|`WWr6F7!S(U+{WtNa-PrC^=LuUth*?;74Ub;E4F@DF7yyn47G@H4S=6@#zmE z(xfG&2|kIj{n`onS8_G3 zT4R1ApE57^dO|UFk{+{lPc#GPrwcO+AdmCMa%~Zul#G1r!a(3i;c&;IOeFl1+=c_* z4pJ>{-%cQ2t*|GD8c7(DD4^hK5n@@MY)H$Qq>n<)Lv}F4+v1*5;*C#U*cNBudsEMw zELOXhwv0)>TLQp4)X(|T0Wr*}Xf=#%qJ+@iST7*64S~5vrizJb+6|cI^i_R%sr(Fl z2AP=>9(SQ8S-DB$(}yx82Uuyr0)uw?q_yy1j<7*w`(+R{IB)Lom6b|*QP^vEZl_ol z_hgSEqYh*qZ-~r~%Luj8^AWD%6I0Qj4PAt$grrp7c(HC-GJqW|hpY8{w_I7g2&8sq^MbeMp_TOcF`E*>9Hb;1KvuB(h zB+0abY?ioyG7!;ya*o~SJ(&Zdr*$OTVY-3B;jYUF*arp4>*@uZGjO<9`be}%hi8Nd z&Fh+t)j-lF&joXJeBrnd5iwfW4|Uqjvw`11#HT~p*>)aI>i54Jwl}^?%{?)2e~U-s zdE{3JNF1fFn!t5~!aXcR5dP8aO*O;T>mn#hL$1eMg84EaMG0Vm6)=GSIsJq(?+ehx z!3#~`%dJSGKABsK4?-u2W^pOQ!q>=MuiCm4vGK*2Qm2Y3XX^R;N=NOS1Isf0ns8DyN8<)#$C+Wh zre|-8G`qAji?59caLzN?SMe#Pm|8nU(Q~6qg{;~c8~`(Vj13eUQfph`B38Uc)K_~n z_!(Ea1qkHpwtzZ6<91>}hJvZUh9lD>@I)?se*Iwi^Tt zU6}Su@AU}YzA66@Uz9hMu%HNlS8=P#?PnT<%$B_epQ9yATZg&H)2$E6y!qUFsa*)>>J+-*H1T9UR2r*Qp_DL z61z~vjYqoaX%u;rcHV|9r5O7nxA`NgRFXc*zGJ5xe_e{Ajm_z%uGIA3x}VoV0+rPGMG@qwNWfNl} zpxx)fIvYfc=XUyBu*XHaf})T5nkQp-^K27awWMY| zMvL!t-VCCm0TbS11h^(|T%D0u4|&!~$yh)Ym^P)O-&S<)>*rP zH$_$Q@XPcmn#!E1&KJ4D6(9O)Qq|4qStUd~;}_pcv7ags{ctB-Q07{^x+Ak%q!q8}Qx`N^W=x>8w)Sxm&pw7~$azJI|x?P%Jc< zd$=sjje!(NGgnaA5Xl8JzoPkgbiH>o$yW@d!p>B~(IU#$xs+Fx z2Q#&*bVOFaw@^viucgSexlP9K_j%^e@3)RgWdHn!j7<9G0m(I~xMf?${D*7O^;cA4 zgzux00a8@r=Fg}^uV1l;$e&O#riW|7AIMJL{aFgR{O>5e_WO{pSua#>e}DgWM2?T8 z)x4A958my}NR`V|S8V zR-o5$z5&`4>wcOlkvmZ`JKWA=>ir2cXn3quZx-0FK&F9=bOt(=n)?6-9O10HRDD_8 ziD!WEY)fi%ntCmo9$F!yX>zCa-4VM=P)I!IYPbY*HL(Dan`F-vKDTNYSJ*4UMWMFk ztn$*0ArC$&L6N9EEdjUoP*T&(&EVCGbL7|{i9 z`Rg!;+2YOwouso+diy;3B|PaJU(6rDVdxGX;7rm?tZ=sw2i#K0#jU9!!er`^*|O7} zlq;@X4))oKW+5)38w{xP`lO(}+|ZXQpQLIvYV3=HQ%pUtt&DU>L*f~rVFeIInn53x zMWjW3WbLB{^R1WPz$)O$`XxOEar|d?D8RL_VuXZ3HtObBr%LbUn&zMh8#_Fke*hT4 zJ9Emt0f*$OqDBiq>Lm0*mjx9zi${?PqNovcUu*B7R25m?b^{Zk1`Y};~$RA!$4BgOR_6# zoUV4)3e;zuc?3Xc*u4*o-xK4I70&N(&IId?if&U?vBvXpuzk@f+)Y-bndLlfecoyB zR$@sOT6=#)BTw>i_&dd9lWA}Ms$}f8#)a8*jPLhtx$zs{+9ZrLB`0%j*oiZWGT&^_ zUjfZ(HL6d}vX3p^CETo)&(+b$J6k$HM^s6aJb7^-TW1L2qU34qYo7|>u1tLV`5K7(14ao3Ma`*beh(peKsYCcV3&t7_E8q%kP!k-x?ps`hn^&6Js@m@5Z;NN_+}D}4{LJV6cm-ue^MRwm&=*snd-f0 z^1aGzzsmiBATq5_T7)#}y^LxZ?o$;ne^x=~UDViK?l59Y6c&ADsmXS(4mY>O;7B6D z2#5NsMZY81RI4_u;Vg-^Dqy5e*cCo)>{E!;91YkqCF&N5_^0LM?kVGX!4#L__zTPq ziW7cWNy<sU4puOrg7LrZ9X(j>*%E~?01HvZAI}s!u?f+5njF6 zvmGb-!oLLQ=9!2>&+uA=I`!|@uWP6}gKimt2VkHNBXbPb8ve`&8lQsh#LGJgCeKIh z#TVjfbnY6)agrx(T!D+{HGq~M9Zr;KTj~L{IR?z8a&^&*ib`P;cQCzPap}J3f$F65 z4|uVwpG7HH@-#ebna4chatEt!&8WuK%-UN(6^W=WqJ@MPuhV2Cqy(DJv-6x7GAbFM zlus00k(jZoLWv1_uw{gv-7>B(q!|`yHZ2r6afLpQP*DNzt9J!dYA;Q8JuQ29hx_VO zhGC(|1(C0IOm2p@ebvs?qkQg%Z5)?o-0hb_4KCT{Cm(`&su{#+bqzTeD5{h1aoAU3 z2)1q;2&_ti#n*EqDe3)nTQqsFdY+gO9_{;;ROp~~bS*tx*l79HU|Nv|a(Op6bsOT? zb!JQ^*s{|`*Af#c7x=t4Y!ObQoLktKYre;dz1!Cs*>`JYu+0^iAD5J_ra+?9OGkQA zUKt?`uCsr+XlMTR{lDUg9sR=-%k((mO=5;=t;6de>nea=x{g32QKv?yS1fZzo7(O( zK(%{