From b6ea8f4d3a5ad90058f25ba80ce283720c758968 Mon Sep 17 00:00:00 2001 From: James Ring Date: Wed, 14 Aug 2019 14:20:04 -0700 Subject: [PATCH 1/6] add draft kernel RFC --- rfcs/20190814-kernel-and-op-registration.md | 658 ++++++++++++++++++ .../device_api_overview.png | Bin 0 -> 43897 bytes 2 files changed, 658 insertions(+) create mode 100644 rfcs/20190814-kernel-and-op-registration.md create mode 100644 rfcs/20190814-kernel-and-op-registration/device_api_overview.png diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md new file mode 100644 index 000000000..32b0471d5 --- /dev/null +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -0,0 +1,658 @@ +# Kernel and Op Implementation and Registration API + +| Status | Proposed | +:-------------- |:---------------------------------------------------- | +| **Author(s)** | James Ring (sjr@google.com), Anna Revinskaya (annarev@google.com) | +| **Sponsor** | Günhan Gülsoy (gunan@google.com) | +| **Updated** | 2019-08-14 | + +## Objective + +Tensorflow (TF) currently provides a C++ API for implementing kernels and ops. +The Voltron project aims to create a modular/plugin-based TF implementation with +API and ABI surfaces. Plugins will be able to create and register custom kernel +and op implementations. + +In order to provide a stable ABI, the Voltron team has chosen to provide C APIs +to plugin authors. This document introduces the C API for op and kernel +registration. For authors who wish to continue using C++ to interface with +TensorFlow, an ABI-stable C++ header-only API is provided. + +## Motivation + +Presently, there is no ABI-stable API for extending TensorFlow with new kernels +and ops. There is no guarantee that a plugin written with one compiler will work +with a version of TensorFlow built with another, even on the same operating +system and architecture. This makes it difficult to distribute plugins without +also distributing the source code and requiring end-users to build the plugin +alongside TensorFlow. + +An ABI-stable API for extending TensorFlow will simplify the distribution of +plugins and allow plugin authors to distribute binary artifacts without +necessarily publishing plugin source code. + +## User Benefit + +Plugin authors will be able to publish plugins that users can use more easily. +In turn, the TensorFlow community will benefit from an increase in the number of +variety of available plugins. + +## Design Overview + +In general, the kernel and op registration C APIs aim to permit the +implementation of any kernel or op that is currently possible with the C++ API. +Where possible, existing C++ function implementations are reused from within a C +wrapper. The purpose of the wrapper is simply to provide ABI stability. + +Since plugins will be dynamically loaded (e.g. via `dlopen` on POSIX), the API +avoids relying on static initialization. + +The intention is that existing kernels should be able to be ported to the new +APIs with a minimum of reimplementation effort. This precludes a from-scratch +re-imagining of TensorFlow APIs. + +The following diagram describes the components built with the proposed C and C++ +APIs. + + +----------------+ <--+ + | | | + | Plugin | | + | | | + +----------------+ | + | | | + | C++ header API | | Plugin + | | | my_plugin.so + +--> +----------------+ | + | | | | + | | C API headers | | + | | | | + | +----------------+ <--+ + | | | + | | C API impl | + Core | | | + Tensorflow | +----------------+ + libtf.so | | | + | | Core C++ APIs | + | | | + +--> +----------------+ + +In this example, there are two object files: `my_plugin.so` and +`libtensorflow.so`. `my_plugin.so` is implemented in terms of the C++ +header-only API, which is in turn implemented in terms of the C API headers. The +C API implementation is provided by TensorFlow at runtime when it loads the +plugin's shared object. + +This design addresses changes that are required to the existing C API that are +required to support op and kernel plugins. It also introduces the C++ +header-only API, which currently does not exist. + +## Ops + +This section introduces changes to the C API that are required to support ops. +An alpha version of this API is already checked in at `tensorflow/c/ops.h`. + +### Registration + +In the C++ API, ops are registered at static initialization time using the +`REGISTER_OP` macro. For example: + +```c++ +REGISTER_OP("Bitcast") + .Input("input: T") + .Output("output: type") + .Attr("T: {bfloat16, ...}") + .Attr("type: {bfloat16, ...}") + .SetShapeFn([](InferenceContext* ctx) { ... }) + .Doc("A bitcast operator"); +``` + +The equivalent C API will be a series of functions that operate on +`TF_OpDefinitionBuilder *`, a pointer to an opaque struct (i.e. a struct whose +content is not made known to the user). The functions include, but are not +limited to: + +* `TF_OpDefinitionBuilder* TF_NewOpDefinitionBuilder(const char* op_name)`: + constructs and returns a new op registration builder for an op with the given + name + +* `void TF_OpDefinitionBuilderAddAttr(TF_OpDefinitionBuilder* builder, const + char* attr)`: adds the given attribute to the builder (equivalent to `Attr` + above) + +* `void TF_OpDefinitionBuilderAddInput(TF_OpDefinitionBuilder* builder, const + char* input)`: adds the given input to the builder (equivalent to `Input` + above) + +Additional functions are provided for setting other properties of the operation +(e.g. `TF_OpDefinitionBuilderSetIsCommutative`). + +Registration is then actually performed using the `TF_RegisterOpDefinition` +function. This function populates a `TF_Status` indicating whether registration +was successful and frees the resources associated with the op definition +builder. + +The C equivalent of the bitcast op registration example above is shown below: + +```c++ + +#include "tensorflow/c/ops.h" + +void InferBitcastShape(TF_ShapeInferenceContext* ctx, // see the section below on + TF_Status* status); // shape inference + +void InitPlugin() { + TF_OpDefinitionBuilder* b = TF_NewOpDefinitionBuilder("Bitcast"); + TF_OpDefinitionBuilderAddInput(b, "input: T"); + TF_OpDefinitionBuilderAddOutput(b, "output: type"); + TF_OpDefinitionBuilderAddAttr(b, "T: {bfloat16, ...}"); + TF_OpDefinitionBuilderAddAttr(b, "type: {bfloat16, ...}"); + TF_OpDefinitionBuilderSetShapeInferenceFunction(b, &InferBitcastShape); + + TF_Status* status = TF_NewStatus(); + TF_RegisterOpDefinition(b, status); + if (TF_GetCode(status) != TF_OK) { /* handle errors */ } +} + +``` + +### Shape Inference + +A significant feature of certain ops is their ability to infer their output +shapes. TensorFlow will invoke the registered shape inference function (if one +is provided) when it needs to know the op's output shape. The registration +function declaration is shown below: + + +```c++ +void TF_OpDefinitionBuilderSetShapeInferenceFunction( + TF_OpDefinitionBuilder* builder, + void (*shape_inference_func)(TF_ShapeInferenceContext* ctx, TF_Status* status)); +``` + +A series of functions prefixed with `TF_ShapeInferenceContext` is provided for +the following purposes: + +* Examining operator input shapes (`TF_ShapeInferenceContextGetInput`) + +* Creating and deleting shape and dimension handles (`TF_{New,Delete}ShapeHandle`, `TF_{New,Delete}DimensionHandle`) + +* Manipulating shape and dimension handles (`TF_ShapeInferenceContextWithRank`, `TF_ShapeInferenceContextDim`) + +In general, C analogues to the C++ methods in `tensorflow::shape_inference` +(see `tensorflow/core/framework/shape_inference.h`) will be provided. + +## Kernels + +This section introduces changes to the C API that are required to support +kernels. An alpha version of this API is already checked in at +`tensorflow/c/kernels.h`. + +### Registration + +Kernel registration with the C++ API is accomplished with the +`REGISTER_KERNEL_BUILDER` macro. This macro expands to code that relies on +static initialization to register the provided kernel with the global kernel +registry. See below for an example of registering a kernel with the C++ API: + +```c++ + +#include "tensorflow/core/framework/op_kernel.h" + +class BitcastOp : public OpKernel { + explicit BitcastOp(OpKernelConstruction* context) : OpKernel(context) { … } + void Compute(OpKernelContext* context) override { … } +}; + +REGISTER_KERNEL_BUILDER(Name("Bitcast").Device(DEVICE_CPU), BitcastOp) +``` + +The equivalent C API provides a series of functions that operate on +`TF_KernelBuilder`, an opaque struct obtained with the `TF_NewKernelBuilder` call. +The kernel builder is registered with TensorFlow using the +`TF_RegisterKernelBuilder` function. See below for an example of registering +the bitcast kernel using the C API: + +```c++ +#include "tensorflow/c/kernels.h" + +typedef struct bitcast_kernel { … } bitcast_kernel; + +// Bitcast_Create, Bitcast_Compute and Bitcast_Delete actually implement the +// kernel. See the section below for discussion on kernel implementation. +static void* Bitcast_Create(TF_OpKernelConstruction* context) { + bitcast_kernel* k = (bitcast_kernel*) calloc(1, sizeof(bitcast_kernel)); + /* initialize the fields of k as needed */ + return (void*) k; +} + +static void* Bitcast_Compute(void* k, TF_OpKernelContext* context) { + bitcast_kernel* kernel = (bitcast_kernel*) k; // this is the pointer returned by + // Bitcast_Create + /* compute the result */ + TF_SetOutput(context, ...); +} + +static void Bitcast_Delete(void *k) { free(k); } + +void InitPlugin() { + TF_KernelBuilder* builder = TF_NewKernelBuilder(/*op_name*/"Bitcast", DEVICE_CPU, + &Bitcast_Create, &Bitcast_Compute, &Bitcast_Delete); + TF_Status* status = TF_NewStatus(); + TF_RegisterKernelBuilder(/*kernel_name*/"Bitcast", builder, status); + if (TF_GetCode(status) != TF_OK) { /* handle errors */ } + TF_DeleteStatus(status); +} +``` + +The registration function prototypes are provided below. Kernel authors must +provide a compute function. Creation and deletion functions are optional, but +if a creation function is provided that causes memory allocation, a deletion +function that frees the memory should also be provided, otherwise a leak will +occur. + +```c++ +TF_KernelBuilder* TF_NewKernelBuilder( + const char* op_name, const char* device_name, + void* (*create_func)(TF_OpKernelConstruction*), + void (*compute_func)(void*, TF_OpKernelContext*), + void (*delete_func)(void*)); + +void TF_RegisterKernelBuilder(const char* name, TF_KernelBuilder* builder, + TF_Status* status); +``` + +### Implementation + +The main classes for C++ kernel implementations are `OpKernelCreation` +(provided by TensorFlow to the kernel constructor) and `OpKernelContext` +(provided to the kernel's `Compute` method). The analogues in the C API are +`TF_OpKernelCreation` and `TF_OpKernelContext`. The aim of the C API is to +provide functions for working with these structs that match, as closely as +possible, the C++ API. + +### Inputs and Outputs + +Kernels must be able to retrieve their inputs and provide outputs. In the C++ +API, the tensorflow::OpKernelContext::GetInput and SetOutput family of +functions provide this functionality. The equivalent C calls will be +`TF_GetInput` and `TF_SetInput`. These functions operate on `TF_Tensor`, which +is already part of the existing TensorFlow C API. + +String tensors will be supported in an ABI-stable way. This will require +changes to their binary representation described in the [tstring design +document](https://github.com/tensorflow/community/blob/master/rfcs/20190411-string-unification.md). + +## C++ Header-Only API + +As described above, the main motivation for providing a C API is ABI stability. +However, some programmers may find the C API less convenient than the +non-ABI-stable C++ API. To address this concern, we plan to provide a +header-only C++ API that is implemented in terms of the ABI-stable C API. This +API will contain classes such as `Tensor`, `OpKernelContext`, and +`OpKernelConstruction`, whose names will be familiar to existing C++ API users. +Ideally, this API will be as close as possible to the existing non-ABI-stable +Tensorflow C++ API, so that kernels and ops currently implemented in C++ may be +ported to the ABI-stable C++ with as little implementation churn as possible. + +## Device C API for Kernels + +So far, this document has not dealt with the challenges of providing an +ABI-stable API for kernels that run on GPUs. This section describes an API that +addresses these challenges. + +There are a few approaches to running kernels on GPUs: + +* Assign computation to Eigen device (for e.g. see `OneHot`, `Transpose`, + training ops). (>200 occurrences in TensorFlow) + +* Call `device.parallelFor` (for e.g. see `BatchSelect`). (4 occurrences) + +* Call `ThreadPool::ParallelFor` (for e.g. see `MatrixDiag`). This is a + TensorFlow wrapper that eventually wraps calls to Eigen. For example, + `ThreadPool::ParallelFor` calls `device.parallelFor` in Eigen. (29 + occurrences) + +* Call `Shard` (e.g. `CTCGreedyDecoder`). This approach is deprecated in favor + of `ThreadPool::TransformRangeConcurrently` but no kernels use the latter yet. + (42 occurrences) + +* Call `GpuLaunchKernel` or `CudaLaunchKernel` directly, i.e. without calling Eigen. +(58 occurrences) + +* `Matmul` op calls directly to `StreamExecutor`. + +* Possibly others + +In all approaches above, TensorFlow core is responsible for maintaining +respective device queues, streams or pools. Kernels then use these queues to +schedule computation. Therefore, our primary goal is to implement a C API that +enables this scheduling. To give an example, one approach we can take is have +Kernel pass a callback across C API. Tensorflow core would then call this +callback. See diagram below: + +![device API](20190814-kernel-and-op-registration/device_api_overview.png) + +Furthermore, note that most of the approaches listed above eventually call to +Eigen to parallelize and forward computation to device. For example, the first +approach above uses Eigen APIs directly. Consequently, we need to understand how +Eigen works with devices and in some cases make changes to Eigen codebase as +well. + +Finally, we should aim to create a smaller API. Some of the approaches listed in +the Background section seem to be very similar. For example, calling +`parallelFor` in Eigen is quite similar to calling into +`ThreadPool::ParallelFor`. Therefore, we will only provide C API equivalents for +the following: + +* `ThreadPool` and its methods. + +* `CudaLaunchKernel` function. + +* Computation assignment to device in Eigen. + +This proposal focuses on these 3 components for now. Due to the complexity and +variety of TensorFlow kernels, it is very likely that we will need to consider +more approaches going forward. For example, how `MatMul` op would call +`StreamExecutor` directly has not been investigated. + +### ThreadPool API + +Here, we can just wrap relevant methods in the `ThreadPool` class. + +```c++ +TF_CAPI_EXPORT extern void TF_ThreadPool_Schedule( + TF_OpKernelContext* context, + void (*fn)()); + +TF_CAPI_EXPORT extern void TF_ThreadPool_ScheduleWithHint( + TF_OpKernelContext* context, + void (*fn)(), + int start, + int limit); + +TF_CAPI_EXPORT extern void TF_ThreadPool_ParallelFor( + TF_OpKernelContext* context, + int64_t total, + int64_t cost_per_unit, + void (*fn)(int64_t, int64_t)); + +TF_CAPI_EXPORT extern void TF_ThreadPool_ParallelForWithWorkerId( + TF_OpKernelContext* context, + int64_t total, + int64_t cost_per_unit, + void (*fn)(int64_t, int64_t, int)); +``` + +Note that we just pass a `TF_OpKernelContext` instead of a `ThreadPool` +instance. Implementation of these interfaces on the TensorFlow core side can +then retrieve the actual ThreadPool object using: + +```c++ +OpKernelContext* ctx = reinterpret_cast(context); +auto thread_pool = + cxt->device()->tensorflow_cpu_worker_threads()->workers; +``` + +For details on how we plan to switch between `std::function` and `void +(*fn)()`, see Appendix 1 below. + +### Device Assignment API + +This approach lets us construct device objects (e.g. `Eigen::GpuDevice`) on the +plugin side. This is preferred over the approach described in "Alternatives +Considered" because it is more flexible when it comes to Eigen assignments and +operations. Basically, we get an Eigen device object and can apply any +operations we currently apply to an Eigen device. + +We could wrap `StreamInterface`, `ThreadPoolInterface` and `Allocator`. These +wrappers will consist of a C API and a C++ wrapper on top of the C API. A +sample C API for `StreamInterface` is given below: + +```c++ +TF_CAPI_EXPORT extern TF_EigenStream* TF_GetEigenStreamHandle( + TF_OpKernelContext*); +TF_CAPI_EXPORT extern gpuStream_t* TF_EigenStream_GetCudaStream( + TF_EigenStream*); +TF_CAPI_EXPORT extern gpuDeviceProp_t* TF_EigenStream_GetDeviceProperties( + TF_EigenStream*); +TF_CAPI_EXPORT extern void* TF_EigenStream_Allocate( + TF_EigenStream*, size_t num_bytes); +TF_CAPI_EXPORT extern void TF_EigenStream_Deallocate( + TF_EigenStream*, void* buffer); +TF_CAPI_EXPORT extern void* TF_EigenStream_Scratchpad( + TF_EigenStream*); +TF_CAPI_EXPORT extern int* TF_EigenStream_Semaphore( + TF_EigenStream*); +// This would just delete the C API handle for TF_EigenStream. +TF_CAPI_EXPORT extern TF_EigenStream* TF_DeleteEigenStreamHandle( + TF_EigenStream*); +``` + +The following C++ API will wrap the C API to provide a `StreamInterface` implementation +on the kernel plugin side: + +```c++ +class EigenGpuStream : public Eigen::StreamInterface { + public: + EigenGpuStream(TF_EigenStream* eigen_stream) : + eigen_stream_(eigen_stream) {} + + const gpuStream_t& stream() const override { + return TF_EigenStream_GetCudaStream(eigen_stream_); + } + + const gpuDeviceProp_t& deviceProperties() const override { + return TF_EigenStream_GetDeviceProperties(eigen_stream_); + } + + void* allocate(size_t num_bytes) const override { + return TF_EigenStream_Allocate(eigen_stream_, num_bytes); + } + + void deallocate(void* buffer) const override { + return TF_EigenStream_Deallocate(eigen_stream_, buffer); + } + + virtual void* scratchpad() const override { + return TF_EigenStream_Scratchpad(eigen_stream_); + } + + virtual unsigned int* semaphore() const override { + return TF_EigenStream_Semaphore(eigen_stream_); + } + + private: + TF_EigenStream* eigen_stream; +}; +``` + +Now, a kernel can create an instance of `Eigen::GpuDevice` using this stream: + +```c++ +TF_EigenStream* eigen_stream = TF_GetEigenStream(); +Eigen::GpuDevice* device = Eigen::GpuDevice(EigenGpuStream(eigen_stream)); +... +tensor->device(device) = < computation > +... +TF_DeleteEigenStreamHandle(eigen_stream); +``` + +Note: `gpuStream_t` and `gpuDeviceProp_t` might be aliased to ROCm's objects +instead of Cuda structs. See Appendix 2 for details how we are going to handle +ROCm support. + +Wrapping `Allocator` using similar approach should be trivial. However, +`ThreadPoolInterface` takes `std::function` and this approach would +require passing `std::function` across C API, which is non-trivial. For details +how we are going to handle it see Appendix 1. + +### Alternative for GPU Device Assignment API + +We can take approach similar to the CPU device assignment API. On the CPU side, +corresponding Eigen object - `ThreadPoolInterface` - has a `Schedule` method. This +method schedules a kernel function in a thread pool. + +Similarly, we can add a `Launch`/`Schedule` function to `StreamInterface`. The +default implementation would have same behavior as `LAUNCH_GPU_KERNEL` in +Eigen. However, we can customize it on the TensorFlow side and implement launch +logic in core TensorFlow instead of the kernel. This way, `cudaStream_t` and +`hipStream_t` only need to be referenced in core. + + + +Advantages of this approach: + +* Don't need to pass `hipStream_t` and `cudaStream_t` across the API boundary. + +* Supports customization of the `launchKernel` call which might be useful if we + want to handle it differently later. + +Disadvantages of this approach: + +* More invasive change to Eigen. + +### CudaLaunchKernel API + +CudaLaunchKernel appears to be a fairly thin wrapper around `cudaLaunchKernel` +in the Cuda Runtime library and a part of their C API. + +For reference, this is the signature of `cudaLaunchKernel`: + +```c++ +extern __host__ cudaError_t CUDARTAPI cudaLaunchKernel( + const void *func, + dim3 gridDim, + dim3 blockDim, + void **args, + size_t sharedMem, + cudaStream_t stream); +``` + +where `dim3` and `cudaStream_t` are structs. +This is trivial to either wrap with the TensorFlow C API or just call into from +plugins directly. + +However, ROCm's side of things is harder than Cuda. `gpuLaunchKernel` might +call ROCm's `hipLaunchKernelGGL` here instead. Its signature uses templates. +Fortunately, AMD is planning to add an equivalent function that provides a C +API. (see Appendix 2 for details) + +## Appendix 1 + +Certain parts of our design involve kernel plugins calling a function in +TensorFlow core of the form: + +```c++ +void foo(std::function arg) { ... } +``` + +We can't pass std::function across the C API boundary. Instead, we plan to wrap it with a struct and break this call up into 3 steps: + +* Wrap `std::function` with a struct. The struct will contain pointers + to callbacks for manipulating std::function pointer. (This will happen + on the kernel plugin side). + +* Pass the struct across C API boundary. + +* Wrap the struct with a callable object which can be used as + `std::function`. (This will happen on TensorFlow core side). + +Step 1: The wrapper struct will be defined as follows: + +```c++ +// Wraps std::function so that it can be called across C API. +struct FuncWrap { + void* func_ptr; // pointer to std::function + + // Function that takes std::function pointer as an argument + // and calls that function. + void (*call_func_ptr) (void*); + + // Function that takes std::function pointer as an argument + // and deletes it. + void (*delete_func_ptr) (void*); +}; +``` + +Note that we would need to move `std::function` to the heap because `FuncWrap` +might be placed in a queue and called later. Specifically, `FuncWrap` +construction will happen on the kernel plugin side and will have the following +implementation: + +```c++ +// Wraps std::function with FuncWrap struct. +FuncWrap get_func_wrap(std::function f ) { + // Move function to heap + auto* f_heap = new std::function(f); + + return { + // Argument to pass to callbacks to call/delete it. + f_heap, + // Callback that calls f_heap. + [](void* wrapped_f) { + std::function* f_std = static_cast*>( + wrapped_f); + (*f_std)(); + }, + // Callback that deletes f_heap. + [](void* wrapped_f) { + std::function* f_std = static_cast*>( + wrapped_f); + delete f_std; + } + }; +} +``` + +Step 2: `FuncWrap` struct constructed in this manner can now be passed across +the C API to core. + +Step 3: Since we place `std::function` on the heap, we need to manage its +deletion. Therefore, we wrap it with a class in TensorFlow core so that it can +be deleted once all references are gone: + +```c++ +class CallFuncWrap { + public: + explicit CallFuncWrap(FuncWrap wrap) : + wrap_(new FuncWrap(wrap), [](FuncWrap* ptr) { + ptr->delete_func_ptr(ptr->func_ptr); + delete ptr; + }) {}; + + void operator() () { + wrap_->call_func_ptr(wrap_->func_ptr); + } + + private: + // CallFuncWrap might be copied when it is passed to functions taking + // std::function as an argument. + // We use shared_ptr to make sure we only have one copy of FuncWrap + // even if CallFuncWrap is copied. We want a single copy of FuncWrap + // because the pointer stored in FuncWrap should only be deleted once. + std::shared_ptr wrap_; +}; +``` + +Now, the `CallFuncWrap` instance can be passed in as a `std::function` argument: + +```c++ +CallFuncWrap call_func_wrap(func_wrap); +foo(call_func_wrap); // foo here takes std::function argument +``` + +## Appendix 2: Working with ROCm across C API + +We need to access `hipStream_t` on both sides of the C API. Since its +implementation is actually in C++, we will treat it as opaque pointer that we +get from a HIP function (on the TensorFlow core side) and pass to another HIP +function (on the kernel side). + +Ideally, we should only rely on extern C parts of `hip_runtime_api.h`. There is +now an equivalent in the C API right now for `hipLaunchKernelGGL`. However, AMD +have said they are looking into adding an equivalent function to the C API in +the near future and are willing to add more functions if needed. + +Note that we have to update `LAUNCH_GPU_KERNEL` in Eigen to call the HIP C API +once it is available. + diff --git a/rfcs/20190814-kernel-and-op-registration/device_api_overview.png b/rfcs/20190814-kernel-and-op-registration/device_api_overview.png new file mode 100644 index 0000000000000000000000000000000000000000..8803ead7976648535a8e29907cf9e7109413e8ec GIT binary patch literal 43897 zcmeFYWn7fu(>8n&N{N(6NJS zz{0Zch5!40?&tme-tX_1n{VtD=geH^IdjZ0GZC8V3dDr8ga7~#D=9wL1_0b(0Kj>L ze;WXR7b+Al0U#JqdM^FKdv0$5&-;bd9OlS}S>hcoJ}D{XEix)a# z1`AZTRD*FJGL2Ie^xrPPCnDeGp}qe#Rgvr~0U`D0+K*VWBy!UJ*KRqPPc!M>uoI{z zr1zb>mU2ox`Wb7iU`ZJ|#vpjOHV*>x*<>k+00NZR?VF4?G zCU>c=@#9E-I*$-;sX;GX8*zg3@19^Hc>ocfAxMM}^k_Z2A_)HXD{!n82iPPm7M84| z^x$>VrUuA4$Es>%>5AoL54=v-k)7pJ$fpSf%R6Gj?W` zR(PPZ4R8AFR92RACK=!-ay_9w%x&UqKObA!!up?u+xz%{gkbrQ_!DI06BDWDOX@P| z-bVX}6%MTxBDlV~&J9zh1oiXGRvdfD`DSK1&emBtTW<%hDG&jiuys9aP`|4>VFCxitWF|N(AJ?wFJI9Bs#cBaKw?-D8{`;{Iy;Ey6`OH60rFZXb^7M zXhge63rtP%k$8>QH?I3fUtW{*{m1yrBbmfliKJGQt-N!YBoPqyu*J=w~^-Z!K=oDA#a)+S%M$LQ?ii2jL#Ty zIZHosH7OTZ$Wa3X+dsem+-SkmXo4TNh*fg+r!?1IP2Gmdgva&K$&4Fu!rU=_%>==w zHH{48Z(a~#Gu^z4ldkS$rsAF)rQ57{1ER56DVHrXo+(zODnFh?5 z1WTol10Zy%2{H(C$O)8W=EcrsSC@I{mU|Dr)CNZtA4WqRdP-`P7Z z1JGdO9$H*tOixVMV8X zSG~kcdiso8TiEDv!`9aqa;6*LrOwS$_qfRZk91_f_UqHb=f3dC(W=Bnf$__3{b7^s zpXworaHE}f@u8A5zHt_ad4jiO0L{#L8ZyXqX?Xs-!mpyoXk=Z=a-VoaEg4;xaO1)M z2pUKam~a!g*v;{ttyM2hcX+Np;mu&^Uc2`SgH?$B@WXb~uH>8+RZr$`{BIY4HX;Lj zOlt$P6Ej00;~6$sjCud9bdob3^cVGxz)XtRSMJ#qF-sL+%(K5N_inD^2XjUqqXETm zm%bKgO=Q~R72d=r1HXsrF3$gX2W+OuFIOx8_RYPdfYv=Z;NQkdAY$kI8~|Lp>o8 ziO!hrmnodsXS(_9f!bREXn2_Qi`C16mx<%u%?>As6Ae4Z^3xd(My(35+VSC@`9S2c z!5sVmxu-Y^O-=!d3tPMA>>FCu|JYZwn%2~om{(c)IZnY8nQA&x#Kd>(uyf^+hE`RN z19D6p^@hJwI94&ISn$4eRb_U&8k{;yEy2ul<^aV7fs*KBf}piOo2NNTF0O;Bt!wk2 zkL+(WfR>O1ORDadg}i^SI~!bMxOIXtb;`T-j2yeKjdk?VRvT+^yrX3h`Nj`BuZ7d8 zy4Ak+4M@zRhV)Jjg2Xjn5uuRVrT1{5b*H1xWn!?9cKLWuYcMZ;GdOiJn0@gf&)faS zOg*X|0{zu9tY~Vt(_QxTxjz1<2+uxWF_Aa*frI)wuvEk@1i|!3Ny;gvBufCWD{*Ey z6Zkc>I$^XlLx8qsS+*^3Za0#8qbE_Ujk5UnaSNUTD2HwvxXSKO!@Rp(RmRZewC%gN z`9=HX%Ug3x_7&H9L_uq21aDtD80a@JxjSLYr1eG};;M zfOgeGd&e5p2NY(7z(oq5V@ycYnIWb?}Y445zFPY<*El9+*@m2^!DLl$FNGO+mF~ z@^E8HKbb;r#8I{fSXUxzlfQMDH1plp(0kxn)#nw|hoSibcV0dusT~+EQw+&x44T#6 z0u+5o``V&?Lh*x(weJIcfhfvclOzW3J2N=>3#fWM3IOiu#%o(Rz>MNgT%bE4M!TN) z4=>>dyCQ;SLJ_efXFVTZsNZ)4=o`Mp7V-vqZz$dZ&}G`4?P*Gcjx&X z&ul)7%|ye?-WoQT{HLO|o;-a|#TSd8;F)eDPOg%ArE@7B8zl8pIg+ijkk(-CStGb2 zOwA>s+$?+d;h*ZeSoY=vV&-)q*of=F55`77Hg-zv%Al# zNnT}bA|Rn*(Lz}F8JK{qpWf&x zFyRdg-c4Y$v2V)>xV!1KUp?5VHRSgY7Y$%!b9AjF|KaGpdgN zktsZGuIO0kN; z-rIuLpg;xt;WqQy4mJ6mSp>Zyo`<=l~iwEG2wam{Kg$_81Gd7tKuF{%HSt zeZU6<*QmHmjhU@=sOtaq05w2!g4L(Vvt7XQGx~8CuXz-L)V0B6M)JT8ZrSJM1@Bg+ zDXDv!l9_5@z!C1L2L@0FjJc~>$nF0Mj=^OFV1Rx2{~15AnIr@M$#oOl{u5S#e_=eC zRtsCDXzb#Fh%I=z%u!9QhcSR>4Hrm_eNYnc7_Hoorwx{25~V*yEP zU-elX$+20Z?98z_(Ju$wdtJ3y=(bpDcK6Y;&p?Ze8$7whx&1fd2lZgAMPv#5vES5& z-oIdRgB2oK0OhgwsdN&?eUtA$Eu#3#u~eR-eJy0xKcm(we3u_e4+9da3A<+fo99}n zqY>DyjzoESnRAANose3)(>n&-zPIXI=E^2I&3bB^iMb5}V_2kbzX5QCUp!f%WrXC5u*}6;~NCYDs5t-nBqee2oLdHFqp#oS$57;bu-R0L4LcGp{PAOH=tAHI8QtKO;=hE!<+82dGDEI4 z2LXjH-N_Zcaj)61v?+|7NYa&fhsS)c)lPE~Gi!}r`wT|sSq@unx-IfJ>LC$tJu;4j z1?y`4=H!4|H(TOnM^G97md)J&=LP7g^j6~2_gDtH9OX5#-_#GZdS!c#JAGUa%$~8& zNo*>{>?(Ur@Mqo9Q|-^IyYoFN6&xS+rqSDcayIUwNW7BtwiP-J#c zPClQh;VguusrDblezW9-FPO!Bz+s|ss{7e@N{k3lj{W=$O--ggM?B-4=ot$$9$9({ zn&cDZNM=d=h{yv9HL;B4rbeEFC(r)pEk2HbC8FHcyfN@I_%4DsymOVt%D}p4L~K{Tejs=_~zbwWovD6YflGV<1vezLtQ*}aT^wxuT6g=>imDb(*v07 z@bpQWRPoh>D3KV8(*S+s`!%Y~<@0P5oIP`{ej|z!@`S`C{r!$EL{|6z1hgxg#9*=hT-N`%hW{7he(|F24UAC0?`_2alt7-YScSCEv-=pce zskBd9z`0xyQ}}s`>JJ88-|eCM^efl%WGVB?-eWR_HIyCy7;Fy{$OGqA&`V>6DWniF zzFOKWtDm3J`r`dj8hCL?D}tm*j!V-lK4fbD46*;xFWxxgc(Z=MrMQc0?3|Z>>YSg$ z2U;mBhipgIJ4rJG3C>ic|B|TD-hP&9W@hQC$&!H8UPGZ7eHtABl%|WzXqYl33*WO= zA*2lKC68{7Q$-!h-<`4IjmCr0@y$P#|4|sYKNM7^ zceX(%BLCcbYr8a=``X*&wErnKcE5Xl&cHzhI~y`Evl zw?^aP0lsIF3NZ}(u9shjvnE5+SEo|YyZ;!C;PHFs&lOVr>P;6i-so1n6yfINJREmy zSaiP*+x%R{U}9LZ{Rll3PDwEpG3Z{P?)><;f38)LC{mup&y20R^~*KNW%ViYMASm6 zYN_l+jH=&t>5I(%Z{>(!r^2TjkuIs4Ul8bmgE`TH1ckaU3q7;w&{? zOR}+!3oue@mu+}BFM95(Eu}``U#`}rW@{5ogqLNOCri?tujxT%M!!ib?MxdviBpg< zc#3mzy-q{(+jI;7Z6L09$&|)04)Ml~# zC?~P!^aM8Aqw9$F4BtsOkBu5j{xYmvPTI&PE{-rU+Q%dTru0C?LYuY+e7bAwsjRQ9 zVl445!cyU#-t?Thyl{}w(0u)FWEH7A{XUy$>I*l0+Zhs-?ksUOxT1ByW}<)CM_R|? zQG>EFH;kW2v3+EyImx#V*Wr^u_2!A$&&)9z@IExcF0R3xNS*9$roXz8l#`5~*uwcz zBe|baqS(aS6oL?9yS3s7vXYTZ6^NN zjRuy(DJ(@KWa^q-LeKZ+`F4u(X|8F3zHi~Tkp1;~;jK-uvYbRf@!HEMT-V`Aom%}m zJvslpshbI%bpnr}N-eF<&^rt*}`m?(w6Ijk$q8nRxYHc=0cadN9 zYHvfBA_J0Dsn_iJtc>(FY(zI4tkHEleuLSFhk z#~53Jo9f5s<;~&S9xru2a#dCQO>T95^ln%nsXWBfc~xv!PA=tAKwSKp7UmZ5JtF|& z*v*yv=Mlfg`F9J9;4i#>dtdAV9_1-^)|__N7(P7c7%H$^w9)Yg+c$pSRbiw1roO9| z>AWf=d2|B@F{&z13*zAE9p;#1y7rPbJ{DHHov#Z;D917X$2g`{wzCkcZd~wuy7|E~ zK(;~TD5O1HveNEiD|X6=CtJi+Qv51u(B5(&OX#DI2v$dys7Gb<_rQgWpkg``X`(=N z%R-7}?{Ah{^E$Bt5{wM@U6n~fw#Y&sk_6XdY79=d7P1aM9M2ghwRIe!l1%N`#wdDa zu{56sNGYj1*;^0<3+|?#iImWx1z8S#4BO|m+vknjuSzojMeHvK|KXWG>=S(=nJyG} zu{eFmo!m+85{l|qXm{nCTp#tPt0nd%mx2ysi0+S`gwK-1hXSe>F0Ym$4^CXMJ1Avn zh03LLLvmszRVO#9nZG@1nrrbbwexk}n*~vZ-(GA0g=rvXKd?1%#6o9Qx%zi?awGav z=-ni}sga&D`+?tLFAH{ORN%r_??EFIs5|*}k^}H**_1`6{xm$r>+eus+=BFw2llZ{ z_JI2rQY=j=p-s)iSgVfakepzMm1K&`YL^0|2v?a08A_NheV4OQ=aOhfwBj#`VpLC@)Vyr8KPEB(1|kHOTx~aZK|)U(MF!v!45)OM29MLwiHTZ&e+& zE?(xpIki)e4SA2;eka~~h?T^vTl##usnrzj=wY|yqLlK4Bkoapf1hDV6YKTrYkmFR z*^eJbJ+EXQ0>=m{*a3;di<}tx5=oj1ercTbDS@YsO)S7p+^8whlK>ZT8v@gs$14&3 zmCp6J(C!^m$4p#41NC1SP4xrZ;UW({TM-9GyA;J2WcldwTX4=6r3iLkWN%7w09E*b z?BTe#vbGP8i9}#;)R{&b!xMBzcV?O%iLnYRhPQRY(o!j)k4RmD68}`!)Xd;XtWI{V z?&800$*m6*lNE5%$N0anY|Ii24+A;uZ%781^bhF8bE}-8{-^gG5S#h_{6In@yQNh{ zoo~lO=^ zk$P}Jq7Pgb*(SDM)ybc76=nMsrq4Av-u)RxHNEy0{UWT_b(4jKMgC;qBhX9Xv9`zj z>sG(D-6@JPb+bG&io4zv%5W#x!8b0k1T13xRMS%dSYIQ)t$!)0a5yWCLeo+xdGSDj zuL19B*N1T|_l#nm4y81BC0u1XG>aeAP7+Y{0GbFH?Ea*SCy#~CCoeH)D3SD@@Q%VQ zyT0C_vOh-7BKOJZSbvIQu!uTxe&vWsRqE$m6-+9#sAh&Qpi6O~wh0 z?_^{`OjoQ_0e68@?$#hq#&U$f03BAG>T6Rv?hr~X5SQK}jC%h*WJ9+?exjK^$ zq)yE7v)#h6vvmkIjt0!t&zL+M=a2`o#$9Mv;1=<=>d8zj2_H?6$tjmL_X|#`vF_&{ zxX`d-5HYi==K6c`=El%d&@+$zC!IJ{Ab;YXoZPw?0AD}B+bkge?Jdp2h^{U*+uhJH zXZ{IuxI=`qEViPtA6GRy5l;T~{HL-xkI^#>PmTX?h6l30Ko&|ro#Tw0T~pGCqVTKD ziiUr)w}8d|Jun@;F4R(VIRv4U7oCxi`JJ2Ha5zj5lte@9AM>=OgBjgW*$LXS@N*PdC+7x+}TSJJ_4=!cOj)i~&d1N=|BS_Rt4qxud8 zBz80h$t?}eFzyFec+yc`spV}*%})J%AyNzhQAe}**&0ioow44jd-F6OWNje!L+HW{ zUjO*<8;JxO<#RG5Bs*7Yd!h{{fD?NjTGwY;Udhh4p!W3@ni$$qJ+-u>^tmQ$B5CaP z6OY3+h8u+*DsP}lJOVx?B2Y&asO78%zCSr)20TY>rrkh`OsAMrbJ^>#<93n%$CtDZ;=YB3oKKImJR4!!OYTS`E>V%9cAqQ?Cnw)|M>-km~|EN=douT$9cKiJ&G_}3}k2sR`pi0Td< z-MkByCl5X5;%;2CMQKH{d? zIBO_k5M^XwXd1iL{@8m^H(u`u`wZG`->K@3R~`vg9Jnz8P^YfqUXY;3D&n*FR|*U` z1dlQo_W7{lzMrI;sx{2}MtL?P#+iX8uSAN-iu6m+w_?DrTw!>By&4Y;B=P z9sC#Z+w5T1gohm!s~go#_A%+N-4b%0h|D}sYSxULe|qL4fBN^nv(8K2ov zc~`PE4|V3`Rc-v-I|W^!H8X2DJfnG8S?$@5M#8qUIWOA^`;6sg43X$N4SLQRNdZ9> zdziMj6;?|cQJ|Rxo2P-^t9dZ09K!&ctwZKgt<~@ta~<8%myKU$qC@${uaRQ3HjZ4n zrNy^ITw$9Nl5U84ey67+Qx=o8d>mr+LZv8=sXDab`y33IRK0*ggQwXc5bcbm;%``S zp%B;0RSk*H9A6WloWiVs-rlYD>JrW3vvD~F;8*htd}eooZz3D`y;3KM@8AnhH=jxc zv2Y$8%t3hMzzi>Zkm%W5xAqYb*`FUf|MVKnDX(8oT*0jm*8NkCd|(R_ecg z!t6hlEqMxq*MYd@03Hixue$TD`|iQp?HDnq5uZzV!-(%+pQ6ZikR-eK8>w-NIGEKY zUej_qb9>H%9O(1d(CoywcPcZ}t8<*08!LO`-p>gaBzVKR9B`t2nqT%Sr% zWR*kyj0SFLG+Sic{ub5{Q%@psoa>pD*-)Fv#4i5OxViEPRvfO;#!1{Or4TNHAj<@s zRXV3qkHmUsH~b`ose&VHBUa*z1g|!^a-WAQ#$N;gSs*Hq+cFz_Cd_&GyXe5iE}ZT+oSeF1)@j8SGg#}&c-{yCPaG^`Gdz+#~0Zs zCPXUQYT#FXhWP=tf*b4zFl#1rC6lyXqjnkpj08Q+BgBf6x(&tnUw#swd5{;=XdR<= zfW^Wb`8Mm$P{KDw8Rp%%C$mK&&l4zHSiP(7UeXF}OU<>dr@4wa29|mnU{kw^ZChF5 zs!bmjVB$z>FYS_WRmoAkN36n-Uu@g0788Er1|PhY#I^}6jZ@s5U6&;E=BtXb1GHg1 zJD<4yc>#!TH9co|!SXtPV}Eu8k@(nCt^T(5K!%;Hw6%!r2<`EwW3p4wjlAaR$f&5O z6+UV2>Y`3C+gC$>qxvszG2=(O5gHbL&*moSjoTQa?b54yd7&NH+4f@ie!(z&{8wZF zQDNE0^;6Xf{m-R=sJ8{8s1tS%1ul#G}d!&di5y-;XCbHhce z@$>W8V1@VRZF&2VA+?u!T5P9 zCRRSL(!Ni=)FAdiP9EIzBR(2EXxz^xq=26tyd^>NjP_EPe@Xg)Q4JD+JiJ2e<)I0l zE5B_$8cej7kqtTflLyz_m-oKDFQnL7?CUQk6OFQN5wKvthOrr%q7} zB9gTDR4PZxh*#VptC73(YQ<8jry42xlXcWI3fSU$iS=}2xcWA9x-aPyd?DF~5pqX0 z#VR#ha6F^8kVNV!9<*B~7b-3y=ocf#{xY z>f7?bbFld&$$B@UhtZ5$w{ku`*+PLdM02NA$P6v=$YLM+6SeRDIEU3VJGQLP%Z-uy5a~oQttOr)H zLwGItBtndDLAx`h#!nPC@LcQ6Xw=_49pYY7!5S@u@6uOoN}_!!eLzIlO7>OneUY!? zRx6MAm_5tv`_6H3fe6e?lGFOWhc3A-75nWCea;Y?Co6bXlL`?>sWR_FV#u+vIBPi= z+A!0Rl>L}&QQJdGv*Qtli|%{ZdrWbU@=);S;j|LR`Qkx$%JrX>{Fcoj*CQG#H9abU-mkVVMaffcC1g$}eCthu)_STN8Nctwi@eeH=2 z<=YSm9o&3{3zg4(cq!PBXl_&)9!&WRNSNxdG-cWqp>F7FTNAWsV_wcP0T<@apuDX ze2y1S8-4G}fW3~yeJxBK98zM|{$`MPrYDm|kaIhd?~E#@Vj{&Ob$ua7tSymn#e|AY_ho&e{BHRt)i@>3vr^q04^iEy24z@Mtr zRiEvt{HP{(SJ{XU(3;oAU4-%0Yq7~8ZOG`pd><;*OThBQMoH@+)*362kU-hMN*cdD zZ?%<#_{{G+(fhk!UW~xq0TpTCYP>DY@-7Zn(Y4jVzBYM9S1^FUb2jI(7@vffWmh$* z4r@m|{F7VaN!+3eIkLBHNM-DS>@+KGf?IMzW0XVTq7+EC53%iw5+g@vTxPOcRYz~gHBxwMpk^2PVg9D95dOHcQPd>iU z35(*>@U@!nQ*F`U9`kYECSWy>3uTA&#QN0nY%8(2VH8JS*T$p!jHSpCy*p+b0sTAX zi?!92or)8Uy>x|RB#v4Y`%hH zt$&&Vivp|P-Qmb&kXUVzUdZj}n0ShZWSh?=G#+{&BuamuRy>syKUhaYu>Ihc!4daU zi?#{A69{79Jl{JBt%CP6%1amZ*C1w-gFg?w-uPMhhvXh5#?RZE6N95gh@G1?YcvH? zi2)GYFU9jB>Gi^r3Sxhx6rvzk%wTc?6#z4|SWJ=MrWC$F@vq6BJ|#PR9=mIyxOg!e zj<`%D>1RcD$&u}Rs%~uwlPPs(!q=QmOOmGa7^|mwQB=ctllgvnRlB!XJ(eWAGVc-) z?DU8Yjo>i(p%dKVnBE=FLPju&;;x)nq`?XI>1pX&+P!vzuH|#3P1NA(8HGFZnk*c- zs_AVlS%7eFi`0RK?}P_mOpu4ym6rI~00Azqso_%PRTYx%)Vn?ECg9EBLLrQrmKTTL zk*S6{i;may0KITq#?4YYv&{+T+C-<|6!gSR`dgBE-ZX7t;hesDt(Nc0EN_f6A=kmDsTVNTAT z>+#Z|W+Kim2A{3U&4r;$3zG}`zvpDQJZ;({MlHO!%F&lgOZ^I=e1Eex+p*$a-0$oF zvum&pvs2dPHtDy}Y?t_w%Htf=r{h!Mx{U0I!)NiB!;oWrVuzUx85_mJyO+D%p$*iU zqG2zhA|x?W<>pT9wO;V?F)|RKGH1Q&(%aedtO7$>ssS42BuPQ7Unt=$1?p38<)j5B z!o_N^8H&f6lDc4WY*)1c(n%oB*VO(Z3p#mq5VPdnT0ipr+vaDfz4f+bC|J9A7KeDk z)NlJygV`SvxZz*?AW;uhrWZ#gK<{PvJ0&XzRo%S2pVoVPKdBePuBm__GZ6(N2_?y= zl3Po$n(vO&`z2Egh|M3zEPntGG&kfx14d<-$x)sgX{Wu#Mai96&(BDz*(C11wBV&z zu6doHI=s6-fQpn$q{ABBif5!$(mc8dx?T2$WYJAd;%MtVSM`F_)LfNLl=x%*VP^X! zLk!_Nb?4)L(yaw#RSut-q(9`~Iff{>ZAic1@87?4IS-=7sYrP>CAIK0_8(0j-iVQ0 zHFEiVxBcm4!rT1G^PlLcRjuW_6CV#7^L>GbBC6Or;zRD9-jw-6^FIv6xXj%?*}gJ_6BAf?+=MX zrFz{i;Q|lx*&glLG#DvIO_$n?SWkJ|I}a>7xo@~ zN;03iYA@R+e|F#{SCtrxLQY{vI=7$8zzu4i)2nLe%`W*TIu}cunDkk2$5x9^mst;L zPRPf+KsR}foV0+=l;IPb&I&u0Y^}rK31V*{!jd}e2naGjO+Of++>i1d-XJFq{yJ9| zU}FE>MO`_2biFEe{(zdDolSZS5y2=n*K*;j0{69M)`S}h*>w&ZgFWPNh0?8}K1e=T z+9qqlw;theEoSDQJMR~>j_K@oyD#SIr`?hpKSciJZ*{TzAZY14G-GBXd)^`dBrrk5 zAFm~$k3KFhY@2~!c8!mH?E=SgIAj(VB=09<{TfQyxbK(B7Lg?B-?%8Kp3q~~wRzCh zN>W~1@}R0jzj?Ps25r($<*QKDd5^PtSf1{Pul)X8qBE5y8Bkd~+xejk?zSOJ{4CJ0 zR?+@M*3dBG1GmD<%E@7?z8`NTsqUa#UWX|^5PKPromp-o+4v?|v}Uc0JG*Bki{cs| zWHNttYd-0~JhPuGHUR5}e0^8kVMM|8N!|c!UQ8{`pQ`Vk2EHST!W3j=h3v76rKijB zE3ZeyFFq3DbC#tHXB5{N)wws-FDT$Nt18XpC{F??TiU;v1-$M`?$igg*^dyoh3u3NY`i=^kjb>P@ z_HQBPR5Z}0CEAl5EwX5EqCuHKB#HFnap_!jP{vNuq>`URzrQXR@|Cq4zl{n6z^<0q z%m{-5P1-cq%gw585sjbx*2O<+L%U61<>XBJ=9@^jiZNE)l8|k2cdDy`X@u6N7`o&pq(7@5Bf49ZT1KdBbYytMZ7)SE=slia$be zm>SjQ6pb3iohAWzO+P;TzEy~165-&HM2)y&GxA=VClF|*&P)+TK;_b z1DFQlO`dg1`pD$Yp0u-m%vBb@ijQ=rs`^jEO77Q|7ndm7B7fVdDhuuuN@0g5*ttKI zlnlRFOd|4R!OEDM->6>2i$@yHOcw_AiHovb^fB}LtxflpFWdlD4?DOJm;Di4k6V*J zqKE4V6V!K2;Uoq3>NiE~BJ8Srw~W&1uQlHe7Hg~XAlB$WP>_7{W7#s+TK-f-e>^O@ z(fg}cjB$#oko*C2y2e)6W{yiEq5i}rt*zsHRztPf_B%@nJQgSCQ4if-Io9enfPI26jrMq;jQQ~JjqLBLzrwR+rTGQb*qLtDe9(gmA$e_kw>GeNODe+v!ys*R&l!b z#({cIog?R`0=LmdOFrImHL0_+*5na`bZt~0`nu0EBQ<=x*X#9yd~I9D zf&JH(Mj{O#9tcejpP9j7W^aRr*k)q)W+&bn%nev=)1MljY|NqKE?Le7yu}H1M0RIN z1a!1(*Z+D2i0CGhkffSIRq=_>o4UOgj)Di#H*G0;G(dpfUyuNqu1@!zR0LQPd5ruP z(*{~zr{3vH$TRP5k$^-G_%{mxkd3{!UW*a%>uq6Yz0Fr9WU9Z4$2u=mzsB`+(NnJZ4wVryQTJJFAeco;cc`6{codXx3}j! zv8)E{@**>5TC1z)0@u$$+Ax|yMc&2oZt7avK3M37_4_`*r#{BwmA!ACNdl$qZ25{88w#1${TXd;GEI=r_JSWze4IjIPUV1+AZAg`e_&gL z8su1?j}Ox>D{3$hB&<=Bul+$VzZwoN?}I}oRg`_77>wxtVwyE&uS6m&nrz|XJ*in_ zosx(Mzy|COyX33YOef+D^{NO08J)STRyDM%Z@Z{(_98UNS|yx%*hpM~jg@tcAjlW7 zZ~DdU#>AxSzkW>zB~c$tQAmwEczSiblz?;l9@v(N%Q8M~0>+I9e!bB#0sC48KCp!R z<<_Gb*EEsn61}n%&x|o{`4Doq-mevrR_$tFX6MAp#MMjsWVQk1ncw&LGk(=mmp&*$ zctgNR05qoqz&execdHNFnA~Ui-$q9CR(m!t=r4x8zU8VP=23w6eB3l4JTLBHv~|@h zh>Q<1bv($a&bjOA&snbj*{~cdu6&Q}=O@H;JV?wptL8q2DUL9>jk1`>;|fi$7eXQhJP=ScXI@^0$rhqX1;AohK9xF~uihe-TKqTEKWsCiHZe@XF#uLF2#yahs z7{)v`<$WHoV~bSNuIpuF^x)A(M9$e@dY)))l$U>?I)h}(Z?MI5Zr!2o+jT^Q-rmXT z+0uxc@irx^gmXFE@;tfxN_Q0KrA}d!h#1UEbYOus`8|%lw7u(mY;-zk0I*^udHhJN zjj!KuSXw#~Z=Wg2Hx%|SR>cp9#>Cz7f3YfmI&-!?jVf)A(%r zTB#gzZ~Lvq!f=`mNbrI!;49I%knu%>5vKf^xbGlkf)(fU#g@j2{RFQsQ^IvNO{Et< zl65{lPI&`2!Q5Q_6xQ<3fRP@F0Dbs^e=ZCsv29!UpMgD z%{GOZQeNF0ZM}=b!>+})z032Tsj~0R#JpZ|RK5;O4*~IpK8xGd52E1c;NIdOukF9o z+H5~RHb$Fk$^UsM)qeQLf*~DuPFQNy-y|!i4CqDa-xCB}??0mf^%AHHnU0?-)0MeI z#%x{=koA@J!B0tGTeon}vwcsUip{EJE>yq_qAE}ybe$&qqD!szAZj5`QuMaqAyBsT zPDbFmHi)L-3V>l>30URQj9)RTF*{isnRa`U5+=lV$tzoHxpZ9hok{suVZ);T*xS1` zjMi95^HkfnENc@lIa=R?&=HO_weC6bd}0wlf(z9?6KXs;J6>lo-Jmx{fxGN&Vvq5d zTfWhmlii?oejT&KfAu}GZgFQ^=Q>{UN7j1XX>962aX2z$=I=4V(Y<>BYh#xpYL*s_mXDmz} zEX-}-v}h3LS^?74u@4E-2RGx*&i__`DE2DpUeC!wNr0Gu^&cimx_j<%H=AovR_y`r z!C1kRKmbjArocGAw<1QWCIJWFgifxD0CY>i&J& z%nf*J(3=3rIS~aVkJb zDfZKBl+M)h+PgW|h5e9V4DW^Z1JR92`RI!unk(3si0go117yKSxFjwVTB0lVnDuv= zA4aiH{7JVp+>T?Q#E%O|h?7r&G=wvP@y2r2J3Xq$;QZl0^?=|27?uY{1nmpe5*!{2 zI470UvcsdTR*Aq)y};;QM_%89dmyPxs#Mv)%j}G6+y!#oLl=8bdro)rD*D})9_=>E#jZMG2VRW0sg8J#(dZ`3^fr|Y9)+jkm+K=0y_mVyh?>!mAH3etIl=>pA$Tj;i<;dm1gkmPa!}j9Ko!~&z>R(6m1LKDv-|#}{ zWmI4NLBp*B(+<36R|AexF%QNjSj@Wl=!}X)J@%H*HUHEEGD}A_%%sC5X|AY~tNC@o zKDeAtv2<{>;i?foI9OYD7Y(+`)aadE#}fq8b93CKbyr)76#;5eLfLC@n{nxYUi_<(aBh1eltP7 zX!gwfmjE8LOBa%<>I*2vy(SJUXOnPX4Hfz3UNQl*@*~hLKfyOpVNa`HfBMtIT>01I z?Go0`>0>cZt*L0Oe7X@#XzXpfs)(Ir?g3+H>Tc{DA0nLqMDN}CZ8t#i7&kl8Qx`M<>b7QCFBo9RK(&FM@M?Azk!^YL_8 zE((kyToDP(uOOC#Ir|eRzV=D>>+5r~*1Cw?bpifbovRD5f0yb+@OF#-f~dgv`-Q}3 z=TygyC8?>_`*Um&jlnH7Aa`D-76{UMiwT8a9W3u!*P500y%x1Ezczf3noC75 zO4KUlQp`2nkPWh<7LT?CSUU$Sd6_S+uFK7&+FNqkdCiu+)q20Kzr(_^hMx(AH$axd zcBY#10;jkrF)iNySv~D>7*F?#ooN9j;;U~({c{Z|-?|=fZ~C?42SqpK|Kv&K(&NRj z9yt7)T*6+CdN>fqM?J@HdLaW?-ZPKgO)!;$q{e+0iMT3=32YS}*OX*Fmi z_QjiSf893Y^8>-~ci_<4 zjA##e`Qz7Ft{+eHZHGeHXs;1FvuYWMwup~dbEFbMcivf2>(|+7ksKr0!GRpG($dDt z_3-ay2_sS)tXO23Hl#3py*CqYbyDT6)^ahJXn;&b0E<(>R6gf<@(1mQ#wUL<36;cn zFgZ@0=(>{+?x#-L3O!%v9bxraG#S^ssaSD$Y+Kru_Ba<;{xd7m@nJIN4vCnTi)rb| zEhxtf1k2sFvNLMa_~^5~B+7zPv)>I)S*zJ8BK+|4w?_zE|fpfXO;h5U8Vb}vA;Sm4WAdEaR$$O`9jo-=iYzh z{*Pt!o`;6L;XSwfc-#dJyaPU9AMf1v5xbo9?)Ajo>5>~Vfo6t>`l-CVTZMU~1EWD@ z<`GR+k!j%0BM-}U<=DjrJX*m)wF$;gB&54TL>hL|a1E`fAf{sn9NMSM>_$jZ_m5LcfK!lrVz2`}^pBlspGTA*38(GZyx<>-m~z zMY`bzr=BNPdL=0i=SHC2{{F|9sx46Mg;V)|T6s37RR6D)=jw3_F5t}eludKCTob*x zUg>#tqW-tl`23Mr&8sPvOY82zgL^zZ-2Q|o4?@h*W(-d?*`>sWMz`nI`$yyy2Bv)P z0$WTs^LAIFq&&wBvlQ*0esTSAlg(ePX?MnE;9L+m^X;TM)?R9HSgsdw7)oWYX#1>%(k}^r6leEA9>C+gv)WpB1>>`{>X;Sab0UIkm|PX0i7Rh< zz3p-e_70VuwzuGLW#frsN(Dbj{0g2uCMZ}$! zQh45^dUIl}^JY!D5@x5MMr4h|uhiFFNt3m~9OFz19W|(ppK~cTHgwtMMtsMybg7A~R~?>UGOMd@DQ#Vf((d1(7V#}d6H32gcVP7%4;8ukMVBc*6qx(Y1k`wZ;>GkjsN3}8XAsbmCENTGW37w6wM|}|K=9w z`~C_pq7QMT!ir7t4BQUsF`)jn?m=(otcmm@9(71fl1Dnm~_e>AE(IElD?A%b=|a2udHSY^S}olP2|B~ z-RxC_^k(Vcv5G|kBv2G*(-nK~d6NcTUaDL5l>L~W2Cg8`Qztz7XXvv8Mow+LXCKF> zBY0x(;fxts<$EMmyN0#KvhBrE5$qfmck1PX zGOBXh_P*RneOz_sBbJCK7BP%o8noz7tL#C+e2P7!EvYFND!g<*gXIW#$H z>@c3~!a~CDXiMYm(|F|p^CunE=;6iWNBJYxMeeQKphxLQyLbJyWKYjIC5jQ>7EXJgw+dQt?&FsR%sNJ!8o)9eLz_8Iy3Qcy)(U!V@9VoaN z-C_W&$-U1F@imnV5>rubj;&;ak0)Q`6u7h?L__ES-r?KBD?Ml{^>nA#Ew@F#YoVzv z*fZ2hy)r(F)gOfzk>F%im7&i)j<>&Bmo|J~yifeO%HKg3xGbOKj)h$JQ zweX)3IcvWN1oiR-K$`LF{{3I*5x1FTHlRT9Xw&E}XEZUje|BW@ojJ~rAsl!!F3Y1s z?Z{MBQJ^wKL^=;iMpZtZ*zNN2l5ZkAzYJ}?qVi2es>?<(D%z*yelAdo4S5vFMx_Y+ zqe197T$oK`$>2djkGwEn^IXy{HaRXAdHah9zTZMf#+D0Q0lKb!XQEvh6LBN7nVLd&~%4 z`(Hz1>71^Q#iTdh6Gcm(j})rA-x#_VozYt?fihAB7xdoz=%4mOEI6JqJ53lxmv0hI z)6E1Flp)V>b`go)B7eJR-dF!?@Xsk9r3mt96(dvA9iOL8+Fg`k$;rmK2z9%JzeHUb z^ZH<&u3Hp$zpH!w>!eM(?knVB#U`31GR{KFLxw1t#jnnJqHuZ{X|`FwJJ*cOkw}!$ zI(roRg`RMaBjav!;8IQ7udc4WSG@hi9Ns5mQbjeVOS9IYbb4Q>-Y-Cm(Id}`*1!Xq zBB8k31$Lr(j|}QXakTbn2uBcLm;F<-l6J)=95NW+x>I1021KSV5JeL+Az{ouWOoOx z|9rX)%%~SKjptXhBpS@?99l21`x_^?J0$0Z^0sN)Wgj|>Ym^l~BuK1F?&|-jR+mLG zKANG8pU94@OF9& z4a>$e7A*+!Q0VjL-+a$_*SZ|HEq=wbvk^thy{vOucmy_Rx#8W~$@#=`-?7d*bRjM> z?dP*AzKXdMgBg<*1J3)_FNbcZt7U_?7l~NWAKU#}Yrs+Mxy5aRoj2;Fxyu;4VVe;` zz3vspK3h7Rl>sPwB0An*?7eYn4SY%&v~z!BH>lUB5r)~aMolz`x%Dh+BBt;o9uind zaUJF5sAlvrHA|9oS(D(9mF|b^MaZDLBgA>iS)4w<2adSW|F~fcf^AJ8;5Oc^_zBJX z4lq##Rn~6>>v;LTyE?akz*2{j=;@6^N*jL3yu}V^x~_%?0!R3S`?*%@G@%f^p2@h) zFL<`*N69|)TAWFhrK9-D_)U|=?S&tnp)9Fm2Bk0R^ulQah@)+UF}Dh z5qCK*>A8RXQuO`1CiQJQN&C--EwF=W{=`C;kMqFs8P8)Rd2=d zcx?Y5^H_gvL3-Tuwrru&6yyi&)X4N0%G*>vI8(LAIvw^lC?aBO^rCeQC0+KdE1vy~ zDB9ul{r$Domga$MAvwyAwC4ipF0bH!d9h~qf>vsfCL%njK9u)6Kf*&9)e<7 zT+g{nb|dQyTh*QwqJn$d-#%d?+Zcc1Sezl68*sX7iUF@Od(rnQWl+9d0g+Ih_>_hF zCUNQc9z{?sE5*0_i;2aiPOq)32K1-Q?Q3_3{96r*Eo{j6S(vIF#=xW(PtnMV;u*r; zCFV5Fs4i4Z7|4LHs>rDGjSl<*{ z>7~U3c6ihT=P@Uu#~;RYD@2B%b8{PqwrPcrV-re@UMqfPvr*4!Mmw%{(rKL0tXU1Z`rRI-byxJngNZGlnZB9^PQijRU= zcir2rP)iYA&9_vy*gnCyXTnpF+b@Up}LxK~M5BhK;mkgXj#^s%W4LfkO$139#z?+1Hf zU8q6F67!d0$MeS9rPQ|?d9vS<-dmTxgABmx<1^)_$l5u0BrzG(&nfw0 zEzrsDPMELU@wCU%w>rq~=dog~K4wAtrn7qu>~&|tOpLMdULf#T=^=cVa5i?x3IYYY zp6l%Fb4?3ewDi}sv~?3xy-|E*>85xnuhdF4OT#CdYutBua(VmW7 zx}Y*+jIi^FZl~IHB$xN9w+a{binc=Lzj&cuQ9F}4I$KgyT&924Qp-;_3LlkZ)S6k> zSI%+XdI45}U(xOC$TEbsFfj+#YHWi&4B}sqqY~fUG5SH`|D+6Ei(qDrWF=+kaFO{q+ z!N$FLGt<*dR$!HFo<5c+TJeI%Bj2<^CvAjh-9oFIN;Z}29izd9KUKthRtRIe7;HvN zlVZjExf;KheUt2S)$XY&DwX-}F+0vMxxvFN49ebaOPldC(8ES;4Y^{w8MpzWb#Zmw z5f2NG^m(pSk9mLgdTC9-1La$T^ken#53|TTx=51qPKE%z4r7ga2BhE!K2Xr&Gc6in ztt$rqp)&JPpP!ur;0rLjz;=w$jpMjm;l>j9x`wlePXk*yD+f+fJe+X!2sx}u@jCBA zw?K>F*j%1XFRpN}I?Naam$b4}HN2Gh(I3gf6E5$w);j+c5A(#kapvzDZ>Bbf^FusH zGT6OKubt+$vQ&K6dR7#~+h3mN!ka&=%(^>P1fzAFt4Poo5(*uu_{RA;gbC;9ZGSIP zaq!Oek7NlZ@Lp-M#Ll|m7K}cd{)5+Fphgd^$qi6CNx9k@gMc!BX#;j_W!RnFp*hI7 z#$l5-ICn5z4LrSv`SkP4?T%Er z-VdY7T=3Qm;fP)nzo?>gk9HyOG2a9P%w>np+{YF@U)52?#lwh_)I12nf_7 z<{)EZGaM4}X5IcFWeC2gWRJQVg&cP7Z0tLXBU{Sc*B?vuv}k$(+1-pZU~!nnDU|Hx z86Dc#jQ_%dAv;FnPAaB^&$|nmp+tjn+_=%7Dp#g|W+?rwd%*!2%c*}K*$8;b?V*Je7BnuYNcvOG2*~<4SoMo3B{HtitIXj*8cx^BaW}e74HEtvl#-Z;J!R;o@DkZm}|GR-9PaLsOEB`0d0L-=%&Dh?6ywF!20K zOhTi>fnzh}l{-05HOP7;TO-8sM8p725MH;Y?ixBInN>50f!sMCA1dLZq_!uwg zKIwGy*qN<=9V;(NoI!Ur$OewfLCmUc7f*6Z_9_L3I<%a6F2|6$6ud;QDydpwTmVg#R*(}ebhtP7WkKX!$!!gbk}>)uVxbj z9=b#K~UBB7mc_whPYFVDr8&Mc(RxeWm=v;6N%l-hQpBjjBOO z{vo?t(rT`cgXO>kKDo3M(!IeK=i`+Dp19f3#PqflS>;_T#$;-;0oCe6vF;HG%`*Oh zMA7Ynk)y8Z3W%wONPnYwjAJ;;NDFQ)>zGnWBXi*T8Gn4S~ zhIY@FB{|kPHghs1u-wchc0cyx9&T~qDcO8QMh4v?&)&6K+U|60XR z&b`c)!6kCH79cCrFJ^@1Gj{QKP*%4lSWjVV(%WfDGFI zt^O!j$HO7gp#h_!tGC>1o@+K%*OIjNo@wV_DXx8MB5OqEEG^`&KOsh;38zFflNp_S zeJ<;x7+(%}IQeRb!)xDhAK%j%&USZbWG)-1P~|Vh@BTfMc)J=79tGi$!g&(OP;$Lz z4Z@!iF?=P`7!BH`J@s><5I!O>)zBsWeeV( z4A}%&L@0v5sh^#3LPLmtVp8aGRdzjb>diDq2lK77IOg<)ytbGpvlc|5#B9fh6-p$M z@9sC>AXERXJ1zQWwRZAUEIBL&@|-qE5a+*e-}DqT)$~V?$I`^X2Tmm*An8jE?n4)5 zAOCD*2$QXqtu!#|At0L=+$&2~Tu&6eux$%s7o@KsF^2&WiG0^{PsgE&@Yf(4W>I;y z(CX6ab;VHba-Q`|bov0~WN8R2!IxeCJP%1MPC#v;6BNC+{J|k#u1A}$?5blmJ2B6? zK`{ka5L}JBMHQC1{@?gV`K~a}!L;YVv*)VzM-wHF&Gvc!`?cVJWQzODOrw+behfdM z%fuY&gi$f!l6GtC?8t}P&e8=ct1@F9jKA!Ddzb83W6vcQ0p;$IeGgM?gGb;3(vv}v z0btU4m*!m6f&1}!Ws6)zd*N`1IcxJ7s6vS89TFBDQT4}#KS9*Q1k4l#i6aUdj8dv1 z=z^)X&(3gR7eKyhoyO%-<-87~Qk~BvA=0`aigy2nBu-q-3O_a69;Gs+vaLRtGMh5D z7C0`=u(Mmo_$CBH(g5Y0l>+vdL({XS!5rMrM|)(3AOxy@gw`qz1lgrGYIBzjz9(_1 zE2ut>^>pb0cHnpmM-LiE#2G7g^u=vC8S6Qk zMK+Ui$zIm56x`tL$wG_d2f2ICsZ(kq5)ZwoV8mpqEMYmjs7)W1eTI#e>Ge@xJ>&0- zaKSK@)3BE#rVA-}Z!R#=cD7~Vw(pCer_mbGG8W$rxvC)@-X~X0%+XZ4W1<;Em(itm z%U-alGWgWtXEnX#1xWO(qFP_$AElu=#y4EK@I=FZmPsyQVJ{`~Ci@J(o;J9yS<)^> z>+o&wHE_WVLxndU>m%79i$4$SC@+xvoDBq|-QXhHl^F!wOWH zN_XVilnW>oUF_vqZzq8MIv9_;o;Lz1YdGNe2`~xhZ)ZC=HluOdmj|c8`n+f>3&?0L zC1lm21vA8Mj&A*51!-XR2RbjGh7RQ8FnW519a>E8+3$$~M4HdPd_Z}55L8Gumo($k z(p=7t(7Q`{pxR_ScT*@GG44AkTUIN%omzG9jX1YWX~~XDQ z`V6;(#AJ0q7O`+~P$`MSAD_?S z(_Qw3If{n?GcfjlH zJ-sY8HIQvPZJATP`=`_Ie--bM_y_Ol?rR$%;OzX=h*GWqkO} z_>gfsu-{y!V|IO?`!?(G6L1II{HUN$UK+RENE7i;4c!`oSU%B*%s=%12m@3g9)M}P4s;2{lF|uP>hS>b8K}RwU zj7Wc65ezfH0LSZET1>i_mco~XCq z(!mUfkfKti^PC&J|s{z9kvFV=o+zj}aR+^)e z|MyxDcy>HZYmpA_rEwpRGs<#hWZB#nVNhKA`rrl+FiLSZ@`$&8be}9f*-**k19l80 zKH0Lkt&>Pmt&mXsWuM-Mm%A*4ZOjrSMZR*`r$Jf& z8Z17L9Te*o9n#GXXO!4~M3x6d^ynaqw#%BcSq=R_VK|w1?GBlk$1iDZ1`eJi?V&U% z|2q(qmsGC0s%^VHj-9L)vTd^>SLRZrtx_auz+vi1GRWClmd<9IvzuDVaQTZ0r^lR? zN)PQLl4Pj`j8&-nRYfoRRqu|n8k(6V#TxdUHyEJ22)h#FD$g0YuS5MW9}y$An+sOYVf*jg zDj572&xhg$?Ux3L04W3N48(EEl5WGl;74&WX?|}P4k$^l=slZ=oncR6nJ7Z7QqiUZ z$`eIsRV4$x-hMIuDIw$u-P7O~pk8jg9F;5y0hO?l_XEZ=eary(ptf;T4q^kUeeFf_ zM95zAG_&qrm{cMwAlbndrq;;bXmMuD8Q%=f<4}Hu94w1bXoNHp&djSrA;vYRZdQIU zFh?3@Kx~pM8o~_+Uoy9MOy}{06#C0 zaX6GUZQC+%_QL8{IAwBUPJB&iBSxi&@yPBMvYjb%&lOd~YdKq*JIog*I!M%xRG#DE%Ih`gckTIQH`J`Qz%d~~uM3mXFsp4CBT_QG zUj~Ob)Y)&G0lII3Caz-4F>U@V6URc&qIG?WuwHp~v-aQA*;wMnA7nxaE0wp*)ry9-tBv+1DxxF&N zo8r@*lBoeuqW|hLR4p!44ZD=kr8Qeyt>Q6|N#DE6Ndn1_MV(D-jE6>~n;TQg7U6wP zCr2G0e1(crglB)n2_c3R2<#kUdZlqK-kpCdNH1afcjk^w0>c^zEE#8Mf~&=3i0Jes z`hKh@J%!QsAAa&f8~R^jzQ%u{q0OM)K!u~l@poIBa4H_=YObBm%+%_uiWVYT?~?J) zLcGt`$qcgDsx;T5t}_6Uz6En-BCEXufB4Adf6dsYNmW|SDgMU|L`&Gk6HNma$xUaL zS%Dc&txx@V{eSfjOLCjQQIqzQv#I*Ij&-FHoXXu*8iZMPuR-~P+DN&HpEBn z1x18jl6(j_R!8t(B@F>kt_c`>Jeh9n%8UJTBZRV0&tdEHP;vZ5YU^h#Vte8dEdrFf zKVW)x4b8oT??C>Po%zIx!+*2B_HbTW{Crzx`r+w`Rn@DrBdhoKMvNQW*Io_{{dHaG zPkOv0CM)eNfx98_2zR4jbguRM;SNbgGPk}iA-$;KS1{dQJa`UF#&MET=)Ige(j@=N z!ku0p$u9*aH*8A=@rzENA`%qBS7BUhQ?Pk4`pH+W&5rMH{zr78?;fZLJo{nHq@zm{ zHW!XVcYQ8Dn|gw9@!B5#86fR?cUdFaVXAoZ!FtvHE9Q%0wnS$v;l)3Ca`_(Syl)48 zSCI2BK0w1fH)R8~7u@|^pDWn_*~T^ShM98P$lso7lw;vu>QkI;84)I1A`@kZrK}5V zBUq1Y)?)2B87_~S7y)rfUhlO(M6Z}Gssjrf2(~s~x*`Wc>y~O2%jo7w>o9^zwr+Ci zbY4WmA6!^zyW;pCzbm>;AuFz<8;gKc!v&ubGbQ9;U5p(L*!_guC6!zNk9!nni%#=R zlRGl?16yO?C%&4a=H>!Kw%spd1M<{`$~DeQHlUI0nWhaV;PP8LT)4X zb|`(B5vbXZys%#^{whB4sXX_MLWY2@98D3h(fdhIkC3ib>cT<#EpYaCpCZk!MiMm4 z`qF@rJUr4haJ{dj1iYC#>a-wrN;{Fqx=fBPY^@$F6vA$CGOpIxPn{9BwjQsm+nj88 z<0I47%aWMfeu5n+bWB_C$K3Yja@8M6irOZzeUtiXogN^stjoeV{6lgL!49{SV;O$@ zFc;C;^=k*)QdacVo`iEFR3)Be!EMDF0nUJ1e(62$HbIciGKjok`dj-30GFNcK;vI$ z8cz7YEeZ1CfMBl*tMP+WtM?cYrO$WLu(!zTzNX@D4`?xamJi?FK{)+!A*0*_*TI|cg8E*PR)m`2ethLvfvx}~NK z`=%4n4HdQjg%ZSMe8qg;3pXEt!2U7X|4Dq`lSYtz+$c-4RSbC2?*Xs)Bn}U!3$4zPFr;P?ZlG2Ssfe4;fd+OWsw$ZHjWGw8^j}eGnjFFsR~$_pH2yhBICsZk7Kg@>WZW1?`C7wy z>qhN4w%Mc^w)(gsIapY$Xzj(J`1g>v!BpXX-j+Arw|FNK+*j%U%lrmW?Ei^JH1F9e z610LHmcBnEC*^znt_Y_bYIt-flN0+h5S#inRG3qOO1U{mepB|*X3ZkvSLd(dl}A_FdkXUO zOtNE`NU!PfOr7k=<<5T8pM2Z-zZavdQy?D}GX%~qNY6aT3>@@V=vOLhb#Sv8xEn+n zJlkv;{5xQa7eOILz6>l?K=HWX?#gZ)M;$-v2mZ3cq+I-l{e*tWKGha~n3;(X|SRo>U zYep|Qn6^Wc3S199{`Wf_;3CV^?yo0$Q|fj8CkhO`_g|$xRA-4l4fhieU;MH8<)Zvm z$I&LKFsGKP+{H8@S!S$JUzrq&HAM_4qD|_A429_4i&1T)Ryka#Ohdz7n;8(=EyhM; z+1FM6Zy5_+!p7qiZp*Qn6w|ISECz`4@IM=1-LHk%ZA{lX4-{>e<;-+UV>MVE*|mCqhY2$vbP*a z5kG8|083_~{Yh28kubyU+lN)*BM|6`RHjNNHOOqA|AB_-4J1GMrPrSyy$%1=c zPe)%wz*(`i_(ta$v{WnO(i{(RF)OCat7qS->5XH3^kc3L?Mi-gNi)f^wEpBZZrvzN zs=+~D)OR@JH7nK#jUN-)ncLQcm+W=Qm9mxlHw1;2Ru*g>$u`6Pi-Xi3)S;^fFqU+FPog7rwS^Dr%o`kW$dC0>2MJOt)Mr zSD+IAu9KPUO}ND1#CJuEe?t$~b94WeD#9(@_tes*ZGXdGeChuf&tIhZ+4rI#?>x&> z>|lM@WkF5Ftx~^}Ifo&VXc+9iA4nd`i#r`JyG%UQPfs^5cKbd^d`P$EpjRf_HZjWY+VZi{_YAX^m#s590bRJjo(eE`8D`wg$Co` z9xVa!`=`SZomP|_`HxB#Zs(;$Y}$waK>wC9P57VqJBNy?*+TuNDim=q57oy_ws^DH&>NXNe%uTv>VGq)SL`g4O>_H7iOeFlugW*bqWg zgxonk=G6d>ROV{eLaK&pqjFXR?}%^Z<1*bBm&!8=$ZLs7SedGg<1NNtCZ++Cm5Y%#c##Z|Hb;;!jn9U;)@?mZvmZJ1XpNA zZzUVz-=GDbz6#91SPME-~*ShBvlO!<8SbbD2#^0|^p)1MNk?bYme)5)3(S_NzRoBJ! zfpYASJyyUCsZpsUz7WvKG#N-0H;)})26)d$u+;;s9!zz$@O}HQWKd@UI(Xh~wKsB= zkL<4M4rn)*+!o2c{ghPyL5*G4+b;>pl_squu9EZsYY|vw#y%4QikW2sGBzi$_s4Qf zxDP$%9u^tDk<-t%_^`x!iV+#ggFmscV|Sko8S0O|{;vPm=1Tu}eRq(%#yvp~HaIg`(=nrW*&Am;~Qdoa%5wKW@8j|~x z*C;zSxpin|s7=DrG#{Du4y3U}k?rx{=QScTH6g2UGF%RNRKbj+jN*XK^vWI2tX-;5 z4;lXy*{yEb+4W>q!HTF?u&G|_G+?Y(EaJsn>m4aEO+b%+8Oa2wr|3Oc9s-Z?2XcZg ziP-(|wFvUNns^l^{_{fnCzoK@aN+8pBy`1=IHP@Bn*Ev}MyNU1;hz$F(k^)|@Iq?zh8&%V+}!)V@K!0u&VP>M0WHp5y|Nw;z|TUo+MN zZ}``{((;@NX}f&`3$$~%R9};GCAN0#+v-B?pnk*8^Ez(5!j-X-cWeNYXS`6DcyR8E z87}7t7C_c-{hJc2>f90dxs{XeZSG*e0y9EsKmXNVv4 zw7>(^%RB5zhkH;5dOn4|K^os-uDhE@;mr{Uwqmoad52Y1Q~=+vih@BYd5*HLJd7p?wP0&g2eOl3 zzAs8Tds@W1biz?dJ?R(=#pNqS5i^C!CL36n-N{{)PBG(;wht3TX0Z?LKO`x8KYW%p zi0#OJVem35zH5HngB;pu;isk)*|WdM9qChfTXkEavDNFHJ+kP9hpS^9_H z)%uSKn9!s_yNd-Qo(Az~=ct_hF@ZIn8xm6$-rMff7zk8<@glj*!mECZB1lm^UqbEZp!Yb;y6DKpi(%xFxct#D$ zp@P-L<6~~vj~H7ET;SZ;f=sFPpWVGJpS+L(Ht}9RSrZ9}j9-yhdpcRi6<&@xbTh&9 z6O{4efU+ho6C?>CRE5tRKnDGQm)E=~aN>{KClsJO=IYln+CfltM4B&93g*)n%v`li z!m!SofDEK@r~4a!$0Ojors}V8#VPkjsW!XsA5l}?Ce0*2qEH#DEdV;a5rrMdLdvPY zkaYT&@Hi4>cVePVoUx{!E=JPiU6mWbbgV!CGo7UO0iC_qLd&Hbin`5=^>$PV&Harb zxBT^xAvYqsw4qreVh)8eo;`H#zfadq<~}h1n_kI^wQt9Hwv&t}7nr-Ixq>`2O8({) z7)#}$(iI=rTQohbR;1wlptFgQ3cnbi+M<1f^l&4GO6BGTv!eCNwP8P&5pY$0;B3yf z=CL2UCV8nm$Jw~b3k|9K)dtFk`IZT}Hw_9>D~39L{LQERpX|N!;W_c#as!P`N0cY& z4ZYt}Kb_a|%!(8QFPC8hYwwyMZRdNw-@Q#vVf^UCmh5?4Ol(p;3{2G^LD}!N6#Ec4 z0sra7mHVVRJl`j(83Rw`or@St?MS2jp9_dPdr4?b1}0Zp3;lFKR`NOl*#$<)dhC1C z0KhNy;@MNZMQe6V-TL{+0#7Gr?{r;$e_zwcY^EzSnF(f8pLQ(yw7}?u{Dn_$Nb}$3 zCCv6Y2;vfaP7X2dyEL63iNOdnx@7?=#yiNr^?~T6a zx@w>rcMzbZJCxvzlVDKO$}vvClwPd?KFA4VF6EK&oGx{&tPu2TbUNRis*_zEA3pVQ z7@9YFYSz$8qscgKcTXb_9=tQGjJxC5`aMpTn5T6`a%ht#hd5r5X0u{r$P%LYA6Luu zt>S9sF4*^ep^1+`ei}mzR|;RH8_HDRtvh>eTeVa9%-JzYr}ki9ac{ z5HI?cuiyHpVGH^BW^qOuQwpuo;|66s7e8&IxD0?nY!p72N4AC9B)5Qp^>SJ z;FODsZ~CpklL-}7^%+@AvpFfVUIjQbpE6h8xR(v>zeyvQw?Q3r^s^B(%6!|?!6X)g zoEE1pLvEC4WQv#s=XmLrXWE!5j@$;|t=1nmryEF? zX_+d;tR%mN=F8ae12tyvH@`k&(TzeNEE8~ecT+a7Icm_^>cMBR zB8iec1iq~UYl4fn6WuV|F0uZ*Yth)`vHixt1s??lPnCH#1EK6kySs$bsRA=;MWT4^ z)xob;pzw(=3p4P^0`IUA$NnMv#J@;y@Japa{g=cvg>7f>DPs3yrkgP2M_$N7y`>?u zMvspm2wW^=;IM1xTi_9g4v%s{haU+-hb!ABet25tm#nSuX7XAt!Au>@F&80{o1Zj8 zFqsQrwWfQbx~wVxx_vo~Z-9`pF0tU>YN=7;8^P2SXER&@y3b!gYm-ci6sn|B*4c_` z7$bO=wa7+MZK!YoI`F!!+sA+QEv{xkeuEP#>G+BRp`heo_wk6tSoBq#{qn}hW9w;Q z;;)bLccv_`tye$|J6Zh;gf9lzSVWk9QC(sOo#N!=@gI`~6SNKvv`@+7Ex^7iG z)KEVKo)X$z$f2i^|5>laXmS=xz1bY70yYfD|3f7vUj!QN6PIrKl>DVMJ$ej=;DOKx zV=#5sMDFcmayPTBCs2spD6`5)6x4n_OW7^^hQkczN?6% z%n_bQ>x7P{VY^{H)k|47fI^zbTim~Lgfk>=r4UclUi6$uz7%WvblTGU<=TSK7;B+# z-NFXVl4+kbIsV-)*esLK6Mh^5UJ>XHsXI9p7x-&0$UOZVk|@yh_(W-K+Je>>Eijg; z>VF8fE?nSi%!pz`K57Vgi0Rs>^R7m(H*mCB zO*tt-jLF|md)$hiLF(h2ddwUnzmMRDt*Y45QSaOl@RHrZjLr*VtUC4uG7IcDm=vM% zW4$4ip1v+wqdJu3-kWkL)7Xkx+Fdt6%)UaV3HdjHS!(y}sg0Lq&t3WdpG-VY0T^WR zwgbh<`@#Yn3WrVx{Dh%a7gnj4)E_BrRm!{uzuI?U5S>Sx^A*q9)2HTYkP_d2qf&SG^LeZ=I)mq>95$ z%cD_NB1OpwFHPyIKl`Nf*&Pg>35z>V;&dN?krDQ$ywCojA46#1_>(V@=dz)ft`Z{O zZI(Ij6i20BYWn*O*GxO_T)VrJ%6-oP>>8QGg6b9g>O1I|>l062CzHxB{Sa6@Ul&;|#s|n%K9p2`0 zrkVH5(qiL$?6$VcNM$)*rs7K{C}uC1pMaX|8BzQsMF!cS+Tc2aO7Z65e|miIydI99dx^V@23A1pi-9s9|n-)BSXg#7f$2tTYqdrV1J>*PTuP_*z=739<0jqh&2 zj%kKr-K^E~H#1$WjwY-}4Y?fC00-F5bk zQlnx7HUp9|xhoFivAI8)#DRm8vj;O-zB@6hm=y{=s$vfmL1zC0B1Lg)x_ivs1~%1m z&$-H9*w<7&<7zT-Zg(#5Te4Gg-ILtrXTG2r@>&PG#M!sE)Iq2NgD7rjmIsU)W8ffy zEV9`jW3|x4I`gZKnGG4*E=_smUwlAs7pa~Gn(a!Xxz2S%V&fRu5wM)2>UW~#9~SbYB55lmbDt+UcO!)s+U_JD1SiCmFz_qk+->yc$Dipq}Hc*4Ua7B zoq}2Vxg9)5Z8YKd9?&a&X4^31`quI56{uH+C6=gH2+;Y!V5j#<`g4WSmYv5Zd3rKA z7G7O+>?h=`CTC3hr}isQskA8l)}J}PQSKLHp2d6{uCtuKKa6x7I0-ob*~5Ksr^q){ z)eCszC)g2JaUG>PPYK`24Zd|NUKJo)i&sl{ZISJBGT0YD=+HDsnaruY`!bp&qpM}5 z1Vi@JYg)J4WF=tfwh}mPK6Fo?Q@o~=m)k(LUd-QanB?6u+`N~8*b}tHit&H1wkOk5 zB5ov?=t6ia_LMkZspyOfqD|Irl(GO%uirGh_}v>nN?ZOcz~eo{&@6DjD=Ep%U8a+; z^yZlzIIv)f_Wo>n9r*DFS*i7xTtQ-^_lJ8pBezE61}pQ7?4L#DKK3=`3W>_2tcIK8 zE(XkHxQ5~itjM6Ki2+chAY<8Fq46e;X9W?q2jhoLIeWcmMQ!_(GY0!#QGO?k+`#DW z)x+tKiL3egMF|(Q5Z!f6nmmx?|DZXQ|9YMaRX_fQRwU*i<_O(K{Ce=fOLlu4q(5wl ztVe$=r~7;rG9LS5?v+!H>Yy=8*!~pV?_({DG`aV=(fcj>Z2ugH=$yQB?s2oGid$3s zRE~=~gW!)tsvgJj8@d#C*UjFf2-5pxI8=2p|HIqtD>>Bbr4>2DDR}X8KWSL0_!b^$ zMZ19=mbvNnDXg(PZ8~w7-(p=mG&LkK6O&FSM+D%k{#P3S_^|tMk zLTS&UVC-2&K?n$HZ}JvQ;BbL<7hxenl3^0CG-hS=uf*jv^tJa1+ctf-B#A8I3`ujW zo4hJ1)98@O|LiKb3nlRIPDC7N)8^*1=wo+8X%n7NWN%ttHVP*{hL0AfNQD7;Ni~79 zVxnM!AJxdX)n@q6gOqQgEj$%}x@i!F2K`z?pg%_scG@0y^7nn&0*>aRq4|4-H$_FbQ|b8 zuRF_20>%x24PUVXPE`(b2XD8Ds(*PGe+R8^{XlFO%QrY=$s1Q|#uf_LeMSck9IzrP z(c4(!P{T3!D|L&Z7&*nrPuGFf6O;+9McZMM&>vU1#|w9`Md9ew?SQ#d_SC`?7hDti z z=leKcYqFIKO#E~;0d1crja{5~Bs|RlvH8TJy4ilov!`v4Z|@jjf~$R#n|-go_V^2? z6MGOE=`UDhOfQoyGoX&;@Sl_7U%|4S!l{b5t?{99*C}qj1@E8BHA0<$g=7a^b&xy$ z-JE6R?e$mqSDwL`De6x9VQAvJMm%FGP8BoepDBXFC-^K`w$W1Z4oO44(FiJ)Ki);Y zCw1=p%|gX()X#52K>Tj?j1^g@$km5?R0_rX-ji$C|8YPt|EP$bgtevCxv{)Z=i0EX z)Dh*Gy;Xj{lHFEPo?V3L1ZfWUHrVS37qo|DV>GHjnbGKHO)6Vw!rr(cj4yJV#iDcH z9#z19Z=eUCwTS9D>Wdy|MAWu@{yZmQ`vu$OAFJZixAY)}oA zci~>R%7Cx?%A%kpn3jm?BHE!jPt3xwX#zX1PaZvXOl5fC#<7YE+ZE{Ln7YS8 z5qHC=srl+|m|I(veasiPUmCS~@g=L^ZvFS}0PR-#O3~6ymU6cEEj{K#f zgVTXLpZKv<#?JWC-uKwPz2QG*C?=N;HC1q?G>iS=`ISXF#9DPBxM)>$&tarpqqLM% zsAJZ0NvPB%ULvZ_t3_haw6O9c&B$lBw+H9zh(e%K-hte|*R2_@N4}5XrorsxOaeN~ z0uq9#?h=iv3Hq&b(Ddta>>@#M1NsnH59F?y+i;z9<0<*vC zfGP1C<15@{vdizBu|mF(4L9uhA|tQ&as#U&j4FF%iK2a1I@f)MuRW`R+S{b0GUw1|4W28QH{Z5Ek50PpY|6A8G zYhc&hM@eoZ8CH?7+#ifv6I*?tPj9GDD0eAsYXBQl2C{X7skY z>p@bVwU2pCR%b`aJ1NbQS>p8@VY6MyQ;SZrCpgRmI@L)C8j}CJj(}l*n~KQi<(BKxAd7P$!i0-ie zlI57+UB?{ooHX0FpYNyDM>pIK@W(R8Ndr6u6PY;%Eu-Dp#C0wQZm(C-1vg|{Sy3jS zebaY#&IqM`k1P8!*Wj@9!<5mQD-!~6IO757w~95yb2$LyR4C}xZ}v}i!+O)UMy78G zpOJI#9KNulBL$t$gkq*&ySu8IQ#a{w;k>Og*5_+NX$yB?*q{22Jw6?9v+=pTG6PfX z%ijI!$e7-baMWA6?h__hg_POP&yT&iz9F=WI^F)6(p}XubY00{yG=RVEme9;;g8+O zpzT!L5}o)&7}tI`Pkvg2q;a#D5<1e=Oj0BGyz)+ucpZduIY&KWcW6>j5_*0uWGn^t z@HE5p>M21-c8czWBiZEsL5_``$3*6{jHl(pmcK8VLOh2L`}in(=jI-#NbC%);RPC^ z^p_^pn6$9#m@k!}jCgC6X|WIN*UI{e<=d=Rfau)&p&*7YZwbd8ppk zF&b)2FfTbznubVTB(V=8JAMv4k8Fa2pGd9 z2C+a5h!q;zaLcBi8(E9S-O#nJJ0NA^MAZYunD!!)dQfb7 ziM{~gFj-@0P(06HYJ0%Y=rYP@X-Yw>BT*?Ygac{UqiE4xZE+R!s#w+l>Z@?$6GN-q zIrpr!=c6Pvqgf=-MM#G&cY#A3$jqtcexRvB8-Du~ir}?vdHv9T>9@Qk4!V$rURm)O zNs`N>97Yt&^tTu~H4HNYCN@pp`_5Lk*t(eRu;c;x?_9z;v|aGa?G`~!mrJggOs2pA zJ)kgg9!Lpejt0nD-Cv|W#4Cx&s(5-(&HP-LM!qvHE_R9B$WIODwX|+;eF5s!`=q!l zc#pGkoaxhpF9)-jt36Cr;^Rjr@LhiOGCF4`eRm>a;uq6HiTw-Cr{eoNN@h2Q{O%D?aiNw^VHnW55 zEf0-(0sGVx1R$-bS$QdCDO116qSEF*SY!kSXX%;g4CK}vM2s>T6fwcu;P>xvz^?ph zjX4Ep)vXbjfj~hQrQ~`w?WHZy6ujs4@o@yx>dPMN9B;(IgFm$UW1X5EJGN8BbuR7; zwZYWV0GFC#3p85@+;p^nwrmnW=9eXS4{#1Nj(Dek`*xL`hf7mol?7scd)gq_2>k)W z%O<>ZInb`Et?YkN90$HQ&vYh-$(@k;j;spa*nviyh(t_8k^=n!zjy!>>n_btNm}$K(E zrv~9)pC>&|pyOU*ZPo6AjEbF6iy)MQadmS86CrSl2x_(m_zHmqQoKgC^XKr7Vq*sh zi7j`_98yhp-h>)g%y!`yoX$Ex9im);o_D;wzU~CUnX~U`Q=#7v~o8!K{9fyqOnRP_gFUgg5{L9nqS}JO_{D&fCKRV59Ab zRfF$O90X}GQV&!|OXpn(0vFOyc}rmbH<}+BX;k8O9-tWd{)&a3OYbNij&YiO!*;~# zA2qAi)NeEZ19|RZ0F#}VQOSH=!&P8-C@;Xpt-N{J#Bf|fLW05jaX%zv&*$yxu~81t z28}c6n&tjL8Z016&zGW-B=9>kNy&~XQ^Ny+g)K_nxvMrvSn`waQ#As&syeAFT)>i#F@*Uk@ z8wVhue>`&X_eX%RvAZYM2`rmb8AL?v#NF8)1~ zp4nQ&Y4tt#q8iuVU`dmUmuw~ zhzO$q#rs2ZooLUEdar(ss(EkHD+kZu>t4kZC|ZC47|0*TnYE%Ohi zDGTp_C`YX5`z=)lV<#W5nV4`MEdOxBL{afG?I1pr#$^ci_x(Ehxt--p)Gk$5 z1H1J6>)qr?xJIhyg8kbb5;7!Y@jhkEH@eO_@G*pR&!1s@94EkQzgn7#l_=SOYT}SX zuxm5R6|vE?zii#T`Arz5RvSa@#U5b;5`UPS*sD@v%nV$lbTggj&MlQ;4ch1Wn+I1} zns*sZdx_d)@{{#yi4e>8zD;K*{O;d1&8xdD^;_pCU_GRycLUaA!}7~Wm2%-O?na4( zJlEN}guFW*5+?n=yOAi}$@fzrd|Bf{q7YwS5P8_d%i@0Qz+Wh*1Q82QZlyS=QZHw_ zCC|OMMB7^~sfTSq5>+5@4HuFQ`|`z8?&e#Hw-0LqPduV2`6e<8z)I)#t2Q~+rJa{K zcD(Rn&&NT z#%;gH0Z+Ds677ku7SFRvfK**w9W@Pq#@Mc7t)1hPWCV`Ae-4=bc+-iuQS}@rE?*v# zKPn^Pet%?&h89-d_{2;4U=TiJ)Gf-oH|wZ4o(XCLvL~dRnrRjH)|A@zLeK@aot+vc zLmRRvcz;`E&t6;m;9NZXsUIk@AC4PXsekEQXFzuk3x0IJWHomXr^#W%VL8GPPaTeR z6gLobZ-hU?-BVC6JKjf4R5DR~{A*bDChMUEThk9+s7Fnr6EhN(6i!Hmn$E7TGXoqj zZJ`@{u#Mw=uwr^m=t2Tnp`cDLC*+{u55JJ^j17r~!TTN@5XA8`CujpR-w|t=n(8)7 zpEVa%a8FF2Q#Q30(>s&oQ`gt;k}0qoX?0#{GAX@NO9d{YwtTNP0vhM50PY|k=OGcd z*?UW9H=xrJbJCEWT+eS8_3Ae$(s8$GSu)}kwetI7hlXWnH{$JHF%SSVO*~nocoXfr zDzwq14&IlOU1W#ywMQD>h;Q$#&zu2d7?75hTkr)5U5Hy6>|Xck{iJz+jZcF`O&-10 zZ)r(wYj0bMO^6el$yStsAYke;ZdG1UgKF3DHV=f#(mwF3?{$xmC%wEx51uUgFDYnE*) zv6tEhR&rMT)$6;r>s?$^qktYAdXs@yylH4*aSvs`B|waft=(wb>;S-;^VrKI!Vh1De6VW z!?*1dU}67Mf7xduA)&7*+`K5+w4}aziDbehXjL#Ua5&3OzE{09Ysb48^^4sNgMF&x zpSxabCN+NbnxqMGa9caqcy(L!Gjri(2AHA#3J@K78`W&#Z_`{}~avR=V#>Vvhc;e22NoZdAjy=Igj? z8+_7HZ<>I|S{BRDaZg=JckpnkQ_kdyYEUNL*`|Y=J$=9jo6m>M3;TzP_;5hhtQn-{ zvPB>`^}gIGL4vX#BSl)DKj0Jv*XA|aNJzWlj!%AB>gfJUAxof=4Qa5iA&Ca+&HoKC z$lxV!Og~lqdXEy#ia>YTzK9i4*mq?0c>BH;7aUz*dkw&kZWGJaZX5Be2O%8*x%vA- zkld%fl@~WnR9f}%E+!-(PKfjvX~-4kYC2a=QTAr_CiY&0!PVU6BrGYBclbdwc=n%x zj$}f@7cc(j8h=}LvwsXQYG8EDK7#&{R}k>{EgEntd2m3|`G+ z7>NrfV&&RmAH{u{mHGQyvVZsQ=id?dcLe?qMu3G!s2yeB=fFCvMs)H2;Ew;^;@=VY gcLe@lkH8%+sQlI=jQ7KB_Sd|8@ydl#-5U}A1&eU3j{pDw literal 0 HcmV?d00001 From 19ad326f69df5cb53acbdc052108d9455729fb2d Mon Sep 17 00:00:00 2001 From: Anna Revinskaya Date: Thu, 15 Aug 2019 13:28:56 -0700 Subject: [PATCH 2/6] Update 20190814-kernel-and-op-registration.md --- rfcs/20190814-kernel-and-op-registration.md | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md index 32b0471d5..6dacd8593 100644 --- a/rfcs/20190814-kernel-and-op-registration.md +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -399,12 +399,10 @@ For details on how we plan to switch between `std::function` and `void ### Device Assignment API This approach lets us construct device objects (e.g. `Eigen::GpuDevice`) on the -plugin side. This is preferred over the approach described in "Alternatives -Considered" because it is more flexible when it comes to Eigen assignments and -operations. Basically, we get an Eigen device object and can apply any +plugin side. Basically, we get an Eigen device object and can apply any operations we currently apply to an Eigen device. -We could wrap `StreamInterface`, `ThreadPoolInterface` and `Allocator`. These +We could wrap `Eigen::StreamInterface`, `Eigen::ThreadPoolInterface` and `Eigen::Allocator`. These wrappers will consist of a C API and a C++ wrapper on top of the C API. A sample C API for `StreamInterface` is given below: @@ -546,10 +544,10 @@ TensorFlow core of the form: void foo(std::function arg) { ... } ``` -We can't pass std::function across the C API boundary. Instead, we plan to wrap it with a struct and break this call up into 3 steps: +We can't pass `std::function` across the C API boundary. Instead, we plan to wrap it with a struct and break this call up into 3 steps: * Wrap `std::function` with a struct. The struct will contain pointers - to callbacks for manipulating std::function pointer. (This will happen + to callbacks for manipulating `std::function` pointer. (This will happen on the kernel plugin side). * Pass the struct across C API boundary. @@ -649,9 +647,8 @@ get from a HIP function (on the TensorFlow core side) and pass to another HIP function (on the kernel side). Ideally, we should only rely on extern C parts of `hip_runtime_api.h`. There is -now an equivalent in the C API right now for `hipLaunchKernelGGL`. However, AMD -have said they are looking into adding an equivalent function to the C API in -the near future and are willing to add more functions if needed. +no equivalent in the C API right now for `hipLaunchKernelGGL`. However, AMD +might add an equivalent function to the C API in the near future. Note that we have to update `LAUNCH_GPU_KERNEL` in Eigen to call the HIP C API once it is available. From 1c87177cb21ad9e8e48a883b162ff0bedd6da42a Mon Sep 17 00:00:00 2001 From: Anna Revinskaya Date: Thu, 22 Aug 2019 15:59:42 -0700 Subject: [PATCH 3/6] Update 20190814-kernel-and-op-registration.md --- rfcs/20190814-kernel-and-op-registration.md | 22 +++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md index 6dacd8593..b07f060c1 100644 --- a/rfcs/20190814-kernel-and-op-registration.md +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -535,6 +535,28 @@ call ROCm's `hipLaunchKernelGGL` here instead. Its signature uses templates. Fortunately, AMD is planning to add an equivalent function that provides a C API. (see Appendix 2 for details) +### Getting Status when using device APIs + +Kernel Device APIs described in this document rely on wrapping certain Eigen interfaces, such as `Eigen::StreamInterface` to provide a C API. Implementations of these interfaces might set an `OpKernelContext` status, which is not on the interface surface. Therefore, I propose that we add a new function that would update a given `TF_Status` with current `OpKernelContext` status: + +```c++ +TF_CAPI_EXPORT extern void TF_OpKernelContext_UpdateStatus(TF_Status*); +``` + +This would allow kernel implementations to return as soon as they see a failing status. For example: + +```c++ +TF_EigenStream* eigen_stream = TF_GetEigenStream(); +... run computation using eigen_stream ... + +TF_Status* context_status = TF_NewStatus(); +TF_OpKernelContext_UpdateStatus(context_status); +if (TF_GetCode(context_status) != TF_OK) { + TF_DeleteStatus(context_status); + return; +} +``` + ## Appendix 1 Certain parts of our design involve kernel plugins calling a function in From ac4ae7f2fac94e5ca23b7d8e9de34a4e0d09860c Mon Sep 17 00:00:00 2001 From: Anna Revinskaya Date: Tue, 2 Jun 2020 17:51:18 -0700 Subject: [PATCH 4/6] Removing device part of the proposal I removed device section of this doc. We never made a final decision on the device solution and then I moved to work on something else. If we come up with a final decision for device, we can create another RFC (I removed my name from authors as well since my only contribution was device API. ). The rest of the doc can be merged. --- rfcs/20190814-kernel-and-op-registration.md | 384 +------------------- 1 file changed, 1 insertion(+), 383 deletions(-) diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md index b07f060c1..948f74637 100644 --- a/rfcs/20190814-kernel-and-op-registration.md +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -2,7 +2,7 @@ | Status | Proposed | :-------------- |:---------------------------------------------------- | -| **Author(s)** | James Ring (sjr@google.com), Anna Revinskaya (annarev@google.com) | +| **Author(s)** | James Ring (sjr@google.com). | | **Sponsor** | Günhan Gülsoy (gunan@google.com) | | **Updated** | 2019-08-14 | @@ -293,385 +293,3 @@ API will contain classes such as `Tensor`, `OpKernelContext`, and Ideally, this API will be as close as possible to the existing non-ABI-stable Tensorflow C++ API, so that kernels and ops currently implemented in C++ may be ported to the ABI-stable C++ with as little implementation churn as possible. - -## Device C API for Kernels - -So far, this document has not dealt with the challenges of providing an -ABI-stable API for kernels that run on GPUs. This section describes an API that -addresses these challenges. - -There are a few approaches to running kernels on GPUs: - -* Assign computation to Eigen device (for e.g. see `OneHot`, `Transpose`, - training ops). (>200 occurrences in TensorFlow) - -* Call `device.parallelFor` (for e.g. see `BatchSelect`). (4 occurrences) - -* Call `ThreadPool::ParallelFor` (for e.g. see `MatrixDiag`). This is a - TensorFlow wrapper that eventually wraps calls to Eigen. For example, - `ThreadPool::ParallelFor` calls `device.parallelFor` in Eigen. (29 - occurrences) - -* Call `Shard` (e.g. `CTCGreedyDecoder`). This approach is deprecated in favor - of `ThreadPool::TransformRangeConcurrently` but no kernels use the latter yet. - (42 occurrences) - -* Call `GpuLaunchKernel` or `CudaLaunchKernel` directly, i.e. without calling Eigen. -(58 occurrences) - -* `Matmul` op calls directly to `StreamExecutor`. - -* Possibly others - -In all approaches above, TensorFlow core is responsible for maintaining -respective device queues, streams or pools. Kernels then use these queues to -schedule computation. Therefore, our primary goal is to implement a C API that -enables this scheduling. To give an example, one approach we can take is have -Kernel pass a callback across C API. Tensorflow core would then call this -callback. See diagram below: - -![device API](20190814-kernel-and-op-registration/device_api_overview.png) - -Furthermore, note that most of the approaches listed above eventually call to -Eigen to parallelize and forward computation to device. For example, the first -approach above uses Eigen APIs directly. Consequently, we need to understand how -Eigen works with devices and in some cases make changes to Eigen codebase as -well. - -Finally, we should aim to create a smaller API. Some of the approaches listed in -the Background section seem to be very similar. For example, calling -`parallelFor` in Eigen is quite similar to calling into -`ThreadPool::ParallelFor`. Therefore, we will only provide C API equivalents for -the following: - -* `ThreadPool` and its methods. - -* `CudaLaunchKernel` function. - -* Computation assignment to device in Eigen. - -This proposal focuses on these 3 components for now. Due to the complexity and -variety of TensorFlow kernels, it is very likely that we will need to consider -more approaches going forward. For example, how `MatMul` op would call -`StreamExecutor` directly has not been investigated. - -### ThreadPool API - -Here, we can just wrap relevant methods in the `ThreadPool` class. - -```c++ -TF_CAPI_EXPORT extern void TF_ThreadPool_Schedule( - TF_OpKernelContext* context, - void (*fn)()); - -TF_CAPI_EXPORT extern void TF_ThreadPool_ScheduleWithHint( - TF_OpKernelContext* context, - void (*fn)(), - int start, - int limit); - -TF_CAPI_EXPORT extern void TF_ThreadPool_ParallelFor( - TF_OpKernelContext* context, - int64_t total, - int64_t cost_per_unit, - void (*fn)(int64_t, int64_t)); - -TF_CAPI_EXPORT extern void TF_ThreadPool_ParallelForWithWorkerId( - TF_OpKernelContext* context, - int64_t total, - int64_t cost_per_unit, - void (*fn)(int64_t, int64_t, int)); -``` - -Note that we just pass a `TF_OpKernelContext` instead of a `ThreadPool` -instance. Implementation of these interfaces on the TensorFlow core side can -then retrieve the actual ThreadPool object using: - -```c++ -OpKernelContext* ctx = reinterpret_cast(context); -auto thread_pool = - cxt->device()->tensorflow_cpu_worker_threads()->workers; -``` - -For details on how we plan to switch between `std::function` and `void -(*fn)()`, see Appendix 1 below. - -### Device Assignment API - -This approach lets us construct device objects (e.g. `Eigen::GpuDevice`) on the -plugin side. Basically, we get an Eigen device object and can apply any -operations we currently apply to an Eigen device. - -We could wrap `Eigen::StreamInterface`, `Eigen::ThreadPoolInterface` and `Eigen::Allocator`. These -wrappers will consist of a C API and a C++ wrapper on top of the C API. A -sample C API for `StreamInterface` is given below: - -```c++ -TF_CAPI_EXPORT extern TF_EigenStream* TF_GetEigenStreamHandle( - TF_OpKernelContext*); -TF_CAPI_EXPORT extern gpuStream_t* TF_EigenStream_GetCudaStream( - TF_EigenStream*); -TF_CAPI_EXPORT extern gpuDeviceProp_t* TF_EigenStream_GetDeviceProperties( - TF_EigenStream*); -TF_CAPI_EXPORT extern void* TF_EigenStream_Allocate( - TF_EigenStream*, size_t num_bytes); -TF_CAPI_EXPORT extern void TF_EigenStream_Deallocate( - TF_EigenStream*, void* buffer); -TF_CAPI_EXPORT extern void* TF_EigenStream_Scratchpad( - TF_EigenStream*); -TF_CAPI_EXPORT extern int* TF_EigenStream_Semaphore( - TF_EigenStream*); -// This would just delete the C API handle for TF_EigenStream. -TF_CAPI_EXPORT extern TF_EigenStream* TF_DeleteEigenStreamHandle( - TF_EigenStream*); -``` - -The following C++ API will wrap the C API to provide a `StreamInterface` implementation -on the kernel plugin side: - -```c++ -class EigenGpuStream : public Eigen::StreamInterface { - public: - EigenGpuStream(TF_EigenStream* eigen_stream) : - eigen_stream_(eigen_stream) {} - - const gpuStream_t& stream() const override { - return TF_EigenStream_GetCudaStream(eigen_stream_); - } - - const gpuDeviceProp_t& deviceProperties() const override { - return TF_EigenStream_GetDeviceProperties(eigen_stream_); - } - - void* allocate(size_t num_bytes) const override { - return TF_EigenStream_Allocate(eigen_stream_, num_bytes); - } - - void deallocate(void* buffer) const override { - return TF_EigenStream_Deallocate(eigen_stream_, buffer); - } - - virtual void* scratchpad() const override { - return TF_EigenStream_Scratchpad(eigen_stream_); - } - - virtual unsigned int* semaphore() const override { - return TF_EigenStream_Semaphore(eigen_stream_); - } - - private: - TF_EigenStream* eigen_stream; -}; -``` - -Now, a kernel can create an instance of `Eigen::GpuDevice` using this stream: - -```c++ -TF_EigenStream* eigen_stream = TF_GetEigenStream(); -Eigen::GpuDevice* device = Eigen::GpuDevice(EigenGpuStream(eigen_stream)); -... -tensor->device(device) = < computation > -... -TF_DeleteEigenStreamHandle(eigen_stream); -``` - -Note: `gpuStream_t` and `gpuDeviceProp_t` might be aliased to ROCm's objects -instead of Cuda structs. See Appendix 2 for details how we are going to handle -ROCm support. - -Wrapping `Allocator` using similar approach should be trivial. However, -`ThreadPoolInterface` takes `std::function` and this approach would -require passing `std::function` across C API, which is non-trivial. For details -how we are going to handle it see Appendix 1. - -### Alternative for GPU Device Assignment API - -We can take approach similar to the CPU device assignment API. On the CPU side, -corresponding Eigen object - `ThreadPoolInterface` - has a `Schedule` method. This -method schedules a kernel function in a thread pool. - -Similarly, we can add a `Launch`/`Schedule` function to `StreamInterface`. The -default implementation would have same behavior as `LAUNCH_GPU_KERNEL` in -Eigen. However, we can customize it on the TensorFlow side and implement launch -logic in core TensorFlow instead of the kernel. This way, `cudaStream_t` and -`hipStream_t` only need to be referenced in core. - - - -Advantages of this approach: - -* Don't need to pass `hipStream_t` and `cudaStream_t` across the API boundary. - -* Supports customization of the `launchKernel` call which might be useful if we - want to handle it differently later. - -Disadvantages of this approach: - -* More invasive change to Eigen. - -### CudaLaunchKernel API - -CudaLaunchKernel appears to be a fairly thin wrapper around `cudaLaunchKernel` -in the Cuda Runtime library and a part of their C API. - -For reference, this is the signature of `cudaLaunchKernel`: - -```c++ -extern __host__ cudaError_t CUDARTAPI cudaLaunchKernel( - const void *func, - dim3 gridDim, - dim3 blockDim, - void **args, - size_t sharedMem, - cudaStream_t stream); -``` - -where `dim3` and `cudaStream_t` are structs. -This is trivial to either wrap with the TensorFlow C API or just call into from -plugins directly. - -However, ROCm's side of things is harder than Cuda. `gpuLaunchKernel` might -call ROCm's `hipLaunchKernelGGL` here instead. Its signature uses templates. -Fortunately, AMD is planning to add an equivalent function that provides a C -API. (see Appendix 2 for details) - -### Getting Status when using device APIs - -Kernel Device APIs described in this document rely on wrapping certain Eigen interfaces, such as `Eigen::StreamInterface` to provide a C API. Implementations of these interfaces might set an `OpKernelContext` status, which is not on the interface surface. Therefore, I propose that we add a new function that would update a given `TF_Status` with current `OpKernelContext` status: - -```c++ -TF_CAPI_EXPORT extern void TF_OpKernelContext_UpdateStatus(TF_Status*); -``` - -This would allow kernel implementations to return as soon as they see a failing status. For example: - -```c++ -TF_EigenStream* eigen_stream = TF_GetEigenStream(); -... run computation using eigen_stream ... - -TF_Status* context_status = TF_NewStatus(); -TF_OpKernelContext_UpdateStatus(context_status); -if (TF_GetCode(context_status) != TF_OK) { - TF_DeleteStatus(context_status); - return; -} -``` - -## Appendix 1 - -Certain parts of our design involve kernel plugins calling a function in -TensorFlow core of the form: - -```c++ -void foo(std::function arg) { ... } -``` - -We can't pass `std::function` across the C API boundary. Instead, we plan to wrap it with a struct and break this call up into 3 steps: - -* Wrap `std::function` with a struct. The struct will contain pointers - to callbacks for manipulating `std::function` pointer. (This will happen - on the kernel plugin side). - -* Pass the struct across C API boundary. - -* Wrap the struct with a callable object which can be used as - `std::function`. (This will happen on TensorFlow core side). - -Step 1: The wrapper struct will be defined as follows: - -```c++ -// Wraps std::function so that it can be called across C API. -struct FuncWrap { - void* func_ptr; // pointer to std::function - - // Function that takes std::function pointer as an argument - // and calls that function. - void (*call_func_ptr) (void*); - - // Function that takes std::function pointer as an argument - // and deletes it. - void (*delete_func_ptr) (void*); -}; -``` - -Note that we would need to move `std::function` to the heap because `FuncWrap` -might be placed in a queue and called later. Specifically, `FuncWrap` -construction will happen on the kernel plugin side and will have the following -implementation: - -```c++ -// Wraps std::function with FuncWrap struct. -FuncWrap get_func_wrap(std::function f ) { - // Move function to heap - auto* f_heap = new std::function(f); - - return { - // Argument to pass to callbacks to call/delete it. - f_heap, - // Callback that calls f_heap. - [](void* wrapped_f) { - std::function* f_std = static_cast*>( - wrapped_f); - (*f_std)(); - }, - // Callback that deletes f_heap. - [](void* wrapped_f) { - std::function* f_std = static_cast*>( - wrapped_f); - delete f_std; - } - }; -} -``` - -Step 2: `FuncWrap` struct constructed in this manner can now be passed across -the C API to core. - -Step 3: Since we place `std::function` on the heap, we need to manage its -deletion. Therefore, we wrap it with a class in TensorFlow core so that it can -be deleted once all references are gone: - -```c++ -class CallFuncWrap { - public: - explicit CallFuncWrap(FuncWrap wrap) : - wrap_(new FuncWrap(wrap), [](FuncWrap* ptr) { - ptr->delete_func_ptr(ptr->func_ptr); - delete ptr; - }) {}; - - void operator() () { - wrap_->call_func_ptr(wrap_->func_ptr); - } - - private: - // CallFuncWrap might be copied when it is passed to functions taking - // std::function as an argument. - // We use shared_ptr to make sure we only have one copy of FuncWrap - // even if CallFuncWrap is copied. We want a single copy of FuncWrap - // because the pointer stored in FuncWrap should only be deleted once. - std::shared_ptr wrap_; -}; -``` - -Now, the `CallFuncWrap` instance can be passed in as a `std::function` argument: - -```c++ -CallFuncWrap call_func_wrap(func_wrap); -foo(call_func_wrap); // foo here takes std::function argument -``` - -## Appendix 2: Working with ROCm across C API - -We need to access `hipStream_t` on both sides of the C API. Since its -implementation is actually in C++, we will treat it as opaque pointer that we -get from a HIP function (on the TensorFlow core side) and pass to another HIP -function (on the kernel side). - -Ideally, we should only rely on extern C parts of `hip_runtime_api.h`. There is -no equivalent in the C API right now for `hipLaunchKernelGGL`. However, AMD -might add an equivalent function to the C API in the near future. - -Note that we have to update `LAUNCH_GPU_KERNEL` in Eigen to call the HIP C API -once it is available. - From a926a55edc019972d90a0e310c81d86e392e02ff Mon Sep 17 00:00:00 2001 From: Anna Revinskaya Date: Tue, 2 Jun 2020 18:07:26 -0700 Subject: [PATCH 5/6] Changed "Updated" date --- rfcs/20190814-kernel-and-op-registration.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md index 948f74637..e6b0b4305 100644 --- a/rfcs/20190814-kernel-and-op-registration.md +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -4,7 +4,7 @@ :-------------- |:---------------------------------------------------- | | **Author(s)** | James Ring (sjr@google.com). | | **Sponsor** | Günhan Gülsoy (gunan@google.com) | -| **Updated** | 2019-08-14 | +| **Updated** | 2020-06-02 | ## Objective From 135293ee13b525e12f64c5606ff5b419d8d4404c Mon Sep 17 00:00:00 2001 From: ematejska Date: Thu, 4 Jun 2020 10:19:22 -0700 Subject: [PATCH 6/6] Moving to Accepted --- rfcs/20190814-kernel-and-op-registration.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/rfcs/20190814-kernel-and-op-registration.md b/rfcs/20190814-kernel-and-op-registration.md index e6b0b4305..155d39b03 100644 --- a/rfcs/20190814-kernel-and-op-registration.md +++ b/rfcs/20190814-kernel-and-op-registration.md @@ -1,6 +1,6 @@ # Kernel and Op Implementation and Registration API -| Status | Proposed | +| Status | Accepted | :-------------- |:---------------------------------------------------- | | **Author(s)** | James Ring (sjr@google.com). | | **Sponsor** | Günhan Gülsoy (gunan@google.com) |