From 868d82aa66e30898f7a938cdcade141e7175805b Mon Sep 17 00:00:00 2001 From: balbit Date: Thu, 5 Dec 2024 15:27:14 -0500 Subject: [PATCH 1/3] cherry picked kernel changes --- kernels/avx/matmul_avx.h | 37 ++++ kernels/avx/matmul_avx_fp32.cc | 7 +- kernels/avx/matmul_avx_int4.cc | 6 +- kernels/avx/matmul_avx_int8.cc | 22 +-- kernels/avx/matmul_avx_int8_int4.cc | 4 +- kernels/cuda/gemv_cuda.cu | 8 +- kernels/cuda/matmul_cuda.h | 38 ++++ kernels/cuda/matmul_int4.cu | 4 +- kernels/cuda/matmul_ref_fp32.cc | 5 +- kernels/cuda/matmul_ref_int8.cc | 18 +- kernels/matmul.h | 85 +++++---- kernels/matmul_factory.cc | 31 ++++ kernels/matmul_int8.cc | 2 +- kernels/mkl/Makefile | 57 ++++++ kernels/mkl/matmul_mkl.h | 37 ++++ kernels/mkl/matmul_mkl_int8.cc | 207 ++++++++++++++++++++++ kernels/mkl/matmul_mkl_test | Bin 0 -> 80040 bytes kernels/mkl/matmul_mkl_test.cc | 220 ++++++++++++++++++++++++ kernels/neon/matmul_neon.h | 37 ++++ kernels/neon/matmul_neon_fp32.cc | 8 +- kernels/neon/matmul_neon_int4.cc | 4 +- kernels/neon/matmul_neon_int4_offset.cc | 4 +- kernels/neon/matmul_neon_int8_int4.cc | 12 +- kernels/neon/matmul_ref_int8.cc | 18 +- kernels/ref/matmul_ref.h | 33 ++++ kernels/ref/matmul_ref_fp32.cc | 4 +- kernels/ref/matmul_ref_int4.cc | 4 +- kernels/ref/matmul_ref_int8.cc | 18 +- llm/Makefile | 47 ++++- llm/src/ops/BMM_F32T.cc | 4 +- llm/src/ops/BMM_S8T_S8N_F32T.cc | 2 +- llm/src/ops/BMM_S8T_S8N_S8T.cc | 2 +- llm/src/ops/W8A8B8O8Linear.cc | 2 +- llm/src/ops/W8A8B8O8LinearReLU.cc | 2 +- llm/src/ops/W8A8BFP32OFP32Linear.cc | 2 +- llm/src/ops/cuda/linear.cu | 4 +- llm/src/ops/linear.cc | 8 +- 37 files changed, 877 insertions(+), 126 deletions(-) create mode 100644 kernels/avx/matmul_avx.h create mode 100644 kernels/cuda/matmul_cuda.h create mode 100644 kernels/matmul_factory.cc create mode 100644 kernels/mkl/Makefile create mode 100644 kernels/mkl/matmul_mkl.h create mode 100644 kernels/mkl/matmul_mkl_int8.cc create mode 100755 kernels/mkl/matmul_mkl_test create mode 100644 kernels/mkl/matmul_mkl_test.cc create mode 100644 kernels/neon/matmul_neon.h create mode 100644 kernels/ref/matmul_ref.h diff --git a/kernels/avx/matmul_avx.h b/kernels/avx/matmul_avx.h new file mode 100644 index 00000000..8f663ad5 --- /dev/null +++ b/kernels/avx/matmul_avx.h @@ -0,0 +1,37 @@ +#ifndef MATMUL_OPERATOR_AVX_H +#define MATMUL_OPERATOR_AVX_H + +#include "matmul.h" +#include + +namespace matmul { + +class MatmulOperatorAVX : public MatmulOperator { + public: + void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) override; + void mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params* params) override; + + // int8 operations + void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) override; + + void mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) override; + + void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; + void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) override; +}; + +inline MatmulOperator& CreateMatmulOperatorAVX() { + static MatmulOperatorAVX instance; + return instance; +} + +} // namespace matmul + +#endif diff --git a/kernels/avx/matmul_avx_fp32.cc b/kernels/avx/matmul_avx_fp32.cc index 2dbcfda4..179d1d8e 100644 --- a/kernels/avx/matmul_avx_fp32.cc +++ b/kernels/avx/matmul_avx_fp32.cc @@ -4,8 +4,9 @@ #include #include #include // intel SSE intrinsic +#include -#include "../matmul.h" +#include "matmul_avx.h" namespace matmul { @@ -60,7 +61,7 @@ void *mat_mul_transposed_fastover_column_func(void *args) { return NULL; } -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { int i, j, k; int num_thread = params->opt_params.num_thread; @@ -112,7 +113,7 @@ void fp32_ref_matmul_bias(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params *params) { fp32_ref_matmul_bias(params); } diff --git a/kernels/avx/matmul_avx_int4.cc b/kernels/avx/matmul_avx_int4.cc index b5ee4c27..65011ebc 100644 --- a/kernels/avx/matmul_avx_int4.cc +++ b/kernels/avx/matmul_avx_int4.cc @@ -3,7 +3,7 @@ #include -#include "../matmul.h" +#include "matmul_avx.h" static inline __m256i bytes_from_nibbles_32(const uint8_t *rsi) { // Load 16 bytes from memory @@ -675,7 +675,7 @@ static void *fast_zp_no_offset_over_column_func_v5(void *args) { } namespace matmul { -void MatmulOperator::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { const int num_thread = params->opt_params.num_thread; int i, j, k; pthread_t thread_pool[num_thread]; @@ -693,7 +693,7 @@ void MatmulOperator::mat_mul_accelerator_int4_fast(const struct matmul_params *p for (j = 0; j < num_thread; j++) pthread_join(thread_pool[j], NULL); }; -void MatmulOperator::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { const int num_thread = params->opt_params.num_thread; int i, j, k; pthread_t thread_pool[num_thread]; diff --git a/kernels/avx/matmul_avx_int8.cc b/kernels/avx/matmul_avx_int8.cc index f1327658..6e524064 100644 --- a/kernels/avx/matmul_avx_int8.cc +++ b/kernels/avx/matmul_avx_int8.cc @@ -5,7 +5,7 @@ #include #include -#include "../matmul.h" +#include "matmul_avx.h" inline void assign_8int32(int *ptr, int &acc) { acc = (ptr[0] + ptr[1] + ptr[2] + ptr[3] + ptr[4] + ptr[5] + ptr[6] + ptr[7]); @@ -381,7 +381,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll(void *args) { return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert(params->A.column % 64 == 0); @@ -478,7 +478,7 @@ void *mat_mul_accelerator_int8_fast_32unroll_over_column_thread_func(void *args) return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; if (num_thread > params->C.column) num_thread = params->C.column; @@ -610,7 +610,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_nobias(void *args) { return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert((params->C.column) % 2 == 0); @@ -681,7 +681,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_nobias_batch(void *args) return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert((params->C.column) % 2 == 0); @@ -791,7 +791,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_nobias_ofp32(void *args) return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert(params->A.column % 32 == 0); @@ -851,7 +851,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_nobias_ofp32_batch(void return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert(params->A.column % 32 == 0); @@ -940,7 +940,7 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_bfp32_ofp32(void *args) return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { int j, num_thread = params->opt_params.num_thread; assert(params->A.column % 64 == 0); @@ -1211,8 +1211,9 @@ void *mat_mul_accelerator_int8_thread_func_2x2_32unroll_bfp32_ofp32_over_column( return NULL; } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( +void MatmulOperatorAVX::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( const struct matmul_params *params) { + int j, num_thread = params->opt_params.num_thread; if (num_thread > params->C.column) num_thread = params->C.column; @@ -1241,4 +1242,5 @@ void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over } } -} // namespace matmul + +} diff --git a/kernels/avx/matmul_avx_int8_int4.cc b/kernels/avx/matmul_avx_int8_int4.cc index eceaecc4..2e7bdb5c 100644 --- a/kernels/avx/matmul_avx_int8_int4.cc +++ b/kernels/avx/matmul_avx_int8_int4.cc @@ -4,7 +4,7 @@ #include #include -#include "../matmul.h" +#include "matmul_avx.h" #include "pthread_pool.h" @@ -322,7 +322,7 @@ static void quantize_fp_to_int8_block_size32(float *x, int size, int8_t *qx, flo namespace matmul { -void MatmulOperator::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params *params) { +void MatmulOperatorAVX::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params *params) { // const int num_thread = 4; const int num_thread = params->opt_params.num_thread; int i, j, k; diff --git a/kernels/cuda/gemv_cuda.cu b/kernels/cuda/gemv_cuda.cu index 82a85bf5..04e2d0cb 100644 --- a/kernels/cuda/gemv_cuda.cu +++ b/kernels/cuda/gemv_cuda.cu @@ -17,7 +17,7 @@ #include #include -#include "../matmul.h" +#include "matmul_cuda.h" #include "ops/linear.h" // #include @@ -210,7 +210,7 @@ namespace matmul{ Returns: out_feats: tensor of shape [B, OC]; */ - void MatmulOperator::gemv_forward_cuda(const struct matmul_params *params) + void MatmulOperatorCUDA::gemv_forward_cuda(const struct matmul_params *params) { const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; @@ -259,11 +259,11 @@ namespace matmul{ PROFILE_END("gemv_forward_cuda"); } - void MatmulOperator::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { + void MatmulOperatorCUDA::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { // TODO: remove this }; - void MatmulOperator::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { + void MatmulOperatorCUDA::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { // TODO: remove this }; diff --git a/kernels/cuda/matmul_cuda.h b/kernels/cuda/matmul_cuda.h new file mode 100644 index 00000000..a1d02c0b --- /dev/null +++ b/kernels/cuda/matmul_cuda.h @@ -0,0 +1,38 @@ +#ifndef MATMUL_OPERATOR_CUDA_H +#define MATMUL_OPERATOR_CUDA_H + +#include "matmul.h" +#include + +namespace matmul { + +class MatmulOperatorCUDA : public MatmulOperator { + public: + void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) override; + + // int8 operations + void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) override; + + void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; + void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) override; + + void gemv_forward_cuda(const struct matmul_params* params) override; + void naive_mat_mul_fp16_int4(const struct matmul_params* params) override; +}; + +// Declaring as static to prevent linker errors due to both cc and cu files +static inline MatmulOperator& CreateMatmulOperatorCUDA() { + static MatmulOperatorCUDA instance; + return instance; +} + +} // namespace matmul + +#endif diff --git a/kernels/cuda/matmul_int4.cu b/kernels/cuda/matmul_int4.cu index 507b564d..87e71c9e 100644 --- a/kernels/cuda/matmul_int4.cu +++ b/kernels/cuda/matmul_int4.cu @@ -1,11 +1,11 @@ #include #include -#include "../matmul.h" +#include "matmul_cuda.h" namespace matmul { -void MatmulOperator::naive_mat_mul_fp16_int4(const struct matmul_params *params) { +void MatmulOperatorCUDA::naive_mat_mul_fp16_int4(const struct matmul_params *params) { const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; // CHECK_MATRICES_int4weight(A, B, C); diff --git a/kernels/cuda/matmul_ref_fp32.cc b/kernels/cuda/matmul_ref_fp32.cc index 548d9926..d92b8a3b 100644 --- a/kernels/cuda/matmul_ref_fp32.cc +++ b/kernels/cuda/matmul_ref_fp32.cc @@ -4,7 +4,7 @@ #include #include -#include "../matmul.h" +#include "matmul_cuda.h" namespace matmul { void fp32_ref_matmul(const struct matmul_params *params) { @@ -28,7 +28,8 @@ void fp32_ref_matmul(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { + std::cout<<"mat_mul_accelerator_transposed_fastover_column, fp32"< #include -#include "../matmul.h" +#include "matmul_cuda.h" namespace matmul { void int8_ref_matmul(const struct matmul_params *params) { @@ -157,35 +157,35 @@ void int8_ref_matmul_nobias_ofp32_batch(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { int8_ref_matmul_nobias(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( +void MatmulOperatorCUDA::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } diff --git a/kernels/matmul.h b/kernels/matmul.h index 0424edee..8351c189 100644 --- a/kernels/matmul.h +++ b/kernels/matmul.h @@ -109,48 +109,61 @@ struct thread_args { namespace matmul { class MatmulOperator { public: - void mat_mul_transposed(const struct matmul_params *params); - void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params); - void mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params *params); - void mat_mul_accelerator_untransposed_fastover_column(const struct matmul_params *params); - // int8 - void naive_mat_mul_int8(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params); - void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params *params); - // void mat_mul_accelerator_int8_fast_2x2_omp(const struct matmul_params *params); - // int4 - void mat_mul_accelerator_int4_fast(const struct matmul_params *params); - void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params); - void mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params *params); - void gemv_accelerator_int8_int4_fast_no_offset(struct matmul_params *params); - void gemm_accelerator_int8_int4_fast_no_offset(struct matmul_params *params); - void gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_params *params); - void cblas_gemm_accelerator_no_offset(struct matmul_params *params); - void naive_mat_mul_int4(const struct matmul_params *params); - void naive_mat_mul_int4_with_offset(const struct matmul_params *params); - // cuda - void naive_mat_mul_fp16_int4(const struct matmul_params *params); - // void naive_mat_mul_fp16_int4_gemv(const struct matmul_params *params); - void mat_mul_cuda(const struct matmul_params *params); - //// GEMM - void gemm_forward_cuda(const struct matmul_params *params, int split_k_iters); - void gemm_forward_cuda_8splits(const struct matmul_params *params, float16_t *split_8_buffer); - void gemm_forward_cuda_half(const struct matmul_params *params, int split_k_iters); - void gemm_forward_cuda_half_test(const struct matmul_params *params, int split_k_iters); - //// GEMV - void gemv_forward_cuda(const struct matmul_params *params); + virtual ~MatmulOperator() = default; + + // Virtual methods for various matrix multiplication operations + virtual void mat_mul_transposed(const struct matmul_params* params); + virtual void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_untransposed_fastover_column(const struct matmul_params* params) {} + + // int8 operations + virtual void naive_mat_mul_int8(const struct matmul_params* params); + virtual void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) {} + + // int4 operations + virtual void mat_mul_accelerator_int4_fast(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) {} + virtual void mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) {} + virtual void gemv_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) {} + virtual void gemm_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) {} + virtual void gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_params* params) {} + virtual void cblas_gemm_accelerator_no_offset(struct matmul_params* params) {} + virtual void naive_mat_mul_int4(const struct matmul_params* params); + virtual void naive_mat_mul_int4_with_offset(const struct matmul_params* params); + + // CUDA-specific operations + virtual void naive_mat_mul_fp16_int4(const struct matmul_params* params) {} + virtual void mat_mul_cuda(const struct matmul_params* params) {} + virtual void gemm_forward_cuda(const struct matmul_params* params, int split_k_iters) {} + virtual void gemm_forward_cuda_8splits(const struct matmul_params* params, float16_t* split_8_buffer) {} + virtual void gemm_forward_cuda_half(const struct matmul_params* params, int split_k_iters) {} + virtual void gemm_forward_cuda_half_test(const struct matmul_params* params, int split_k_iters) {} + virtual void gemv_forward_cuda(const struct matmul_params* params) {} + + protected: + // Protected constructor to prevent direct instantiation + // Directly creating an object of this class is not allowed because of the empty constructor + MatmulOperator() = default; private: + // Delete copy constructor and assignment operator for safety + MatmulOperator& operator=(const MatmulOperator&) = delete; + float interval_to_us(struct timeval *start, struct timeval *end); void CHECK_MATRICES(const struct matrix *A, const struct matrix *B, const struct matrix *C); void CHECK_MATRICES_int4weight(const struct matrix *A, const struct matrix *B, const struct matrix *C); }; + +MatmulOperator& CreateMatmulOperator(); + } // namespace matmul #endif diff --git a/kernels/matmul_factory.cc b/kernels/matmul_factory.cc new file mode 100644 index 00000000..31c3b5ba --- /dev/null +++ b/kernels/matmul_factory.cc @@ -0,0 +1,31 @@ +#include "matmul.h" +#include "avx/matmul_avx.h" +#include "mkl/matmul_mkl.h" +#include "cuda/matmul_cuda.h" +#include "neon/matmul_neon.h" +#include "ref/matmul_ref.h" + +namespace matmul { + +// Declare external factory functions for each implementation +MatmulOperator& CreateMatmulOperatorMKL(); +MatmulOperator& CreateMatmulOperatorAVX(); +MatmulOperator& CreateMatmulOperatorCUDA(); +MatmulOperator& CreateMatmulOperatorNEON(); +MatmulOperator& CreateMatmulOperatorRef(); + +MatmulOperator& CreateMatmulOperator() { +#ifdef QM_CUDA + return CreateMatmulOperatorCUDA(); +#elif defined(QM_MKL) + return CreateMatmulOperatorMKL(); +#elif defined(QM_ARM) + return CreateMatmulOperatorNEON(); +#elif defined(QM_x86) + return CreateMatmulOperatorAVX(); // Default to AVX +#else + return CreateMatmulOperatorRef(); +#endif +} + +} \ No newline at end of file diff --git a/kernels/matmul_int8.cc b/kernels/matmul_int8.cc index 3d5f2ed3..03e50dc3 100644 --- a/kernels/matmul_int8.cc +++ b/kernels/matmul_int8.cc @@ -13,7 +13,7 @@ void MatmulOperator::naive_mat_mul_int8(const struct matmul_params *params) { float effective_scale = A_sc * B_sc / C_sc; int8_t *data_A = A->int8_data_ptr, *data_B = B->int8_data_ptr, *data_C = C->int8_data_ptr; const int8_t q_min = C->qparams.q_min, q_max = C->qparams.q_max; - CHECK_MATRICES(A, B, C); + // CHECK_MATRICES(A, B, C); for (i = 0; i < C->row; i++) for (j = 0; j < C->column; j++) { diff --git a/kernels/mkl/Makefile b/kernels/mkl/Makefile new file mode 100644 index 00000000..da6dc430 --- /dev/null +++ b/kernels/mkl/Makefile @@ -0,0 +1,57 @@ +# Compiler +CC = g++ + +# MKL Root Directory (adjust if necessary) +MKLROOT = /home/elliotliu/miniconda + +# Update LD_LIBRARY_PATH to include MKL libraries +export LD_LIBRARY_PATH := $(MKLROOT)/lib:$(LD_LIBRARY_PATH) + +# Compiler Flags +CFLAGS = -O3 -Wall -pthread -std=c++17 -m64 -DMKL_ILP64 -DUSE_MKL -DMKL_ILP64 -mavx2 -mfma -ffast-math -DUSE_INT8_INT4_PRODUCT -DQM_x86 + +# Include Directories +INCLUDE_DIRS = -I.. -I../../llm/half-2.2.0/include -I$(MKLROOT)/include + +# Library Flags +LDFLAGS = -Wl,--start-group \ + /home/elliotliu/miniconda/lib/libmkl_intel_ilp64.so \ + /home/elliotliu/miniconda/lib/libmkl_gnu_thread.so \ + /home/elliotliu/miniconda/lib/libmkl_core.so \ + -Wl,--end-group -lgomp -lpthread -lm -ldl + +# Source Files +SRC = matmul_mkl_test.cc + +LIB_DIR = ../ +LIB_MKL_SRC = $(wildcard $(LIB_DIR)/mkl/*_int*.cc) +LIB_SRC = $(wildcard $(LIB_DIR)/*.cc) +LIB_SRC += $(LIB_MKL_SRC) + +LIB_AVX_SRC = $(wildcard $(LIB_DIR)/avx/*.cc) +LIB_SRC += $(LIB_AVX_SRC) + +SRC += $(LIB_SRC) +# ../matmul_int8.cc + +# Object Files +OBJ = $(SRC:.cc=.o) + +# Target Executable +TARGET = matmul_mkl_test + +all: $(TARGET) + @echo "Running matmul_mkl_test..." + ./$(TARGET) + +# Compile .cc files into .o files +%.o: %.cc + $(CC) $(CFLAGS) $(INCLUDE_DIRS) -c $< -o $@ + +# Link Objects to Create Executable +$(TARGET): $(OBJ) + $(CC) $(CFLAGS) $(OBJ) -o $(TARGET) $(LDFLAGS) + +# Clean Build Files +clean: + rm -f $(OBJ) $(TARGET) diff --git a/kernels/mkl/matmul_mkl.h b/kernels/mkl/matmul_mkl.h new file mode 100644 index 00000000..903b0939 --- /dev/null +++ b/kernels/mkl/matmul_mkl.h @@ -0,0 +1,37 @@ +#ifndef MATMUL_OPERATOR_MKL_H +#define MATMUL_OPERATOR_MKL_H + +#include "matmul.h" +#include + +namespace matmul { + +class MatmulOperatorMKL : public MatmulOperator { + public: + void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) override; + void mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params* params) override; + + // int8 operations + void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) override; + + void mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) override; + + void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; + void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) override; +}; + +inline MatmulOperator& CreateMatmulOperatorMKL() { + static MatmulOperatorMKL instance; + return instance; +} + +} // namespace matmul + +#endif diff --git a/kernels/mkl/matmul_mkl_int8.cc b/kernels/mkl/matmul_mkl_int8.cc new file mode 100644 index 00000000..145eff68 --- /dev/null +++ b/kernels/mkl/matmul_mkl_int8.cc @@ -0,0 +1,207 @@ +#include +#include +#include +#include +#include +#include + +#include "matmul_mkl.h" +#include "../avx/matmul_avx.h" + +namespace matmul{ +void mat_mul_mkl_int8_effective_scale(const matmul_params *params) { + const matrix &A = params->A; + const matrix &B = params->B; + matrix &C = const_cast(params->C); + + int M = A.row; + int K = A.column; + int N = B.column; + + const int8_t *data_A = A.int8_data_ptr; + int32_t *data_C = new int32_t[M * N]; // Temporary buffer for accumulation + + const int8_t *data_B = B.int8_data_ptr; + + // Shift A instead of B by adding 128: + // We shift A because the mkl interface expects A to be unsigned instead of B if row-major + uint8_t *data_A_shifted = new uint8_t[M * K]; + for (int i = 0; i < M * K; ++i) { + int16_t temp = static_cast(A.int8_data_ptr[i] + 128) ; + data_A_shifted[i] = static_cast(temp); + } + + MKL_INT lda = K; + MKL_INT ldb = K; + MKL_INT ldc = N; + + + const MKL_INT8 ao = -(A.qparams.zero_point+128); + assert(B.qparams.zero_point == 0); + const MKL_INT8 bo = -(B.qparams.zero_point); // Adjusted zero point for B + MKL_INT32 co = C.qparams.zero_point; + + float effective_scale = A.qparams.scale * B.qparams.scale / C.qparams.scale; + + float alpha = 1.0f; + float beta = 0.0f; + cblas_gemm_s8u8s32( + CblasRowMajor, CblasNoTrans, CblasTrans, CblasFixOffset, + M, N, K, + alpha, + data_A_shifted, lda, ao, + data_B, ldb, bo, + beta, + data_C, ldc, &co + ); + + // Post-process the result + int8_t *data_C_int8 = C.int8_data_ptr; + for (int i = 0; i < M * N; ++i) { + int32_t acc = data_C[i]; + acc = static_cast(std::round(acc * effective_scale)); + if (params->bias.int8_data_ptr) { + acc += static_cast(params->bias.int8_data_ptr[i % N]); + } + + acc = std::max(static_cast(C.qparams.q_min), acc); + acc = std::min(static_cast(C.qparams.q_max), acc); + + data_C_int8[i] = static_cast(acc); + } + + // Clean up + delete[] data_A_shifted; + delete[] data_C; +} +void mat_mul_mkl_int8_alpha(const matmul_params *params) { + const matrix &A = params->A; + const matrix &B = params->B; + matrix &C = const_cast(params->C); + + int M = A.row; + int K = A.column; + int N = B.column; + + const int8_t *data_A = A.int8_data_ptr; + int32_t *data_C = new int32_t[M * N]; // Temporary buffer for accumulation + + const int8_t *data_B = B.int8_data_ptr; + + // Shift A instead of B by adding 128: + // We shift A because the mkl interface expects A to be unsigned instead of B if row-major + uint8_t *data_A_shifted = new uint8_t[M * K]; + for (int i = 0; i < M * K; ++i) { + int16_t temp = static_cast(A.int8_data_ptr[i] + 128) ; + data_A_shifted[i] = static_cast(temp); + } + + MKL_INT lda = K; + MKL_INT ldb = K; + MKL_INT ldc = N; + + const MKL_INT8 ao = -(A.qparams.zero_point+128); + assert(B.qparams.zero_point == 0); + const MKL_INT8 bo = -(B.qparams.zero_point); // Adjusted zero point for B + MKL_INT32 co = C.qparams.zero_point; + + float alpha = 1.0f; + float beta = 0.0f; + cblas_gemm_s8u8s32( + CblasRowMajor, CblasNoTrans, CblasTrans, CblasFixOffset, + M, N, K, + alpha, + data_A_shifted, lda, ao, + data_B, ldb, bo, + beta, + data_C, ldc, &co + ); + + // Post-process the result + int8_t *data_C_int8 = C.int8_data_ptr; + for (int i = 0; i < M * N; ++i) { + int32_t acc = data_C[i]; + acc = static_cast(std::round(((float)acc * params->alpha + + (float)(params->bias.int8_data_ptr[i % N]) * params->beta))); + + acc = std::max(static_cast(C.qparams.q_min), acc); + acc = std::min(static_cast(C.qparams.q_max), acc); + + data_C_int8[i] = static_cast(acc); + } + + // Clean up + delete[] data_A_shifted; + delete[] data_C; +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { + // std::cout << "Running mat_mul_mkl" << std::endl; + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_32unroll_over_column(params); + // mat_mul_mkl_int8(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { + mat_mul_mkl_int8_alpha(params); + // MatmulOperatorAVX fallback; + // fallback.mat_mul_accelerator_int8_fast_2x2_32unroll(params); +} + +// avx fallback operations +void MatmulOperatorMKL::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_transposed_fastover_column(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_transposed_fastover_column_bias(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) { + mat_mul_mkl_int8_effective_scale(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int8_int4_fast_no_offset(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int4_fast(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int4_fast(params); +} + +void MatmulOperatorMKL::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) { + MatmulOperatorAVX fallback; + fallback.mat_mul_accelerator_int4_fast_no_offset(params); +} + +} // namespace matmul + diff --git a/kernels/mkl/matmul_mkl_test b/kernels/mkl/matmul_mkl_test new file mode 100755 index 0000000000000000000000000000000000000000..b78b4a4e653ba6876efb58c682e104a8439c67c8 GIT binary patch literal 80040 zcmeFaeS8$v^*_E7HV`p5K?BALC6=X46l@|7%_`Ik?CNaYNE9N8f&n%`s6YZFLX-q_ z6Uua5plVxeZJS!{hpH{K`mxbk+)Z}!#8VQU0w{!lILm_sMA*cz`MuAb+08O0t?lRY z`+oi@yfSm|x%Zy?bndz5o^$T(8ms-jD2^Lo(EbfDeBZz!oH#}Wq-*encM??nmu^Ti z3^Ak|k_}0QfdCotm#!v#<=q;kCWPKGT0W|f$|Yzu>)#`^e4%%!c6w8}{{FSzs-}eA zhB!4&<)SH=NDur=WY6&X(93E(^o}gY>hoy%b?;fDRCwqeS&qi0Ls8T9@039*yfQQ( zT03dDsNBM{YP$ZNsf}mo9V$+3q`x$6JoWDs6`djU)|XqTmD9g1TDj1h>Y#r!J}Juk zck0X3%6WFGc=YetT027TP<;=f9F5D@{ihA5;bE;l{rHFCaoRkD-b5$#@4Jf^&7U;k zyNmP2Enc*=c-6R7X_LlHnsDcef;;bG>C~=7{1Hvv|6rB@)8ZkR)_gQC$g{sbD)SN;Oz?F)Vj8r~QDul?Y$ z^#lJ=Kk#S!srNPvOkef7`pNGBd?^0(|916*XJ9}58Vdx8dL)SR{opC?r(f@*OkeG9 z>jw|l51u>wY3IIv@R<97|F9qU-Tl=2Qa|vbe)`4tlm8U*_cbpW{j}$JKX@AZfiLW* zJ!|^GPdj8^^Fph%FZexxC*n{4@6mqRv!x$6F@V1Cf7VZaG~e(oeWi$}ApcOq2t#^@ znixkJ@$WbI`xdxP$^;fQXzb8n$jx=EnwPs^(b9Q~7d@SC$hFPP$j!@Np8v$66-D{W zXJ$-WT(C5M=Dhifk(T?!l7gkVD~jeVFUrkTv-X+MF@N#A6}eC3FIkeiBCR-W#rV4n zOXe+JT;MP)Se~D6C|a}x$R5pGocYueydRuVuy}=ak=BwYisvoQ%bn*~R=j9A<$Cae z8ATHvMJ0v#HitDWuON5jqP+a2RtJ)n&ZFWpA9*k}g?h1S-uy*RrKWxlVN|~&m&)WW zTDqXX`qf-_p~Q+sePm>#p|xRt!SW&$ol%rJDL1!h;qro&x%21cEW1uKe{ z=g(V$tSj=D7ug)BGu5$h-tyd{Vwz8)3&)eWj)hO=E||9njma%szG!LD z0)ild89~p~q(W4XPa|1?@-#Yz5afpii-qlwabgfgp8O)kxeL!;m&^HA4{<@1*2h0Chwu4$19Dqi~4I#yz448XVK z$;IeIQ9hoF3q!Jl27h8{F;NaWtBMpF*io=NKMZ)HU`ZiGQVH3LqCCePcTjW^MW1k> zwc$jE3cr8awA}G`rWo$G+on#-z3a~LcTR{1-_;xL1^0&Tx^tqMAv+sLCfsShCWe(d zEjwG7YRgRp+O%uZQd1++B2yy3Ku2;62X?0ZH$rk6i2sA|e;{BS-Z2PWhd+7~d@w=- z4A&zRi#LUbu$-LX29zGa{sy8X{p0Ze74XWC{F3?YlIsn?@D5;VdbHsUEqyKkMYk-v z&Hy|gXt3oh9%HE2;_v+AXArNV4g0is{GvreNWwm*#f!f|k`4+weG##PN*(Z1^qH)`>DEd4seC@sEs>yr>I2N~|v;zh|XKx7|i zFl+G=RzAvLMV$VHfkB&t9Tab;G4Ak zg*tep24AIvZ`0swb#RXcU$28VY4A-tc$)^_tb=!G@Jb!rppEAi9Xw8hZ_~jOHF%W{ zo}$4$I(WJUuh+pdHF%Q_K39Xc>fnVMyiEs>)8<{#!@sZ2ONS1=R>Kp}!8d7e!)di# zh|ViDc#ICdO@qhj;2sU0po2GQ@I)QFO@pWC;OW}AHcbbQ)7mZS;7M9MQwOiiR{N8! zgKyK|b9M00dReH0cRZltDb&H!?J9hg4jx)BYjyC@dRec7S8Dh->EN5Rc5c?en`Wr> zR_frJwDxS#!P_+WHXS@AL&a02gRj-#9v!?%gV*ceg%7HDnsjiF25;5D6E%374xX;j zgQA0HYVZyne69u$=-`DK+;B#1SBRcN>m^1H*Vap%4!%~ylc0ld(%^|Yc%=qU(!sZB z@Dv@~qrua3@FoqOu7kH}@Jt=NQd=+CI(TTkEY!h8EncXDd$j#xl@8vd!Pn~Gq4ly} z2lqUtj{7DZeC}KozF7wkt(Qt2JhWc6=-`Imsd%>O;7ye(yh;a8d0mBj^zi>u;q`j> zA5?gg9{#2ZZ`H&9sKVQH@TNbja771S`-BSb(81Fesqlaf?s->*8_ufl5Iv`B^bn(i zXKL^`9el0^Ptd^&HF%;99$GI+dbqY;QgraO8lE&Ae3J%G*TE|_xTu3~)8LspxJQF$ z>)=fqe69}Orom&h^Ll99Lh(2qyignW(7GY1I2L~-Ytp|Yjn5E#B>qBhQ7fm1x1Og` zs*VwYYxrMV1pM|e)L>8|;CuwUBLc38k}Ne40Z)#|pP==N+LI6g59uwU0lOmt-t?Io zA-x5a3H@u0fI|ci|7(kYYoaqtQzGD0N9bQi1YFg(swsg8IMp5gr}c|qp*f?pm9R#ctQmH`UrSp1bj#YJShTxLj*h}0v?)kN=}P_N9x(qBjAyGCouwkQw09Z z2zY!1JUaqDECN0^0zNzfzAyrQa|FCF0)9&bd{qSe8xioe5%5Um_xcEU$fiWaHbuZS zQx3znIRc)@AnKzs0)ATrd`krUTM_VW5%AFw@Tv%S=txG%o(Q;hgkz=ZBjDeOz|#}~ zPl|y5C-|Qc_|FLZX9WJ)2z<_evrBZIjS*eZhxQl@VuhzDs;5bG`eJHW(VmIF*ljTM zjClo5epI@F@-L(me|t|)&w@=XP8+Ykr8iCjgTJ~rP8+cQFTHWvc>S;T#%aU#zt9_} zjn@BMZ=5z$|MK2AZLI$Jy>Z%5{SWuXX(RPd?Tyn0>c6WuPJ)sDTfK4GNc}hV#%Tlf zNA<>eieC(d`%4?Czr8n38>qjfH%=R;zq&V08>as+y>Z$o{jc`MX@m5?&>N?X(f?d; zoHj)N^4>TJM*jJ|aoQOD5BJ7tL-bGWjnhWxzpFP+8=(JNy>Su<{Wtc;ze(|^-Z*WH z{)<<`{ZFKLdpO=R=4ZPMh6Vg6i-FNc&zPrnUK?Mki!ac{AJ@en(ZwIs#qZO_zpsl= z(#7x8#lNkKe^VD9u8ZHGix1SrgFCM6?`2*5GhO_YE`C%OKcI{6*2Qae@$I_!ySn(F zbn)No;=k6#f2NCX(8bs3;$^z{Q{lMivexH`cNoWza8ewt*s-dV?|`)Q3F9f@1|-;} z9yzd5S@D%TF`{%#`3PX;Q;;4%YArA_dL#a&{0>Ye)x#Lqf;;(9<&DezonkDNzIRoQoHCp*^xQZn(9ey z5ZzOrMT>iez1?UqT3Xcd`HtDJ`Gb{Dn4AF7aCVeDw zUs@|pmSw8-XGov1^{(u|T-cnooakthyyBvk7frSoyq%*&@99LY+S?VC+~|(ok0u$C zo0IFE$^c>geLXhw(c)UC%vsEz6uTHs%e1r@FzkR}p^z}ui zK{l7wAiB~On{#z^oPjUjj*KWSx{>xi?Sjf%$d7Wsv~M#o8PyERPXXq)7mYOZiRMRN zL6fhd3D@U(C7KS#V#N`@oJ{s>pGh+mIr!}>rcp2ci`p*LDQoaf^_+>p{8mwGMAuAH z3|3#fy8804Iug|3#$c=TO~n_YnEwZi2H2ovBYvh_oo8Be354D?VEbS~=^oa3Jz-8y z5zS3}c`1q^PfruSy~AIKi1LFj%8j||UG`^`d;r7qoAn{7A3&qRETS`E<_yt2^6ff+ zoG07FXV1>X_(x28xhWffRFAS2NVK9BQ;M=IGMc2!kBlZNGZ0Nh(=4V0r4C^X&s_B5 zecHP*JNNvYT5wG1MiSQ7M5-MVYj?difCf6F;v`>w5*-%St~RA%riW1&GiaKIQo`S< z14;^&1>P8j15L8KD=ES5Dl(;L>q&VNiL^eXYH_o;$V1CRYGw;WN#vTmfvDtqcdU#m z4awf*M(1fPkVkql%%_TF!Fi4=emm6QU;a#O@B?at%WL~+@c4)Z-vyya-{3W!YJ+Fd z8&i}N(cmO?zLhVL_HVTJlGfgL-X~!U1Kj;nwY|S&bw%$(7r!#VPHLkzz?l&0!tJ%Y z$J~wyQg#Q_2LGPkXfQ^2VY=PDEghq*@gaAqhvZTe{T^n<=B(osz%k-b3}SU4!M62OKMDdTmZQU74X(eFxqd<}ZqW7V6kItz*@! zP0?TaLLGZcLwb%ce}f?{1=7htT0HH{qrK~`N2y2qY_37tdK&TvWLDQxebKYn{nWKo zgtbJA^Or2M)eu^J&xCU~@f8y(0xJGJGEq;<=&erd*BB?yNr-Kl}**Nj# zFGF6xW~{!WRo?@IG*&N&?s8Kc7HzHiGSCiecfVsIuhlw>DK4F81WPkb?pjlBN(Gb! zhK2&Sg#zga*j-qH8PZjCzj+LG(JV)Tpj`*HdlsMfcRJ!utEss4K;? zbFnFoEB-Ua>wOZ)L8FUb1dVl}(cl2%ao!=~5M3&9h!r4?cAP|rHu&<@=(9#1N*a~G zXvg7A(@RqTn&Fg!FRmTX-)Q4;82I;AtN87%2s(KV_=UBg6TaM!t%A`CklEd>oI zV54gOfHx!?g4w0udt}hC@P8fmi(k>uOne5f&&4-U!y<&_uDIxNFHq-Jn=%(zq_<3O zv7s(EZNY+IONc?cHRyVJ6TvX3o$^TKrdAD_$e^!i&?JT-`bJjG&~j6{2G3-0Hkp8K z(x9v8?T(K9JE63gG7)b3n;OopD>uD{WDqZwXpw0y;*|*R+G)SnF1=&ghA==Kc>Ut# zml=q3yGr76lt*I|FbvdD%TK@ssE^cc`NDY!Wvflu)B$QLVWMIpP^I8a|0tZMg+ZGZ z3IZf*AUU|^c~cA`a+lUWl0+9JV{7-j}U^8jj9P=AF5$0}S@af7jV59cf=#61bvzcYS;u-KAYwEc(Ly4ko z&t9d*@#Qf8e!q2+E)=6Q_W!1lW*4 zxmQD#&@ZY>sQ+42*Q=;(E+gtpS2A308Hvy`D!{fsr(sfOPNE>OZmQ>g^n^xg6!nN4 z7tn?io74{fjnZ7vmmmtNln^06{Ep#M1&DzvLRElJzPO-b+(mEjw-m*pg_4vUEtIHC z*Fp)3m4#GEC@fd;7357Sy7OBfQ$=VYcTNp95akKvj@V#C%Ebh>dcItQ{1Dz`?olJk zq-zV@99f_N0)Du_2Z+2sULAqkFarKgc49lPTERNusuAVNdFp&`9S3QELK5;$uq@*H zS>gyAajjLmS>(L(A}i55z$+q4%u+L{{rCYQkjAO=8&JqUv^QQx@pmI@l#E*)Xv zoc_*d5ark4(R;KC%kNUhwdZ3T)Z8O$h`Hk+&c7RN0M9S?;N0r;;mqFp*qn$Hx9A>s z`~&24k9>jc2IrMWQ3;O5EuwT3hl?VDx~Ggm@01j@R)h%DiX4663h5&w&@+AHcG_6{ z%g=@MRaw^@HRz0iO{dkhp%U3D8u$%CnqbGe*ANw@D-|Ao!&yKqt|yG5>v8nJVp39D z?5@YNMfZdXWWwD3XbWCe>9`ea3h5n!+hXbwy#4`#^V0wTCvD#2(M&Ft4y5*2`K_01 zj#|nYu)B>;oTK>-k5a8@NWF+{pj8S;Gfi#$*865USDQ`_lWQxC*uTFu>jIT4yg>a<8*R;gV`ZWP?hdz`_jr|za!pmAew5FH&Nx7+vIRdHkAv`Edq z>W8X~Z#$27bB%(teSjReu^z}ga-g=Ju+|qnBDg0EcgMzwjt0?j8T;~tw|&*N_hMb{ zAK-w0ymXR#LX3sG;H-{9GBWwRPsdGu+t)52i(uYYHe7JW4(FPv_-2%z?DAD3$w?bc z2lnCEDheqgL_X16SM-E7TH%r6e{>uyh{h(hLGGBE4$jH@D^xXA$gwkFGTjPDD!>B z;+tGkjb_1EY;z4XbHyg6jS}5MzUzVnK}u%q8zzZfQ@)Fra_lq++dU+Xl1aym=b$SBoGOl!HDF^Gm72io+hyg?jr zUyOLi{V~4dUAKGWuAJyUL8H;D&(W{X|HQ)zN;ZH0*W;n`%c)8)zjSY?WEAj7<~6KS z`kZNKoM-1+T)$rr=45xhn~j(96t)yd_IJeeI>UjC_lUS$0zM6+QBW81b*v5yEGdr0x;49Kt$*UE#SB&o8%GQlW}TCMc8E^j)_up zD)!{lJ|S(E$BKjAhOLlzEUpJ;iYOSTf{Eq73Cc$&XPeS2{MH8hfXA>KG}|0r;^gUq z+jzHCYNkrFUk2|VC0yyYa?L_>r?}ZCJRKLmj-K{)M#0(5@f*f6p$g?P9AHwusig_* z(O~1ecBwUk^RxZc=d~U&4)uA<3QR*qr?QrqJ9?Ya1m?4S^sX644*(c=SsRsN(AC4X3i|idfq4d!s^f)PBebnP~j7(_b zzGZ2!p^G9yZ_D}=P`D*4Qe8Qo*8e%HSnyWSQ95u=3 z7?$AY2-yXW5`Es%M=w)HP_Bb8%keo26Sc~;e)`&nZBAr`Ph=!qV1Stcerq*krSYOm#bkHK-s78J znh=!hYd-@!b&oLE9j#D=yn}*1VRh0mRJjdd2vqNwHd7#w{FbQ_0ZqO{?b`jvP<=O6ti|VDJbNz!M||GA*+(dY2TOPM4&VG?lU8|wNa>; zm;~}%(-_9wsiA5!5D2#z+y(8FyOY3JM^4(D$VOqv<}0ji^kD7XgL>XEt!4FOP1pCdZ9s{`r7)bu5jGO-oSVI@aSswB*;>Nb!JnDnGgtVt)9#O?)x zVUsqeQ`IN(V?XMUPv)p2>1)s5{E{}5#hcaM$mK|+;llttuk~YG)hld1eC<;-maTz%PA{(pzvQ0AVilc^BerQSd)!Pj1r@CsW%6nc&6VdbuZ_`9UTvj0eQ6QpX<9lPGy-XZ&!wHN^# z_5on>syLw#B~m)%&K0dO6!OT;D~?iV{AGI7hZ4r0lLL=25KeH9QKUg`ntdK`#AnyZ z9kchy^|L`Y^ERKQhmRgR>2Zm13^vLEP-`NimXefe*}vj&4AP$0hPx!CMZTc60?P+= zpJq8I)IilxpoP^%byFLtPPGkyYXmAQvk%IFxMeBwswA9d zZ)ro3TD1qe{PQXaDcSoGX=J;+dWOD+oZeyN#JDOkd`?}D&J#Re4t+C3WtFVJZC}#J zB2J|gKweWhdRN?A42dXJi&qYV(we5bjo;JwyTHF2Z;(B+dZYuM(?|P@Ut9|XuUN(9~r-zeUvK7HBuX&GgcxoV#>{`4n+@}y= zY3;1Can%6Q9%-)QpMR6t{jh~mS~^-8?%@O`t%znwDSSEMkQ)8b$c;SeqLS-pT~2Mm z)<+A8{xP!fAmINm5?i{4s`>n<==2_rEdsZwGj5;)Bo5iw`5s0(oc`_HZowSjE7B1C zs?-C8B9o62AWD$1G?%toQaFT)aVU1Cjt~1MXlCoI8Lo8-YD2>xOvkuwpq|7m=L_WKtbS8MlH&ZW* zn{DoRgUG4<;I~JiCD-A2fnIi6l3S743x5-HSMf>t^z9f2;J{JdD|e0cKnHR*xr1y?1<&)*9` z$lXK8g@Qo!CtQ24fv-49rLyeqf=pqi-Tg!ol#0@qq}fpA0Mcf*9qQ>>dyk>$RzFR0 z2!4vJq(42(l8RS}t|t@3igQJu!oMgEM>x{YO-m46(-K9qEb^9WL=%;Dz$Cg#4Lvhr z)|SQ?igQTKI1fOorUY;2lW`p3{X6gyEK98#Y#v=8wN9!h>`1W|jbY}96vBFK8rDgv zK^cUEGq0q2YE=1}Bxymql*N>Wtefu1G6{y%UFA_!iP<6_l2u8{Y{Ka-%}l+lEJKhC zsTL_eUXrOS7RQP1XK-40GEo^p{X^9#s8;Pieel}$ zKACW4S8w}C{{+_{(Orlymtt59L*CQ+wq2XSP)DnbStcja%?3jk&S;IhV}FZR?Uz3qXy_~W%0 z56l==fHDWU?wJ$69A@M>pu;1@z)9nJ{;@fJ4r!Z2*F8$$fQqt%-fF(0R*Y1fC|MKM z!oXj2{LDdZy=(eKof{O4sc*v#NF?eunw0O&LKuuI!sR9!^(NFPe@7LY7W6hZq%Zs` zK1T$;|AJDcZ+uI>Holu8@Lhqrv2T1Mzc#*_U-$NR<)i)K+qM7e_IE`DzVFQL4`1om z#y2_w-@!Tk;T!k0@$JP?HZ;FK%Iy!|$NRo+f1MHd4m$e7=l(w>i>u3O&mMk;Wt8VQo`N zhjD18-BDo76T{w#)_T_<$tzhu5Uap(>E9^qsXOhC3l^!k%5Lu9pF5Ay_ASJnqMsSZnj6|(Nf^4}&p8l|k)L}`yGWo%J4uspAcu8b`^}F?#8Z<5 zsma-cGsWL=rV!nQTn5yV$rn;D+nk->A-RO+qGe7R_{&-?K0#=nZLza`K)pev_ZZk<^# zw>{DzcV&BhCrSo;2FOp3^}yq&xRT%QWGKxZh&i4!Z`q}TN_{g40X>Qb?`zzaVLu|J zt~?Zm~$byUT_z}V27!1q0R(WZYQijSZk4L zou_baE`h3|v)h=)T@ajYJgjYNjCb?pgF%y`qZ8yXA#H*4GmyiCyWPt;^F@Bcc~UAm zE+c6&DCInoCg0^=7DY+>`b-)?N$>TUWTd2D_n8z;NzOi#22#?(K9dGfQbwOiF_d&y zpGntI(#?G)4W^`S93#HkQZM_-KRXqK;ZLRGOkkydruKdW3Cfqd(Ny=4lUQ``CxN<{H#{b&o0>=m&5!ltto<^wWB4jj0k?VOXX)5 zqzv%0E231qMcK~sfS(mJe%4TSpT#X;P0av9YY?0PkoB94qdm=u$pN_UIs;&5q;v?$ zqcn_nYr}ZAY7`EFpSg&iy&vLdGen${Dt-h^{Pq@NWh(C?e&!;6MqF&rlOgVv0q&LI zdx$Z)2!^F`uwJfpff^slwcZYKt+j*`$23@=!LY7d$!Zum6z@$^z0do@+l?C;*K#?D z1F6g_PGep;Y}MoJ=pr9BI&eEhTGwA?P&kJ5!Q^gGPrAE{e@={Q)Gn1#-Hf-6QPKAX z`(aendMIcy((5J?quL{|jUn~2_k1)N;_5~3rz)!=zXR~8yBMFM-2{B<3i#A+=czV9 z>M-x-%cFsT<}&$$;4bN5d(>&_P+%ove_&K~7HN<3Qsf%KE2u!gYO=03-t8`nqQom~y?-@v z03{yoC(%fWfA1$Tni4DfNgPOtDnI*bor5TGK|dv9DACqW;&qfbwx7hol*sp!cs(Uv zWZdSf4T`12{rw~kp~NaR(OG+gZg|ZX_zk~giL^(=-xZQ3{ObTr?eQ<7*Ja#_@!z5` zt|{7^ZZRl#H!=qHB-Y#e3qfgMU`z1^1A7iHs*~_s6($69v~mi))REtSwta*HbYzmly-B%tr+1umM1p7KLJM zq4Fteu@A(0u z;|O`Hcuz#x-6JY!Wzy;c%Yg6EPbq4Vn&?_b2RV;=ptB*vUK`n?jL#7~01p&;Hz5fZ z4!CMN$FkOExMP1p2}*|Bv`*wMYVu)*d%{gtvR?fTxS1D?2U>J8zNqJaSUi$OfHCD) z#YXw6^Q(*lp%28euPgopCW>1cIk*_X??&fpL`l_@R+EXdTJDI4{bJQk zoXzQt5xK(|=AC@SQ;=b-=JRX51HB#ef*?HJH18Uc(Y z29b5()6n2xR%uKrCk+D3WeNt?`6-9-a8Fn#xX4cMeT=tEfOuzB3`+XityNKivn!_T z`xQM!w;)GZVioBzC}C^fh2R+aHN?X~BEJu67o0NO>nhPV1h)vzPX_T7FCwpt zdBEAEi#UdvTloiDMaK>d*e+zmYJG&q4!&8G4v5dT8^{ea`Jf1SQ#5zuhTF4e$xT>v zK5q&@bzlIn_{_7XdJqQi*~f7Yis?t~uF^`Iv@=5ro{3}f6wVV>g6lcbLB5F91&!HHscwKb+qk-3{$+E0pQ@>iQ684*GQ^eQbJy&|#SiFA?%FyOo%RUreNZe4jRNMYw%Kj=L2s&ibsU>r}<~;=mGrDc{W{i{U!h` zcGnXzqHE(81TC)Vn1k7-MA4O-ZMkQ9CSP%$$l!bCzg8lduNXzF%Q0)F`R6p8d_@OR zF_#&-{NhBHv;%rl2zJcg$~SZWK0TNE?)5BE)%>Vi5;d05c6!__UP0W#73Z< zaEBy69g=$~LmKQXHp1$j44c*oZc2gn!QnVev5E`tN)Pz!g_=}{dZTOcFC z|B;*=UlPpTvYQ2W{6Marif?AVkuNddNag^UXyF>GRE)a6o8AUTZ#`%gS z5DPeZ0&^Fza=S%tHGDkara03S&sUJ%SL2{M{+Zgs93?%>c^f3QHV-6*?HRN+F=iek zu}yP9V(Hb40_RaO&Ki{L2OJbwLCLqN$x|tLv6|edCYMrj6q42b8PdN$58N!oufpdM z5e^@Hs8~R)*uJ66bg35?MJcPEJhx8S96uvp6PxOJ0<*?X8Ml$gu{qr5~T-t`y zd(KHg>N4-)-6v7A)%9SeR9Yz&zhrZ*qXj0ob9iaaThrVlu9N0$VN`vQ@<}^zB+7sX z9r7^e;KN|Vn{DtlCas664{%xELIL;mQ8+7A3B1kg4H&_*s^P3=+vDww=ISN5nZfA{ z)HGrpGHv5Jn>!JS#|C$w(ZU@9waoF5zb$+$j?iugK$g$uo)Alr76>aOG+;F`2e@CL zEi^xp7w&7UP>S&e`Qu~t^3@|g`HN*2;1G7ki;{^U*9rtZh>z~Xjt5sW3TjSX6eK?d z^28X8JWDozi5Nj(d0w)42R%MVj3w2`-9sm1FCA)sE)mb7=hE@CJU5qCH3?7~e@Gi% zchbX&aL>p|$BJOFxfB!UBHS7Ny+ zh7WssNFMH!pW9t)B3H;vQ(9_^;M!P8s9e9HWs>3gFLKZ#gW+`pK*h6+uJ3_I<0f4% zAyQEny8LrY8BHj#_n`quAtWIi7Z2+3F$c2O&1d~*{7hG9J zsU8;DvG8k~ZJGd2FkiJwzz4Yc=74*`AS5Pnjo`H)VPcvvYyO<%0UK&1#wM>$iUK)2 zjjL6YQxRAWib(q5Ts)dnTR{IDqw^-X5;bEDfJRe{gXB5K+KAG9-jT~L8R1lT#wV8) zz*&VkT}~}cJc)duccG-<61h2^u==sfa_sESDR5FAG5Z9O zx!et>A>=tn&~tM9>@&XVDOCr2^OtP?R9>BY8h$mKYf*mle%6?Q{|AlP?VDfB8siT) zW|w&{zi~gv9Bk;JQP_V>ysg>W6^$7NKk{&gz>%a{nmzno<#o;uNP4d|8keiJKPPFX001g(KVog}`4O=3ZM#WCuRS1US? z|D)aRxGGEUWC}8p2US754Kn#QxHgA(&U(@8>ubmK>37TzfkSYUkM7*L9!54YjaylbC4mOIupA|J#Vq>QoT02FucD^GptT8W){aT z^~i-iwk|xow{@zs`yK`r2i@X#B4#$Mf9O3=MOD%KUi#PO_v7CArSp{X$G@`q{Wc89 z?{8P<*YXIP-)8_)9#oSDQS#r_&e1tIyJJSdwM?#tWMPQH5 zn;!_s=ZYG9HCB5)eDg|;BpZ(Il?_LGD781w=0dcT-vx0`;ND8-RFvSnL9Db(VSL7Iwb~+cccs!91WTGaoA&Ep_?FVs?dV z?t6X7S6spDYv*ec+%mX;#k|iLSNdYe)NfuwGvz_6EUTN>*ora!l&a_rb8gY zto|V9uQaO}u9cXjf>jWP+5mj~uWbO@TtBbJ3(j{NX)?&M98>!4&(9%ZuI1hD8ENwW zN{da#{dJH@6Z4B+Z#Ny3JK^8_I4Pi(8HRj<0A#i3KjMvQaWO7#%AW~VQgD3zOw)nz zdcoe-LG&{lm~ngDbPz09K6L&hP%9KafB<~SQ0pvXMT>jtD5;sZRpWakUjMD`34^_X zXy>U>$-A9hqxcOezya-9x8RzdDTB&zR&m`wlc@CDPPAJ(#BV(W+S!S_F^%}Efg6~> z)n{-QQX3%HW^fAX@d-=v;$hfak{5d&WS|ivVlZit4n`Q^d*;RRC=N9GggJq|jGCs6 zv+s>Vs<9eWH@f%E3#jRjjGAgtd#gFHW(%mH4yN1TUXH^7q-1zq6t{=P2dG5Fby5xu z_LV&D5j+$h0a_&VgH+T>bS=~UAbL1`Cckwry48xc;zPG~;QYoR_gUt|n+!=>bl|k! zg?mcT8SeNS?3^FuSP*}=&C!|RXcAH9Z=WtO-t3!SlGqaReNBY(c1dEZ>W1r`w;yVS zl0=a7e)69^LXNtl;c#9u8U=!TDWGPRj|ZtBy5AY3oUJhkw8+O-9ES66a4!OY91pg_ zOjFk4494;0C%`w}ivvA z%+592xxIF-0qrk~zZHXWi9*By)Sk2?N{7&zL@ykCn-eM2;yfNe-TUDxT9Vkn?QxzM zBwyIyNKI}+Uz-v`-C-Rda5Es3#LBDv0@$$2=UsMYCpzb)u2l2ox1t|t!C7xxv^zG< z;x-QE4xx=;Ww@1ZA7Z#47NCKc9j|*jXXUD4qj1y0uwl)LRtIAIzY>ck7=k_`_F1*S<9aMCzVI5*({vbSvN4UN31Vux2Tzv?MmrLu8#YX zJjH9kmXU});2A$0$0*zivT#uK%YV5YDoRo)LnsAH$f=d9W;V&5N1$TL#tj1I*Z{j> zatHD#a6qw>#M;eoJ%HCq3)c?N^jgwLL)F_tg4cTEaKZH>Vn<*T@bzrtH|~X6)QRnj z183hZdo@!Cwj64`Uh)U>mrUyNoG-RC!m_*=yK-nQL<0pD4Iq zsHECyaYDmLpVELT+Wk7-^no!nigrVv5I5@__~4?4!nN0=}Ic@bvzyS7xuA`RY{uP1zqk_$cWGpI)6^hw7Q-)3hswE!EK8Y z+{*?C5)27V776=mJ--#kh5JYGTffARfkSvfk#_Dh4n9~;&>q;i&)8DB0`>>R3QkIO z;HiB1CGZWOz@6{0yQe6nVBB3pY4;FH3qGaCF?zt~56aBCL>pOHa6Dy>ATY8z16)}g zpTG-zLJXfUm`}KlPq+bR)B>IuDc#Tu?4$kp{x;gHg*N$%yp`C4iu1DDEq7v&7z%LbcH)pV2i`=m=959g+VI4)#Uf)pLZT&orPyX)WMIb7biI+!lngS z7GNJ?zp0mBxV@gz3HXKD@H7?s$md-YY@li>p%iLZOOF;)D1h5P^K%ZN07kn6Oci+EK+%xW0(CWk+3rU)}1s@yL)= zh{O>dR9#~O;fztpZb=wGw+O=7qx9K3D7&vcr#H*7EFDG~P>-*DQE!$H^jVsUhm-R= zFVb!S>-we+tqXzB;6&az8Uys0XO27&daDE}fURv0 zMAvO#nsrbozBS!l^g2Y5Q`6mxUnh3B7v4JULy38g81D0q44eQo7NK!@#yEn;4)qBm1AdZUaW}iVbi@Sq(GB@r z;1K@y%YOvR{9{pzd^OO3KN;Z${CRvebL>Y_d^N-DkaiH>0X{ax2sS$ftQSzP!PlN) zBnACw+?dWVg3^f}gH5u-6w{`IS28wBPbk3fWS786S_HGE2P^N9KL)=gwml~}+temc z33}yS!499072NF;rhrH83SN;**C35{nhW#@V3>mDgZ#z-ND?fhaBmhpcd+oGEV-*} zwW(G9ZLpf9y+LZvBOD2$6mKT3b_(o@No<8XsA#Yz9DTzlloTGJlb%=pGx5@r!b8MI zcMuj*3y=2Z&4TRn zt!jgg09s!VUPtW;@b;2%GIVdlNR>k@g1Atd%xfD{Z#p7(cpp zvjtCr_(=wAJpT2~b_$q~9PGqaOZQvZW;=(@3D#iH=gkPh1Xx<&$0$*1X#qChf}`-{ z-}t$<*)}kStZ>!B+SqYb_VT~bM51KVuz9g*gsAG@8p0ltBlLaF#v2-_4NLSkp( z!K^Smv}J1f*dA98%NdCbUdp|Ucv4|K)m-M6&Dfg})0Q8_ylizId;>`yZK>7~zhtaC zK%49Vsym<}5lYHDC;~-W7zxE17?J~TASs|BsSP8+2Ua@Ex&X#ibupSW(p`~C3ODEHmn2@0uiFnswBiscb^hU*7M*215C!mXHue4=?r&lo-7Onz6ix3@q+jiwG}|c>ZSa+TTddEkh=deJ$6x}aPV6MB^w_5eBNZb=Z_Y^ z3Jvw4Pe=~Juv{|wg4zY)I#q$lnua3D;r51dv-WCwN)%=dq6&h|@}E(AI0U1(3PXpf zn@Bwbm>R;+LHvWLx1tp=H3Y4ucOLrv=M_f)L&*7MT+!3ar_zoa>d>CrxE(gG(Z<1Z zb~l&?K22jI735`RFXXq!Mv=k`+LwAvb~OabrS(>5{Xok=hJ6hu-ldJmK_009qE^WI z8eE+U%D|d!u7UlbIe3NN4d!|4qk7>7s?}f@TmYKBhIS zi&}5F2}kJgR%Jo`8;fAvgJhEk)#2nQ49{fPCAX!xx112RsVpOdmH`Dfs zos$%kJ-Gt8Ks=IE$ow7|I|)$1dPXmMolJ%8QcMBK96rf z6%6BkNbX&rUD6ZQvBtqDaxsIeC%XtNA``GeU>6z8wz{xggx&vw2)0}PByfdUMNVYc z_xwW%Oh$DItRi6L{}ZbSw+G%l>0}hyNJf#>YKLGH0f(h7FayY!NYWcXu-a4uNB~^S zXeJ(Yn(;0Qg}*%lYsdwafnmil!w2KBUtxyCi9@UrIl)%wdm{Lt2_BtbJlNidks7g! z(9T%a$P6C`ad_!9eBgS$ki66Co{g_^GP}p`Fqa{_hvP7e9L(-P=8k4Ea>U=7;b^6R z@fM48oSJi$;t9wKY8Bs#@lM!XaEo%L+b{sbEF8kp> z&IMx!waJ8amIn67zA2)z z+m{8EjsBM2k!;kH93V$B#tw;^LB@C`fqwxs9oKGG!*`YpA9XN%+}z8-(ijJe!b!q! z2G4;>BORlHUp%5yITm-SX43dMahcQ}Kl^Hl<(mMhNoQq z8BQhr`wNJ!`x7j%e&9Uf+FXgM-45^n#D8dcTaO+bcL_;zNkyD)-bLdooU% zc2kpuzHrpmjSC(4uxfz5d~0)Cp2OD+U<;b&Hr{J@bl_`lc1Iid9&MFyriV2eB`=Zf z2@1Z|rZ&(5)DpP2Lf>AYlPt~>$kjn$Tx03LZI~aCT!7Eo+zcPS;qdK?$6%CqW4s~s zoNq~G=afd0x~m~|hrXWx+j8DygZy!x>cS8o1?vJ?B5kbGij2Sk-1jgZ2`9p$aeyTb^a&Y(FDS94H*ttvHpp}G#ttH12}(n|qLB+( z&-#QpMM2VqT-20$IC!+kLxLB5(}vkkc2LTf^iX`&Sw{MtFPTd4gm?oeUAV#b73s9X zw0}%pcutNEekl(FOB@LcLUa&^!Xc#h8-|mFFf!PVp#Z29n<*Z<=&{dE-@Q7i9Yc<> zW5_{v4EZceel}PGmRfj>9YA^@tgG>(Ed45}`KkHoj*;Ypo~9O(@|FcIl>QDkr0q!g0c;9C{voXL}2?#5E`bE!mA7$ zzBTzOEG3<+39OcL3|lwDMpdXlh7Cp?s5{QoOLAWZwvj1AN-5<6suVzBqJ72kbQ2 zk6xniV|YHC{lfr-pQLB@RCtg)J#<>P9g07EplBpm3S^K=AbsspP(>qgoxXH|jGt)w zI|Vw1R?_47OQ0KxP=NGfUx4moI=#V)Rf_#{0?|g z6R_jE)f2uCVQwc5cev2@4r3&${Ac4X+TenMyNaY=VgAThz$8fF;eG_ogbh|EYr{vh z7_QLcJQfL9kW_~{ps?5)#HPS>Wsnn&Y5EKK^}MmrSBxQj1vp0RAhHNF18N*ZP*ccF z3f6;GCePt=N?3LAMW6`+-7fics4ZTHy5eJEGl5SK_TcZNuO^?IsGy`EJHa?+7qF6D zfT=F>?XBry>j1udn?XvA($Q`7xIm9H8Z!ZR2wMmcR;{eeM8G%&@PTo_P5^sxunP){ z3wWd3uo|2pg~bR6Bv(S$K?8eBc;lgqKOJ{Sl7( zqyh!kK)v7^s(Msa$WvZY8jy<@g!Mvb7di_;(G(+SXJHWdU^;D3BMb-l1gw!-DMAfZ zd}m=d;Rz{Uj}jiR2&zlfw~_c~wGhDFKl0CK8$?Ov0f-BmE;7*1Qr3e@7 zw1mAqhpe}7Qln(T)yZ&$oz-AqMbbZYSUa5a39^BafH}a1iKCn>pU#edn^t@LU*-1f z*>B@i3>Wvm8s)>;CI&aX1M$V8!Y%I@ajP^+Fkj|3rV&%XYQR3(gxzr`4t({vO$4g} zcMKjGt^CG6^~WPNY5bur%paZtf54I}d*ABHfOdH{TW%SyzpLx4(e5_A%@$MY5iGMN zEGeufI-kK#z>ZQLT90(jVJcNrrG3J?q5~~Q31Y5hR zK89zxaCxe66kky~5NJq;TyHes=i*kIZq?jzRSp|SI)u}@=^NU33-0?!d7MIWpo{`^ z+hYyh;JXv}?Xyj{!f=4O{~T)K*`{xhJvI((b42)jYsGzG`W*w(FBhU5{6g{l+ByNB zz~SgHjAMUL<(XjZ--TN=cxRr;B-MgANQ{Hh7+e&&z$U;n(iuMm^8}p#EpS=^=M3n%Y*;iyVy_`qmCV3p;l=_EI|TwvM$irwupA6P{KN4O z*4SFo5;L(3-Zg8GN~6{RKHC%Kvm|P?FlnZd#Ee3+$HJW@QT~G}fwe%WQTZ=s13apl zq!)Zt7Fo#%#txN*jSPXpgcvAEG$0d);7X##p~WU1^=DNaI~3A9lOqdBUz%)3#xj^} znZyIFv1S8im|&wt9ETNr=eh)kb;9#0iv3q#zI>RSh|{np)EX*T&fr-cNU)4wirrk&pE25pj9v29+0bZ zf|`GePl%ay87y@I`MBW6h_+i2Ir+!8<2KGi@6qQ)Up^~$%#5Bq9yA$vB-%pni`n(a zf#bol!3_iOVMJ`dCGWxYN_g2PD>H}0Q66Xs6Y#bq2f^=2=Go;tv1o7$Wp4z+VQ20_ zCgqV~pVPY~yU~&i|5XhivdRSg0=jY*RZ6}M_R_0E!=L48H&^CJ0LU(>+KVvv}2C@ zuomNo>RF^3atk37_HMLTa6JqbcA2pmID$^L$z7J6_qP}nbllsHok7=^!Y<1Y%E1D*Uw?8P>Cgtoq&2+*a&5zL7qfEW}FS$Q64jXz?(P@fci42eE) z0Vj=5;gx|S$JHX4JTqV=ZC_&C>pL;WUS|M?(cur07aQv-&UQ0=7%g2TIR&#qu9K|(q!WTjtjL4Jfj73!bWb8!V7n zK{O0^YLR?92oDCpF~MWVe>=zk(V5E zxHS2xdhxQGepa-CeZ<6Ve3*VxoGxB^+wdWdI`P?0)k~aRqIq9mmoN7fpBbZ~@!^}g ze1#uP#GOmrw&XX^JxsbkC%COWlC>&B`t*#Q_z?%eYbmH_*|*lwM@{fi;}r?GfQc)5 zIKXutm}bH?%Y@JA5@Zd0e|#~@9naBE`ONviYW}$B2ims~EYcyV!&O>kap!O^TckN3 zu-lrrJtoaaG;RwEYhkt^Gon)T&kR`eZ++6VIuDKv3T#j&D-(SwhZ&9`g?u& zv9)V{*ANJeKp2Zt8aAk(j>ZO6u5M6&{U+O>hL6Gq*XswQZrqbd>$z_M>D7>L_57VcKwii(bW#lw&US=#g!2IA@9WS_3OU;J-b*(ZOaB$R83Kf$n_ahoFg{fYW6cOgH+Ovt zgpmbgNd6NgpHh?Os>w(2QYIl;S%ybn{$wB)_AMfOvb_&tVJwAy+Y;0j!BrRqfjB~x zam{t^e~&nh|L+jTnF!+Wew8>LRQc@Jp^vrn>Ff|+*^AMOq>oCV{`cu4{-4uF%QJn_ zhq-UQGFYXLtV9sNLP)KUMRUKw=wmxx%0wVj3i0SepZ=fu?`L1xhOblg&OZk4|Ly$4 z64J-M@{i|V#xC(b{TLk0e}X#ycqtQsOew^p@A-ew*x%yE20`+~|L=?cBP^%?g`I(% z3PL->2MFtS2HU2uv+k4V^P{14e*}Z}zc=r<{PT5xc;&w`@4UM1ts^n-BVmNXygxCV z&3iju$`l||R^ic?yrBIe8rJ`R=08%NRJKRXKL+pr?fe)28}tA7TQUDR+Wg0>^N*J@ z1;~_Dc>JUJPr|QZ_ObJ;=Zjjb{%g+{KP7(P*he3`$B{z)f>~|o11@#{%10H?_qFJ2 z!yi?=_%+WFoj?=*>e&f+1wy(#gTJ^f?q63;SK#y;U(m_$enZMv( z{M~@lE%X}u^O|rR9?=HX3Df8xd`!l#yg-dm5CaIbIWM+ff_ zxMycggr1@WiirJq_dv;Tm24XyB~q!qboq8~A&|huoy^97!DEip6Nt7@0n!j1r&3U7 zeB!Gqwl@~iq5DhqAdhcc**)n5fHr|S&*PGs{7Dc$d`CVk-u&P#l#%_G-Dp1n*$>*m z3XLEsch=k@pTiBRbHSkmX@H6V+HgAK)yMo9~PFqUIPgt(8o2gK@cSGN-J4*F_A2-Enx|7kN}y`T&-3sY4K`z*&nj( z&>)_Pi?IzN%}YW;s7Oge3VA3gZOx}NZwVI93FcLR+x&ngda3C%KP!lN_$7wz_soy0 zyQ96b!A|?f^L(^AbI+VPGjrz5nV)wq=@<34T%KHb*iFX(pQe+yeO9)6V{uD{t1LQsrP2;9Yi^b3OMu{B0LXD{Dl*ebbJ5@Y-2cd_!aWs2g0N6yZ?d?9lrNFuz49DB*o{>#z5YI z&$vGIFHYJ3$GODymE@Uj*>&W;_kP#9`z!Po;;yH0F(8iqwTGT=duqCf?rFyL)*w3o zGJYhmXy=_^O}Ht3aWQdmAydgxlifXyS%9sSlot#vxH@n|NJhtLD$|gMW_l~f{BX-*B+n5!@hsm z|D&~KQ~TFdm1Xz;XsCD!q_l=#@e+vO8QS-3JI?QXaRgHEJ8((Ok2mfu-qV34D*nBd znP;%O{0_bcJ@c`j%HcEvewIFh4G_L_Uw2=$6d0;D=L}7_z%## z13PKJhV|Drd}oaw?SrRsCG35?<6J2#s%|Tz4~fxM-z4sO%HRp(gtgl+ z5XJLi`b6^Bu?D1R^vx_z#2ODyvRKx zm2{!f`Sh;L2na^b#&r+m+5T~Owi{Sr-@bWj^4Q8}aP#J8j;D(!fAG01X;GHG8iZGL z@m1>slP_CA{oay$R#LM(wspzhl)T%ZX?NiJKi`BX_!2e-pKso_^7R`;`@xlf z&Ba)}#^u#(r$F4g>zB9^7$+@{+;}gvdiDG1DtrsNNnW)2q<`C90vF)^+E*_F312r8%h>zh2c-@M7*xLPhLPS#>bSSA zV<~~P&n`gBrv_Q@3lz+b5bdq2=D`~Pj(iKjG&*YC_4KO;kO&_uLD})?s{Lh8Ly%|O z83~+`z!?dgk-!-V*h#>lVT_`UemZ{kjy3Q!(L=G*aiykxu&yJL8VaUEgL*KnM`C^9 z5nWrOYx=Uay4D|wMkBES9pTByh~A_k)@dCgWR3K#TcfRu1XD6fC>ai>!+kQeS+5KZ zk6fl#>6IND+b>(ARadWE3GPFgs1X@TREI)ZQ`J?Wcr-H<)7P%m*HvAWjBnMN6&xoZ zG(w3Ek%*3Q+Y*%=hqi!gEZsnC*Hzz=2quF=sp{?FWZX!^5kN6LED|y$Bf)F*Cu(c* z^Rhy3u%!kkA5nAuRaf;!Mtso%NXsQ-MTA|nbIBlg2YZhUb&GME(mS{2wrZQdOR9f#3Mxx=qH9BfWk7v?} zOj_@Yhf{hio~G&~4H+KMQBBeCP&k$jrX%rKHH)sX_=#XDrNn2^z``!#a6HmSyixya z)^sqqDG^Qv)A8h*HO;O5=8Z;2Q`fa^&Hm2HP&|fAf*7iJ`4#$nk2u%H-G|^$n&R6U zjtmT@KYUVLPxZ@5VP~~>wni`4XH|HmUZdsZ3q>-h0Z4TP&zUkHd3KrG(4Lf zqT*F~IWOY;e1g%$U~rz;nl>CsrZd5)p2sJ@9-6N2Ub9AKF&GMkqr85MbTSxACE}=Z zqd%BR$A`m7gO}Yt$-H!^l^Mx!zX9%4vR>E?o2aLfPpVgzLmS?tDW`?M5sMq~{{B?h zqCiTG&1s`(qLs1CkdYpQq3nYK)m(NNZ_Z4CNnV$m4Wv*lkRcJVjs8q5WYmn*7_~K- zSTY`sB1cT9j16hbDe|k7j80Kr4Xi%$vN})YoO=C;^h~jwl&(?*%&9N3srePwkXK)) z(4JULQcKE{i>yM^+~sN1>8}Top3zGN-?mbIf9b}_D*R`qdxBD*vYjZP_eL)>gYy^r zRPECV(lt(Jd+?ps_Fx)si*Xg*g|#E3fuAylW6g=8TZd+r@sHWGEyFc3zZUDOY0n)z zxc2M!{ruC7um$wP@(qIkaC&z1HtxHoo@dJI>&P3#_<6797!X zmH7ZqaY^ZdQN?KvXOTqN@yFnU@poKM^oKVgGq`I->*>U@F#My}<#PJ~#{j1Q$8X5x zmgDZUCjcJ@d;!pdZCkB7mwO0>(*Z}Yq}B*H2z(DcBVJ&0E_WPo3a|lm6Q9lH(twEp zEZzY>N|1zumAEXR8}JaI4tOqvGm2HYC<$pxE_Wwj0`MWgaX=lyoB;FzW&v9PrvPsP zoCZt)&H&y4sNs^suK~IN9|hC_vw$AJLx7EdM*stW#{qi)-KeAl;BvrGzy`oEz;?iK zz#hN}z%*bM@J_%f!21EG0Ve=w0QUoGL%H0~0NsFZ0P28L`5wR)z(&AszyRPrg4jAc z44BxS%T-{}I*X%UJ%AIpLjHg~I12a%;0$2pS&09IT&@dn8t}`2p4*Tv;3(iU;1u9- zKpiKAD;9z;pbsz$7yxwN4mkt%0Nx2W0r(K0hNH|`zyRROfNs3q(T)1o0i%FZyCFBg zX~0qNp8=cz)b^l!ivaI|o&dVBM%)89{-@9{zzM*I0jF`R#|&Wh>!?rcu8rQ0{1AKq z`~U;vs4u{YZzFxc>|f?`i?IEs{~(vU5^&;&$OkpM9LVJk0*?Lz$_qFHSc&r? zJ^0La7vL1&DB$SxkQ1Q$5abQm_)jPYAl}c*RbaI^@eAk^!Iz=O1b>PA5}Zao0*?JU zmn(z5PyY+}1G--UJ?a1dLHPi+*C2O-aM*ib&awSk%`xJ3T(D?C8SP~d{zCkwyK=cp zh@xcyIBVK+{CxNw*aVr7v2eN9yDF|;QoglpRQvQNR$pGb{9<5P{F?x?pMkCtLO9}^ z#%~|tDZX$-vj@M?P%a16Vh+pe+BkmgXy>0NOqp zZ7aV5+ z?j~2+#`0UGm&GGq`w?%XKcD~4fJ-a5H-JkhI1P0WRdDA5H>lt$f$LLnR|40g-~zzi zq~Q91>sD|hz;!9OJ-`JN+&FOU3hr^>S{2*@;93;i%fK}%xEbKCRB*+x^$iNH0yv+7 zy8<|mf@=is3I*2%T&02=1a7&4yA?QH!BPKqp@Mq|xC#r7^ppC(;$SX!1Ioj7rFn5_ zhebaxa*}?wkbaVmP=9tfoUbF)_e}#Qb&&Mfjs9&B^pXv4pvRjRx|otv*cAv%AZ!SJ z((!%~M*Le5wjW`iqA<8h{OB7sW8nV*(&2f$Cch2+?;>OoGDYy6pzXtc*#%^uneKzO zh04<@8r+Y|%mxRN7W|$>oInzN9`u^p_dR4A?k-yA()Sjxb5-ssX>xgXmNvN>&|G~E zmj`VgVHP8SXaW3+(J;GHx!hauYn#YUz@+UV8zbz*2A9YA6_|-;mu{Mg*79}*(^8H` z|MW8lpQEKQ^t|!~G2|RCgRRUqMB%3vcc+qq&x6_QrY!-{h zY*stsQOeeco64EI4FU(R%g#u(k`V-bZp(ghlCv}ABiFJTZXW6^~!gf)QPs#k-+^%bBeUUwkuYc^ry z2s>yK_9VhI+;^eofq2q*#%B|D6k&rlVQzFj_uGVBfv_o?uvUZ>Z!MgcK7=(C2qPUH zMOYMJk3*NWwnYoIc6X_BSi5>j(QtREdx>*+#O-npC)|sj!-MWc&f%UWcad{AQ1l5j z?}b_i+D}_KwVwmvlWx!D&WB&)k%^8mWAQi##*1+U*>S{~z}RgO#et*zxX~UT1&+SR z@?rAx8;8tKE6Yze_`JXIuXENJU~fy6~$O!S99 ze*pUDw4;B7(YKOqB>H05oQLnooRNBWfL2F&=~p1E{t1Dy9wKF&CiWRR9?T#KaHiowl|mi0og&;me=R^ z1>O1OJqr3tjD6p>qaT^2uLR5G7+W91sRw$MvY`C;IhY+MS<-m@{x9Wne?-QD5nS5D+|VbXs|Z|zrojQl_1nBoXk;|QJ&A&B2 zS26jt5r2~XIPuTsa=q|VKeET!K{g#0cSrFpt_t+RZEU1Tv7{I|1=k)xhBmG+bBPsG$)BHE-hNb8XtupMEJxDxm*>6 zf0n}g7MD8TH0$FH%HNA8CttF27or0nUoYk~cqT`3_W}R#Kjv~z*yZ~cW-qTM9j1Jd zb8HW^vl-@sZ+sC`&kN){vmIf}@vFexb|LIt@jZB$?FUf(yF5rZKSpla6oEDJh4(sIyxlhVN;wxX{zVU7jCj^x8nc2wkjWnQ# zQ{^)*!pB6sDcpZXZcNZm3+UIpd z{tuE5O*zTYeh4}F+^syvmFH3Ad9U((zw-Q-;&~BIhA&3IzgRpUe1dz<7tb?CxaR`# zTp=J{^k%M1w0k*EEfe^6Kgh+&J0a4Ob3WSdCPz!p!c{(vaB%=a;ok+VcRSzg?G!knc0U5=Lz z^~6vgi01=BusQK8^Eo4)Zx-SIPdslCd`~wRe4{W&kIwOp0tN)^5ilX(sDNVvjte*; zU{=5>0jC9=5l~yEG3ORg7tkYMqksVcdjw1fI4aJ9 z3+NG$ho1cZvKdaB5Ozw&+|b;-Mz6fCHxo-|bZ>2SO|_@0rZ&T#aiEK1*s70{MG^WV zdQ+*DlADAf?n;Gd(@A^T~`(Q*GDA%JP0<8yBE(}5lq zySR7(YkiKTjuozR%8M?#;G*l!Tk3Elsp)z*Tq`fXs=UYr#s+>RXbF@>hr{tZ*J1(=dSZx+oD@P9!Rffjxx}ec z2>Vy!(9Uwo+!wv5IZBF(7G7OaQ3_A#=5tE_Yl-7Vrvu!EMnb8imWibz1F>))UbskW zI55o5!{Q9tptj+f>-1*tDj{>dYA70y4OI2TGriGpRf7kJ)L<&T1u&h7h1U**W8q{Z zgvVeqG`M!8p}q<)1XK;kzfCoImHhRG`l}Mjc$yBE$75^zGtsDC)!&l%+ zN`6S3-ihFUsbpv+PkANIxb{Z|Rz_+Y>Q{!6X*vr_BB>5(D_JI2B70Td+N#XetG2}I z)!O2ix)PI zF8j0MG8a%9DXomOKvup~UCv7RPpVX`y^Xgx^Bvh~2SWA;L#vBmN}&)v4R4*;mF*GZ zTiMUuEb!YD{6jW;o)q-jKk#1U%*EOCsG% zSuCvS!dw)}|0)~&)xZ-UIc}Hz-)%PZdj-8J{BPLM|An9rh`JTo)E*Z2Q3d~Z0-sgz z&jU|#2nab)y^{N9hCfe};|Y$13{bCBE=|l!B;F0YL;JWU$0riM68J)LYqG&_vcZ4O z20vtj|6|}OpQEijCvS-i-DN}nO&dIomxxcIhcoctncS4XkB9g(+uMZKuD;#}y#u45 zD}Tl5t^W(jr@WZxKVgoiB>ggmKVMVF6PE+8p&e%rN=gpb0)H{`uZwYWi3t8v;LirW zr-vhCKXy#uJq;ZHDM7Cz;!@~IkI)lc;Ol@VKH23QA@zI>@RV+k;3L!hs=&*6qr?}X z3nV^r{>je3fMpBAi^MEoGw@5D7ihW|2gsm{F;FEw`U;MbY8nQ;E{QDgcE;bW$@yHf z2zv#1%D-EYn->GsWf%oFiuq%iXfW3Tf3|a}CdcbC2|8y=eB}H{miKRoUJ_a0^MYQ^ z|0I1SEIO5Ie1a3RJu!HBA<5zeoVnvVI0r+_!t|Fo%4b>2K+LF$n^_3kJ-cMIj!~o_XHo$ zBOJ^2GU2@t_3qX%D^m2~GKRlUYgF)cLQmX6Fl>Jxe0DIrHAU|KmY|=wpFj5slAi!i zau`$e{8MO9h3a=b!-M{tg0JAS)rS5a8~nEgA30BEXC=Vvs11Du=2fg+T*DKTdY)l; zkqHaX_$^m2Vi=I0sUbxCA_bl+W} zdVI$Q?{;%LDfP`Q?87C%lb%c{da@JvLj1pGgFj${|24zI?p@6@zf+Xq81R(txRUN? zP>F@ojWfJkYuv&auro67J}vR0{U!zff3@NBfuNV`TT;$9AW_PXT-TEJ?fU|sRq}Zr z2988Oqu7&r;7On5x|-zw0K<#SS-{@|Pj*|b`&EdrUkd)?3jd4HkrDq1A)M zP8VWhtR=P|*@eF5gvY@!7dGcy)BWgrEKb$FuWk@D4IO$65b>h44;vR=O37vv$D4bj)#6 zg1<)!4D%CmtFh2o$c}{>9{L&KbZ-kg{6ic1pW5JmEBNT5zn1;}GBh+Q@6WiH@2mR-Pr+Sh57ShgqQ{d%3 zjKn)H;`vM@S#laXmk4hF_y=KU+%p`H;U9Ci3wpVqff{1&u)xc8e|8oXAuT%R@7~1e z*tuVLcK}cMOep&EJfr8d*8d|b3d^BtSz-KtDJRu|Qi%LZUX6?%cDS2k*cn@Rw*s$A zA`841c$!DZ^+Xx;fZ)@(ieqjQ0Y7JWNi_@K0lrYZoO>~skK9L+?Q0wGL@)PMWP9A~ z6!iCV-n}A;ZwPviBF}e$e-PuG?zcIitjEeva=Rh-z0i`GTMs;|mj?bU_2gmTnLQEq zmc3&Hl1YZ=wATN>67+JPN6r_TF%M*Xg#4xc{2B19yh^$ysQ0DR09&(T3y+$)JlI7Tb}bS;awP!2S7g}B^t!!v6$`g*3G5P*;W{EV z)?eGy;Wzx(v>1kFhzi`aetoCE%jjxa*X{?lW%D&n9c|4Z#>R)?ZxwM`TduV6q_AbTBo7Z%hP6aX_!tlA6QE`vL7 z=pr9>VaSI{%i9nO(j`Z-IFw3327S@M@K9rOcX!jeHp5$8TfG|U($&!{HLI1^Yg+W(&C23nS5^dbj-nOpUsKKUi-Fyjq2f{-`#vDnKY2n+O_WhZK zZLJTZgLEVj-6pI#@07ju;Y7sn@ott3e2p=J24wzmliyiu@IM&caQGxMUQ?HU(U|2T zFo=d48!j9ZYz-qc5~L~!MsdBHT>?tIp}(&uHgnpnO|!Dx2xA>%6sjlFe9Y5sMXxaT zawtnWerDTE#Zz>uriN8*e>}N0nCvq`nZBSg7>xGM%~8Zsld9eFn&UEJFgNaeQB9?u zRpWvTx;pL3;MZjLIFw0;M=)TcTlzjk4W53(X!HfsK?D2a$u#U)Zz?7Iv$bl9Y)-SX ztyO;Imo&%2O0nlFbi?dGN}f-c9ZJpiRBFARE-Ol<$GwT6!7ce_$7fND8c!kXAaZFM z1GCIGxH5j@j6G?_GohPJeX(e+_6}p%k(Fb`w#AXDuwqdV#75q5BkuuWT z5@FcK{wW4LGx48>#II?+hps$aeF zWZJ3=(#qm2Pam7$m?gxiGE``1+84s*qIRu8>3Mv11E@kbnsLilNJuxvf3(vOiA;s=Z!P#LjqKe>9UC%nw9}w}vKMnC#t>31`BA;ZDQ4 zqM=$#pS?++qqPN!k|*A+L5l)$@yhm6Z}ylUy64Cn&Rg4Hji;I2=G&1sWAHPQ&Uh{E zi}i<8Dbu(ICcA)Mo~AjRsxZEj9)%jiHFFkLtQJS%Tvumn+U8VT-X)3!YSCiXlvG-3 zxIFFUR%gF-G&dg>+gd0amWG_fy{^HO+EKJgeJW6a)KQqhN-?~v=8PEUrcZ;#)vWE1 zBIk`jPks&0nu)ocnJP{tn=M|FwlOd=V&K-{RGfzO>21ca$7<1xFpblYj%8-1@L_%C z?9fxm2fHK>c48GRZ{^K)t*uLyFlG%OMdeuQohxh2xQ?C6G_DP&GDDcTG&f`HgEcI( z3gsoN0>Mb$tVXF^s;1M}BJGPYw?(lz7+W))Ta@{-Ok*dqFNMcx^BR%-0h+oJ zZ_y^T7HU2DwZK*dwnox};>P5;C9pD6WR3!S_+fco&q)i#!P%7!>N{yw&K#?V(U^6N z(17J&Hbn4K_wMyD^<F_F#1%LixmW*4C`qHkjYLir3qZ#xxxy^(n!<~!Sj<X`UlqtXbxv<)m z6xcQ>+?UT|N>4B@{Jr$(^Vkd4=N7u1-g>h?A8ncTuBNuoY>ax5VShsepUMcvLg9I; z4IRonL%zAAfO%%uc7jG`PRHtlS>%s$I>0IcrKhT`3H4%C%a}7`oOh#wFYd)tY0QO( zV325^q%8!jH#8W;267T}nH1I({MQ0k3+ojL8G^bE+nyf7-&t=2Hey3Tj6e7)B=)o{ zxk)u-8d7E`hUc6?LPQsR`|f$vy4-{V=017019a)|bCIrijMt%ij~h2!Q%0 z>qbZ-i7{M1E!UducK*dwD4aS}FEGX_xXeCJk?H#u^N&xIM4*)F5Jf5csamtV!yhNx zOuJ?WhDvtkn`fWC{c3ML9c>6E*>)^mC=90p8-tx51Ct-Kw_~e6v$cJKmGjfJFJLx* zwRa~AW>gSZ=N^v>IZ~vLdINx;YCyMnfurWt0M~GJstL+BY=%qw!$c$F>l${DQ$< zNGvYX(FZf%*s#&6Q`?5{Dk)$($st~TdeXo+DbT z!|_UXSAvY+P{cze^jL}Em*+kux*i@nb1vi4ecj|F6tC-XB>u~Dml8q8)2;F4`@9nB zl2AZ3{tdwYU)-r6)1MIOODOYSD&nj84+4)PzGnXAy%ZAC_btd#9!YjGegz@n=-4+U zDbtttTu3PIwIEt@s{C&Oj_!<*@#Q@c65hyS&EA_K=_LFD!s!kQ`7H0DkWeN_`KNr! z{0p|)UU;a8W_-L~MNZz2A@9Xd3ATgv;TDRI?t)~Psf+7@pYY(n@|=^aFOSeCH#&Ozwu&DD50bi(Q)#ZI|(Q7XiYyL z;!Eg0LHbW9@#XvJ65cG{hY`t|>C5N;p~RQ(g-R%@QDI{b z%Xkv~DMTim%)fk}QNEWJP||Od5e4*!^gRk*#=k_tOLh{j7V#A_ZN139%2&pd^duh@ zM`6PvepnI;sM2=X#D7V|Z$Cl&UYq!riuW`^Cx{;t@zrwPCExQ +#include +#include +#include +#include + +#include "matmul_mkl.h" +#include "../avx/matmul_avx.h" + +using namespace std; + +// Function to fill matrices with random int8 values +void fill_matrix_int8(int8_t *data, int rows, int cols) { + for (int i = 0; i < rows * cols; ++i) { + // data[i] = static_cast(rand() % 128); // Random values from -128 to 127 + // data[i] = static_cast(rand() % 256 - 128); // Random values from -128 to 127 + + // std::random_device rd; + // std::mt19937 gen(rd()); + // std::normal_distribution<> d(0, 6); // mean 0, standard deviation 64 + + for (int i = 0; i < rows * cols; ++i) { + // int value = std::round(d(gen)); + int value = rand() % 7 - 3; + data[i] = static_cast(std::max(-128, std::min(127, value))); + } + } +} + +// Function to compare two int32 matrices +bool compare_matrices(const int8_t *mat1, const int8_t *mat2, int rows, int cols) { + for (int i = 0; i < rows * cols; ++i) { + if (mat1[i] != mat2[i]) { + std::cout << "Mismatch at index " << i << ": " << mat1[i] << " != " << mat2[i] << std::endl; + return false; + } + } + return true; +} + +void test_mat_mul_int8() { + // Matrix dimensions + int M = 64; // Rows in A and C + int N = 64; // Columns in B and C + int K = 64; // Columns in A and rows in B + + // Allocate matrices + int8_t *A_data = new int8_t[M * K]; + int8_t *B_data = new int8_t[N * K]; + int8_t *C_avx_int8 = new int8_t[M * N]; + int8_t *C_mkl_int8 = new int8_t[M * N]; + + // Fill matrices with random data + srand(static_cast(time(0))); + std::cout<<"filling matrix A"<(A_data[i * K + j]) << " "; + } + std::cout << std::endl; + } + + std::cout << "Matrix B:" << std::endl; + for (int i = 0; i < K; ++i) { + for (int j = 0; j < N; ++j) { + std::cout << static_cast(B_data[i * N + j]) << " "; + } + std::cout << std::endl; + } + + // Create bias matrix + int8_t bias_data[N] = {0}; + fill_matrix_int8(bias_data, 1, N); + matrix Bias = { + .row = 1, + .column = N, + .int8_data_ptr = bias_data, + .qparams = { + .scale = 1.0f, + .zero_point = 0, + .q_min = -128, + .q_max = 127 + } + }; + std::cout<<"Bias matrix created"<(bias_data[i])<<" "; + } + std::cout< + +namespace matmul { + +class MatmulOperatorNeon : public MatmulOperator { + public: + void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) override; + void mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params* params) override; + + // int8 operations + void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) override; + + void mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) override; + + void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; + void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) override; +}; + +inline MatmulOperator& CreateMatmulOperatorNeon() { + static MatmulOperatorNeon instance; + return instance; +} + +} // namespace matmul + +#endif diff --git a/kernels/neon/matmul_neon_fp32.cc b/kernels/neon/matmul_neon_fp32.cc index e056b922..e37c1d05 100644 --- a/kernels/neon/matmul_neon_fp32.cc +++ b/kernels/neon/matmul_neon_fp32.cc @@ -11,7 +11,7 @@ #endif #include "common.h" -#include "../matmul.h" +#include "matmul_neon.h" #include "pthread_pool.h" struct fp32_thread_args { @@ -120,7 +120,7 @@ void fp32_matmul_bias_cblas_gemm(const struct matmul_params *params) { } #endif -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { #ifdef USE_ACCELERATE fp32_matmul_transposed_cblas_gemm(params); #else @@ -128,7 +128,7 @@ void MatmulOperator::mat_mul_accelerator_transposed_fastover_column(const struct #endif } -void MatmulOperator::mat_mul_accelerator_untransposed_fastover_column(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_untransposed_fastover_column(const struct matmul_params *params) { #ifdef USE_ACCELERATE fp32_matmul_untransposed_cblas_gemm(params); #endif @@ -260,7 +260,7 @@ inline static void* fp32_matmul_bias_optimized_gemm(void* args) { return NULL; } -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_transposed_fastover_column_bias(const struct matmul_params *params) { #ifdef USE_ACCELERATE fp32_matmul_bias_cblas_gemm(params); #else diff --git a/kernels/neon/matmul_neon_int4.cc b/kernels/neon/matmul_neon_int4.cc index d43453e3..f0af437a 100644 --- a/kernels/neon/matmul_neon_int4.cc +++ b/kernels/neon/matmul_neon_int4.cc @@ -6,7 +6,7 @@ #include #include -#include "../matmul.h" +#include "matmul_neon.h" static inline void dequantize_block_q4_unroll2_no_offset(const uint8_t *int4_w, float *y, float scale, const uint8_t *int4_w_2, float *y_2, float scale_2, @@ -398,7 +398,7 @@ static void *fast_zp_no_offset_over_column_func_v3(void *args) { namespace matmul { -void MatmulOperator::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params *params) { // const int num_thread = 32; const int num_thread = params->opt_params.num_thread; int i, j, k; diff --git a/kernels/neon/matmul_neon_int4_offset.cc b/kernels/neon/matmul_neon_int4_offset.cc index 5cac2a74..fc472afa 100644 --- a/kernels/neon/matmul_neon_int4_offset.cc +++ b/kernels/neon/matmul_neon_int4_offset.cc @@ -6,7 +6,7 @@ #include #include -#include "../matmul.h" +#include "matmul_neon.h" static void dequantize_block_q4(const uint8_t *int4_w, float *y, float scale, float offset, int block_size) { const float32x4_t vd = vdupq_n_f32(scale); @@ -279,7 +279,7 @@ static void *fast_over_column_func_v1(void *args) { namespace matmul { -void MatmulOperator::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { // const int num_thread = 16; const int num_thread = params->opt_params.num_thread; int i, j, k; diff --git a/kernels/neon/matmul_neon_int8_int4.cc b/kernels/neon/matmul_neon_int8_int4.cc index 8b5bdf42..94482792 100644 --- a/kernels/neon/matmul_neon_int8_int4.cc +++ b/kernels/neon/matmul_neon_int8_int4.cc @@ -10,7 +10,7 @@ #include #endif -#include "../matmul.h" +#include "matmul_neon.h" #include "common.h" #include "pthread_pool.h" @@ -1293,7 +1293,7 @@ inline static void* fp32_matmul_transposed_cblas_gemm(void* args) { #endif namespace matmul { -void MatmulOperator::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; @@ -1342,7 +1342,7 @@ void MatmulOperator::mat_mul_accelerator_int8_int4_fast_no_offset(struct matmul_ pool_wait(pool); }; -void MatmulOperator::gemv_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { +void MatmulOperatorNeon::gemv_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; @@ -1374,7 +1374,7 @@ void MatmulOperator::gemv_accelerator_int8_int4_fast_no_offset(struct matmul_par pool_wait(pool); }; -void MatmulOperator::gemm_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { +void MatmulOperatorNeon::gemm_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; @@ -1406,7 +1406,7 @@ void MatmulOperator::gemm_accelerator_int8_int4_fast_no_offset(struct matmul_par pool_wait(pool); }; -void MatmulOperator::gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_params* params) { +void MatmulOperatorNeon::gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_params* params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; @@ -1439,7 +1439,7 @@ void MatmulOperator::gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_ }; #ifdef USE_ACCELERATE -void MatmulOperator::cblas_gemm_accelerator_no_offset(struct matmul_params* params) { +void MatmulOperatorNeon::cblas_gemm_accelerator_no_offset(struct matmul_params* params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; diff --git a/kernels/neon/matmul_ref_int8.cc b/kernels/neon/matmul_ref_int8.cc index ab7ae40c..fdc4eb9e 100644 --- a/kernels/neon/matmul_ref_int8.cc +++ b/kernels/neon/matmul_ref_int8.cc @@ -5,7 +5,7 @@ #include #include -#include "../matmul.h" +#include "matmul_neon.h" namespace matmul { void int8_ref_matmul(const struct matmul_params *params) { @@ -158,35 +158,35 @@ void int8_ref_matmul_nobias_ofp32_batch(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { int8_ref_matmul_nobias(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( +void MatmulOperatorNeon::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } diff --git a/kernels/ref/matmul_ref.h b/kernels/ref/matmul_ref.h new file mode 100644 index 00000000..1a666550 --- /dev/null +++ b/kernels/ref/matmul_ref.h @@ -0,0 +1,33 @@ +#ifndef MATMUL_OPERATOR_Ref_H +#define MATMUL_OPERATOR_Ref_H + +#include "matmul.h" +#include + +namespace matmul { + +class MatmulOperatorRef : public MatmulOperator { + public: + void mat_mul_accelerator_transposed_fastover_column(const struct matmul_params* params) override; + + // int8 operations + void mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params* params) override; + void mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column(const struct matmul_params* params) override; + + void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; +}; + +inline MatmulOperator& CreateMatmulOperatorRef() { + static MatmulOperatorRef instance; + return instance; +} + +} // namespace matmul + +#endif diff --git a/kernels/ref/matmul_ref_fp32.cc b/kernels/ref/matmul_ref_fp32.cc index 258610a4..5913b7bf 100644 --- a/kernels/ref/matmul_ref_fp32.cc +++ b/kernels/ref/matmul_ref_fp32.cc @@ -5,7 +5,7 @@ #include #include -#include "../matmul.h" +#include "matmul_ref.h" namespace matmul { void fp32_ref_matmul(const struct matmul_params *params) { @@ -29,7 +29,7 @@ void fp32_ref_matmul(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_transposed_fastover_column(const struct matmul_params *params) { fp32_ref_matmul(params); } diff --git a/kernels/ref/matmul_ref_int4.cc b/kernels/ref/matmul_ref_int4.cc index 0f456991..c597f179 100644 --- a/kernels/ref/matmul_ref_int4.cc +++ b/kernels/ref/matmul_ref_int4.cc @@ -5,10 +5,10 @@ #include #include -#include "../matmul.h" +#include "matmul_ref.h" namespace matmul { -void MatmulOperator::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int4_fast(const struct matmul_params *params) { int i, j, k; const struct matrix *A = ¶ms->A, *B = ¶ms->B, *C = ¶ms->C; const int block_size = params->block_size; diff --git a/kernels/ref/matmul_ref_int8.cc b/kernels/ref/matmul_ref_int8.cc index ab7ae40c..a0449942 100644 --- a/kernels/ref/matmul_ref_int8.cc +++ b/kernels/ref/matmul_ref_int8.cc @@ -5,7 +5,7 @@ #include #include -#include "../matmul.h" +#include "matmul_ref.h" namespace matmul { void int8_ref_matmul(const struct matmul_params *params) { @@ -158,35 +158,35 @@ void int8_ref_matmul_nobias_ofp32_batch(const struct matmul_params *params) { } } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias(const struct matmul_params *params) { int8_ref_matmul_nobias(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_32unroll_over_column(const struct matmul_params *params) { int8_ref_matmul(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32(const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_nobias_ofp32_batch(const struct matmul_params *params) { int8_ref_matmul_nobias_ofp32_batch(params); } -void MatmulOperator::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( +void MatmulOperatorRef::mat_mul_accelerator_int8_fast_2x2_32unroll_bfp32_ofp32_over_column( const struct matmul_params *params) { int8_ref_matmul_bfp32_ofp32(params); } diff --git a/llm/Makefile b/llm/Makefile index 94c34bb9..6960c0cc 100644 --- a/llm/Makefile +++ b/llm/Makefile @@ -5,6 +5,7 @@ CXXFLAGS = -std=c++11 -pthread -Ofast # CUDA flag DISABLE_CUDA ?= 0 DEC_SHARED_MEM ?= 0 +USE_MKL ?= 0 # customize define DEFINE = @@ -21,7 +22,9 @@ TARGET = $(TEST_TARGET_GENERAL) $(TEST_TARGET_IF_CUDA) $(PROFILE_TARGET) $(CHAT_ BUILDDIR := build/transformer PROFILEDIR := build_profile/transformer LIB_DIR = ../kernels +FACTORY_SRC = $(LIB_DIR)/matmul_factory.cc # For dynamic dispatching of matmul kernels LIB_SRC = $(wildcard $(LIB_DIR)/*.cc) +LIB_SRC += $(FACTORY_SRC) INCLUDE_DIRS = -I$(LIB_DIR) -I./include -I./include/nn_modules -I./json/single_include/ -I./half-2.2.0/include/ LIB = LDFLAGS = @@ -55,8 +58,15 @@ $(info Detected CUDA_PATH: $(CUDA_HOME)) else CUDA_HOME = /usr/local/cuda CXX = $(CUDA_HOME)/bin/nvcc + GPU_ARCH := $(shell nvidia-smi --query-gpu=compute_cap --format=csv,noheader,nounits 2>/dev/null | head -n 1 | sed 's/\.//') + + ifeq ($(GPU_ARCH),) # Please modify 'arch=compute_87,code=sm_87' according to your GPU architecture/compute capability (https://developer.nvidia.com/cuda-gpus) - CXXFLAGS = -std=c++17 -Xptxas -O3 -gencode arch=compute_87,code=sm_87 --forward-unknown-to-host-compiler -Xcompiler "-pthread" -DQM_CUDA -DENABLE_BF16 -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -U__CUDA_NO_BFLOAT16_OPERATORS__ -U__CUDA_NO_BFLOAT16_CONVERSIONS__ -U__CUDA_NO_BFLOAT162_OPERATORS__ -U__CUDA_NO_BFLOAT162_CONVERSIONS__ --expt-relaxed-constexpr --expt-extended-lambda --use_fast_math --threads=8 + GPU_ARCH := 87 + $(warning Unable to detect GPU compute capability. Using default compute capability: compute_$(GPU_ARCH), sm_$(GPU_ARCH)) + endif + + CXXFLAGS = -std=c++17 -Xptxas -O3 -gencode arch=compute_$(GPU_ARCH),code=sm_$(GPU_ARCH) --forward-unknown-to-host-compiler -Xcompiler "-pthread" -DQM_CUDA -DENABLE_BF16 -U__CUDA_NO_HALF_OPERATORS__ -U__CUDA_NO_HALF_CONVERSIONS__ -U__CUDA_NO_BFLOAT16_OPERATORS__ -U__CUDA_NO_BFLOAT16_CONVERSIONS__ -U__CUDA_NO_BFLOAT162_OPERATORS__ -U__CUDA_NO_BFLOAT162_CONVERSIONS__ --expt-relaxed-constexpr --expt-extended-lambda --use_fast_math --threads=8 endif # LIB_SRC_CUDA_CC = $(wildcard $(LIB_DIR)/cuda/*.cc) $(wildcard $(LIB_DIR)/cuda/attention/*.cc) # LIB_SRC_CUDA_CU = $(wildcard $(LIB_DIR)/cuda/*.cu) $(wildcard $(LIB_DIR)/cuda/attention/*.cu) $(wildcard src/*.cu) $(wildcard src/nn_modules/cuda/*.cu) $(wildcard src/ops/cuda/*.cu) @@ -82,10 +92,37 @@ ifeq ($(shell uname -m),x86_64) endif else # For x86_64 platforms with AVX2 - # For Intel machines with AVX - LIB_AVX_SRC = $(wildcard $(LIB_DIR)/avx/*.cc) - LIB_SRC += $(LIB_AVX_SRC) - CXXFLAGS += -mavx2 -mfma -ffast-math -DUSE_INT8_INT4_PRODUCT -fpermissive -DQM_x86 + ifeq ($(USE_MKL),1) + $(info Using MKL kernels instead of AVX kernels) + # Set MKLROOT (adjust this path if necessary) + MKLROOT ?= /home/elliotliu/miniconda + + # Add MKL compiler flags + CXXFLAGS += -DMKL_ILP64 -DUSE_MKL -m64 -mavx2 -mfma -ffast-math -DQM_MKL -DQM_x86 + + # Add MKL include directories + INCLUDE_DIRS += -I$(MKLROOT)/include + + # Add MKL libraries to LDFLAGS + LDFLAGS += -Wl,--start-group \ + $(MKLROOT)/lib/libmkl_intel_ilp64.so \ + $(MKLROOT)/lib/libmkl_gnu_thread.so \ + $(MKLROOT)/lib/libmkl_core.so \ + -Wl,--end-group -lgomp -lpthread -lm -ldl \ + -Wl,-rpath,$(MKLROOT)/lib + + # Include MKL kernels and AVX kernels + LIB_MKL_SRC = $(wildcard $(LIB_DIR)/mkl/*_int*.cc) + LIB_SRC += $(LIB_MKL_SRC) + + LIB_AVX_SRC = $(wildcard $(LIB_DIR)/avx/*.cc) + LIB_SRC += $(LIB_AVX_SRC) + else + # Use AVX kernels + LIB_AVX_SRC = $(wildcard $(LIB_DIR)/avx/*.cc) + LIB_SRC += $(LIB_AVX_SRC) + CXXFLAGS += -mavx2 -mfma -ffast-math -DUSE_INT8_INT4_PRODUCT -fpermissive -DQM_x86 + endif endif else ifeq ($(shell uname -m),aarch64) ifdef CUDA_AVAILABLE diff --git a/llm/src/ops/BMM_F32T.cc b/llm/src/ops/BMM_F32T.cc index bff8b7d4..653be693 100644 --- a/llm/src/ops/BMM_F32T.cc +++ b/llm/src/ops/BMM_F32T.cc @@ -31,7 +31,7 @@ void BMM_F32T::forward(const Matrix3D &a, const Matrix3D &weight, params.opt_params.num_thread = NUM_THREAD; params.alpha = alpha; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); for (int bz = 0; bz < a.m_dim_x; bz++) { // if (params.A.column % 8 == 0) // TODO: debug this @@ -82,7 +82,7 @@ void BMM_F32T::forward_weight_untransposed(const Matrix3D &a, const Matri params.opt_params.num_thread = NUM_THREAD; params.alpha = alpha; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); for (int i = 0; i < m * n * a.m_dim_x; i++) { params.C.data_ptr[i] = 0; diff --git a/llm/src/ops/BMM_S8T_S8N_F32T.cc b/llm/src/ops/BMM_S8T_S8N_F32T.cc index f5ef8701..ba2bd5f7 100644 --- a/llm/src/ops/BMM_S8T_S8N_F32T.cc +++ b/llm/src/ops/BMM_S8T_S8N_F32T.cc @@ -41,7 +41,7 @@ void BMM_S8T_S8N_F32T::forward(const Matrix3D &x, const Matrix3D params.C.qparams.q_min = -128; params.alpha = alpha; - matmul::MatmulOperator matmul_op = matmul::MatmulOperator(); + matmul::MatmulOperator &matmul_op = matmul::CreateMatmulOperator(); if (m == 1 && x.m_dim_x > 1) { // merge each batch params.A.row = x.m_dim_x; diff --git a/llm/src/ops/BMM_S8T_S8N_S8T.cc b/llm/src/ops/BMM_S8T_S8N_S8T.cc index 1dfb534e..1835366a 100644 --- a/llm/src/ops/BMM_S8T_S8N_S8T.cc +++ b/llm/src/ops/BMM_S8T_S8N_S8T.cc @@ -41,7 +41,7 @@ void BMM_S8T_S8N_S8T::forward(const Matrix3D &x, const Matrix3D params.C.qparams.q_min = -128; params.alpha = alpha; - matmul::MatmulOperator matmul_op = matmul::MatmulOperator(); + matmul::MatmulOperator &matmul_op = matmul::CreateMatmulOperator(); // process each batch if (m == 1 && x.m_dim_x > 1) { diff --git a/llm/src/ops/W8A8B8O8Linear.cc b/llm/src/ops/W8A8B8O8Linear.cc index 33e6098f..ade99e94 100644 --- a/llm/src/ops/W8A8B8O8Linear.cc +++ b/llm/src/ops/W8A8B8O8Linear.cc @@ -55,7 +55,7 @@ void W8A8B8O8Linear::forward(const Matrix3D &x, Matrix3D &output params.alpha = alpha; params.beta = beta; - matmul::MatmulOperator matmul_op = matmul::MatmulOperator(); + matmul::MatmulOperator &matmul_op = matmul::CreateMatmulOperator(); // printf("W8A8B8O8Linear-m,n,k: %d, %d, %d\n", m,n,k); if (m == 1) { diff --git a/llm/src/ops/W8A8B8O8LinearReLU.cc b/llm/src/ops/W8A8B8O8LinearReLU.cc index a965b007..89700556 100644 --- a/llm/src/ops/W8A8B8O8LinearReLU.cc +++ b/llm/src/ops/W8A8B8O8LinearReLU.cc @@ -57,7 +57,7 @@ void W8A8B8O8LinearReLU::forward(const Matrix3D &x, Matrix3D &ou params.alpha = alpha; params.beta = beta; - matmul::MatmulOperator matmul_op = matmul::MatmulOperator(); + matmul::MatmulOperator &matmul_op = matmul::CreateMatmulOperator(); if (m == 1) { // let's loop over the column dim instead of row diff --git a/llm/src/ops/W8A8BFP32OFP32Linear.cc b/llm/src/ops/W8A8BFP32OFP32Linear.cc index 0702e21f..dbbccd06 100644 --- a/llm/src/ops/W8A8BFP32OFP32Linear.cc +++ b/llm/src/ops/W8A8BFP32OFP32Linear.cc @@ -52,7 +52,7 @@ void W8A8BFP32OFP32Linear::forward(const Matrix3D &x, Matrix3D &o params.C.qparams.zero_point = 0; params.alpha = alpha; - matmul::MatmulOperator matmul_op = matmul::MatmulOperator(); + matmul::MatmulOperator &matmul_op = matmul::CreateMatmulOperator(); if (m == 1) { // let's loop over the column dim instead of row diff --git a/llm/src/ops/cuda/linear.cu b/llm/src/ops/cuda/linear.cu index 6e65c980..4868ec67 100644 --- a/llm/src/ops/cuda/linear.cu +++ b/llm/src/ops/cuda/linear.cu @@ -32,7 +32,7 @@ void Linear_half_int4::forward(const Matrix3D &x, Matrix3D params.int32_zero_point = this->zero_point.m_data; params.block_size = QK; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); op.gemv_forward_cuda(¶ms); PROFILE_END(profile_name); @@ -69,7 +69,7 @@ void Linear_FP16_int4_ref::forward_ref(const Matrix3D &a, Matri params.int32_zero_point = this->zero_point.m_data; params.block_size = QK; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); op.naive_mat_mul_fp16_int4((const struct matmul_params *)¶ms); PROFILE_END(profile_name); diff --git a/llm/src/ops/linear.cc b/llm/src/ops/linear.cc index 3be00536..63c2e6c2 100644 --- a/llm/src/ops/linear.cc +++ b/llm/src/ops/linear.cc @@ -62,7 +62,7 @@ void Linear_FP::forward(const Matrix3D &a, Matrix3D &c) { params.opt_params.blk_size = BLK_SIZE; params.opt_params.num_thread = NUM_THREAD; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); #ifndef QM_CUDA // not support yet if (this->has_bias) { params.bias.row = this->bias.m_dim_y; @@ -109,7 +109,7 @@ void Linear_FP_int4::forward_ref(const Matrix3D &a, Matrix3D &c) { params.zero_point = this->zero_point.m_data; params.block_size = QK; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); op.naive_mat_mul_int4((const struct matmul_params *)¶ms); PROFILE_END(profile_name); @@ -147,7 +147,7 @@ void Linear_FP_int4::forward_fast(const Matrix3D &x, Matrix3D &out params.offset = this->offset.m_data; params.block_size = QK; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); op.mat_mul_accelerator_int4_fast(¶ms); PROFILE_END(profile_name); @@ -201,7 +201,7 @@ void Linear_FP_int4::forward(const Matrix3D &x, Matrix3D &output) if (this->has_bias) params.bias.data_ptr = this->bias.m_data; - matmul::MatmulOperator op = matmul::MatmulOperator(); + matmul::MatmulOperator &op = matmul::CreateMatmulOperator(); #ifdef USE_INT8_INT4_PRODUCT if (!x_int8) this->initialize_memory(params.block_size); params.A.int8_data_ptr = x_int8; From f69c9dad973f7e3df090c4c4d797ce35c26bd31e Mon Sep 17 00:00:00 2001 From: balbit Date: Thu, 5 Dec 2024 15:29:48 -0500 Subject: [PATCH 2/3] dont commit tests --- kernels/mkl/Makefile | 57 --------- kernels/mkl/matmul_mkl_test | Bin 80040 -> 0 bytes kernels/mkl/matmul_mkl_test.cc | 220 --------------------------------- 3 files changed, 277 deletions(-) delete mode 100644 kernels/mkl/Makefile delete mode 100755 kernels/mkl/matmul_mkl_test delete mode 100644 kernels/mkl/matmul_mkl_test.cc diff --git a/kernels/mkl/Makefile b/kernels/mkl/Makefile deleted file mode 100644 index da6dc430..00000000 --- a/kernels/mkl/Makefile +++ /dev/null @@ -1,57 +0,0 @@ -# Compiler -CC = g++ - -# MKL Root Directory (adjust if necessary) -MKLROOT = /home/elliotliu/miniconda - -# Update LD_LIBRARY_PATH to include MKL libraries -export LD_LIBRARY_PATH := $(MKLROOT)/lib:$(LD_LIBRARY_PATH) - -# Compiler Flags -CFLAGS = -O3 -Wall -pthread -std=c++17 -m64 -DMKL_ILP64 -DUSE_MKL -DMKL_ILP64 -mavx2 -mfma -ffast-math -DUSE_INT8_INT4_PRODUCT -DQM_x86 - -# Include Directories -INCLUDE_DIRS = -I.. -I../../llm/half-2.2.0/include -I$(MKLROOT)/include - -# Library Flags -LDFLAGS = -Wl,--start-group \ - /home/elliotliu/miniconda/lib/libmkl_intel_ilp64.so \ - /home/elliotliu/miniconda/lib/libmkl_gnu_thread.so \ - /home/elliotliu/miniconda/lib/libmkl_core.so \ - -Wl,--end-group -lgomp -lpthread -lm -ldl - -# Source Files -SRC = matmul_mkl_test.cc - -LIB_DIR = ../ -LIB_MKL_SRC = $(wildcard $(LIB_DIR)/mkl/*_int*.cc) -LIB_SRC = $(wildcard $(LIB_DIR)/*.cc) -LIB_SRC += $(LIB_MKL_SRC) - -LIB_AVX_SRC = $(wildcard $(LIB_DIR)/avx/*.cc) -LIB_SRC += $(LIB_AVX_SRC) - -SRC += $(LIB_SRC) -# ../matmul_int8.cc - -# Object Files -OBJ = $(SRC:.cc=.o) - -# Target Executable -TARGET = matmul_mkl_test - -all: $(TARGET) - @echo "Running matmul_mkl_test..." - ./$(TARGET) - -# Compile .cc files into .o files -%.o: %.cc - $(CC) $(CFLAGS) $(INCLUDE_DIRS) -c $< -o $@ - -# Link Objects to Create Executable -$(TARGET): $(OBJ) - $(CC) $(CFLAGS) $(OBJ) -o $(TARGET) $(LDFLAGS) - -# Clean Build Files -clean: - rm -f $(OBJ) $(TARGET) diff --git a/kernels/mkl/matmul_mkl_test b/kernels/mkl/matmul_mkl_test deleted file mode 100755 index b78b4a4e653ba6876efb58c682e104a8439c67c8..0000000000000000000000000000000000000000 GIT binary patch literal 0 HcmV?d00001 literal 80040 zcmeFaeS8$v^*_E7HV`p5K?BALC6=X46l@|7%_`Ik?CNaYNE9N8f&n%`s6YZFLX-q_ z6Uua5plVxeZJS!{hpH{K`mxbk+)Z}!#8VQU0w{!lILm_sMA*cz`MuAb+08O0t?lRY z`+oi@yfSm|x%Zy?bndz5o^$T(8ms-jD2^Lo(EbfDeBZz!oH#}Wq-*encM??nmu^Ti z3^Ak|k_}0QfdCotm#!v#<=q;kCWPKGT0W|f$|Yzu>)#`^e4%%!c6w8}{{FSzs-}eA zhB!4&<)SH=NDur=WY6&X(93E(^o}gY>hoy%b?;fDRCwqeS&qi0Ls8T9@039*yfQQ( zT03dDsNBM{YP$ZNsf}mo9V$+3q`x$6JoWDs6`djU)|XqTmD9g1TDj1h>Y#r!J}Juk zck0X3%6WFGc=YetT027TP<;=f9F5D@{ihA5;bE;l{rHFCaoRkD-b5$#@4Jf^&7U;k zyNmP2Enc*=c-6R7X_LlHnsDcef;;bG>C~=7{1Hvv|6rB@)8ZkR)_gQC$g{sbD)SN;Oz?F)Vj8r~QDul?Y$ z^#lJ=Kk#S!srNPvOkef7`pNGBd?^0(|916*XJ9}58Vdx8dL)SR{opC?r(f@*OkeG9 z>jw|l51u>wY3IIv@R<97|F9qU-Tl=2Qa|vbe)`4tlm8U*_cbpW{j}$JKX@AZfiLW* zJ!|^GPdj8^^Fph%FZexxC*n{4@6mqRv!x$6F@V1Cf7VZaG~e(oeWi$}ApcOq2t#^@ znixkJ@$WbI`xdxP$^;fQXzb8n$jx=EnwPs^(b9Q~7d@SC$hFPP$j!@Np8v$66-D{W zXJ$-WT(C5M=Dhifk(T?!l7gkVD~jeVFUrkTv-X+MF@N#A6}eC3FIkeiBCR-W#rV4n zOXe+JT;MP)Se~D6C|a}x$R5pGocYueydRuVuy}=ak=BwYisvoQ%bn*~R=j9A<$Cae z8ATHvMJ0v#HitDWuON5jqP+a2RtJ)n&ZFWpA9*k}g?h1S-uy*RrKWxlVN|~&m&)WW zTDqXX`qf-_p~Q+sePm>#p|xRt!SW&$ol%rJDL1!h;qro&x%21cEW1uKe{ z=g(V$tSj=D7ug)BGu5$h-tyd{Vwz8)3&)eWj)hO=E||9njma%szG!LD z0)ild89~p~q(W4XPa|1?@-#Yz5afpii-qlwabgfgp8O)kxeL!;m&^HA4{<@1*2h0Chwu4$19Dqi~4I#yz448XVK z$;IeIQ9hoF3q!Jl27h8{F;NaWtBMpF*io=NKMZ)HU`ZiGQVH3LqCCePcTjW^MW1k> zwc$jE3cr8awA}G`rWo$G+on#-z3a~LcTR{1-_;xL1^0&Tx^tqMAv+sLCfsShCWe(d zEjwG7YRgRp+O%uZQd1++B2yy3Ku2;62X?0ZH$rk6i2sA|e;{BS-Z2PWhd+7~d@w=- z4A&zRi#LUbu$-LX29zGa{sy8X{p0Ze74XWC{F3?YlIsn?@D5;VdbHsUEqyKkMYk-v z&Hy|gXt3oh9%HE2;_v+AXArNV4g0is{GvreNWwm*#f!f|k`4+weG##PN*(Z1^qH)`>DEd4seC@sEs>yr>I2N~|v;zh|XKx7|i zFl+G=RzAvLMV$VHfkB&t9Tab;G4Ak zg*tep24AIvZ`0swb#RXcU$28VY4A-tc$)^_tb=!G@Jb!rppEAi9Xw8hZ_~jOHF%W{ zo}$4$I(WJUuh+pdHF%Q_K39Xc>fnVMyiEs>)8<{#!@sZ2ONS1=R>Kp}!8d7e!)di# zh|ViDc#ICdO@qhj;2sU0po2GQ@I)QFO@pWC;OW}AHcbbQ)7mZS;7M9MQwOiiR{N8! zgKyK|b9M00dReH0cRZltDb&H!?J9hg4jx)BYjyC@dRec7S8Dh->EN5Rc5c?en`Wr> zR_frJwDxS#!P_+WHXS@AL&a02gRj-#9v!?%gV*ceg%7HDnsjiF25;5D6E%374xX;j zgQA0HYVZyne69u$=-`DK+;B#1SBRcN>m^1H*Vap%4!%~ylc0ld(%^|Yc%=qU(!sZB z@Dv@~qrua3@FoqOu7kH}@Jt=NQd=+CI(TTkEY!h8EncXDd$j#xl@8vd!Pn~Gq4ly} z2lqUtj{7DZeC}KozF7wkt(Qt2JhWc6=-`Imsd%>O;7ye(yh;a8d0mBj^zi>u;q`j> zA5?gg9{#2ZZ`H&9sKVQH@TNbja771S`-BSb(81Fesqlaf?s->*8_ufl5Iv`B^bn(i zXKL^`9el0^Ptd^&HF%;99$GI+dbqY;QgraO8lE&Ae3J%G*TE|_xTu3~)8LspxJQF$ z>)=fqe69}Orom&h^Ll99Lh(2qyignW(7GY1I2L~-Ytp|Yjn5E#B>qBhQ7fm1x1Og` zs*VwYYxrMV1pM|e)L>8|;CuwUBLc38k}Ne40Z)#|pP==N+LI6g59uwU0lOmt-t?Io zA-x5a3H@u0fI|ci|7(kYYoaqtQzGD0N9bQi1YFg(swsg8IMp5gr}c|qp*f?pm9R#ctQmH`UrSp1bj#YJShTxLj*h}0v?)kN=}P_N9x(qBjAyGCouwkQw09Z z2zY!1JUaqDECN0^0zNzfzAyrQa|FCF0)9&bd{qSe8xioe5%5Um_xcEU$fiWaHbuZS zQx3znIRc)@AnKzs0)ATrd`krUTM_VW5%AFw@Tv%S=txG%o(Q;hgkz=ZBjDeOz|#}~ zPl|y5C-|Qc_|FLZX9WJ)2z<_evrBZIjS*eZhxQl@VuhzDs;5bG`eJHW(VmIF*ljTM zjClo5epI@F@-L(me|t|)&w@=XP8+Ykr8iCjgTJ~rP8+cQFTHWvc>S;T#%aU#zt9_} zjn@BMZ=5z$|MK2AZLI$Jy>Z%5{SWuXX(RPd?Tyn0>c6WuPJ)sDTfK4GNc}hV#%Tlf zNA<>eieC(d`%4?Czr8n38>qjfH%=R;zq&V08>as+y>Z$o{jc`MX@m5?&>N?X(f?d; zoHj)N^4>TJM*jJ|aoQOD5BJ7tL-bGWjnhWxzpFP+8=(JNy>Su<{Wtc;ze(|^-Z*WH z{)<<`{ZFKLdpO=R=4ZPMh6Vg6i-FNc&zPrnUK?Mki!ac{AJ@en(ZwIs#qZO_zpsl= z(#7x8#lNkKe^VD9u8ZHGix1SrgFCM6?`2*5GhO_YE`C%OKcI{6*2Qae@$I_!ySn(F zbn)No;=k6#f2NCX(8bs3;$^z{Q{lMivexH`cNoWza8ewt*s-dV?|`)Q3F9f@1|-;} z9yzd5S@D%TF`{%#`3PX;Q;;4%YArA_dL#a&{0>Ye)x#Lqf;;(9<&DezonkDNzIRoQoHCp*^xQZn(9ey z5ZzOrMT>iez1?UqT3Xcd`HtDJ`Gb{Dn4AF7aCVeDw zUs@|pmSw8-XGov1^{(u|T-cnooakthyyBvk7frSoyq%*&@99LY+S?VC+~|(ok0u$C zo0IFE$^c>geLXhw(c)UC%vsEz6uTHs%e1r@FzkR}p^z}ui zK{l7wAiB~On{#z^oPjUjj*KWSx{>xi?Sjf%$d7Wsv~M#o8PyERPXXq)7mYOZiRMRN zL6fhd3D@U(C7KS#V#N`@oJ{s>pGh+mIr!}>rcp2ci`p*LDQoaf^_+>p{8mwGMAuAH z3|3#fy8804Iug|3#$c=TO~n_YnEwZi2H2ovBYvh_oo8Be354D?VEbS~=^oa3Jz-8y z5zS3}c`1q^PfruSy~AIKi1LFj%8j||UG`^`d;r7qoAn{7A3&qRETS`E<_yt2^6ff+ zoG07FXV1>X_(x28xhWffRFAS2NVK9BQ;M=IGMc2!kBlZNGZ0Nh(=4V0r4C^X&s_B5 zecHP*JNNvYT5wG1MiSQ7M5-MVYj?difCf6F;v`>w5*-%St~RA%riW1&GiaKIQo`S< z14;^&1>P8j15L8KD=ES5Dl(;L>q&VNiL^eXYH_o;$V1CRYGw;WN#vTmfvDtqcdU#m z4awf*M(1fPkVkql%%_TF!Fi4=emm6QU;a#O@B?at%WL~+@c4)Z-vyya-{3W!YJ+Fd z8&i}N(cmO?zLhVL_HVTJlGfgL-X~!U1Kj;nwY|S&bw%$(7r!#VPHLkzz?l&0!tJ%Y z$J~wyQg#Q_2LGPkXfQ^2VY=PDEghq*@gaAqhvZTe{T^n<=B(osz%k-b3}SU4!M62OKMDdTmZQU74X(eFxqd<}ZqW7V6kItz*@! zP0?TaLLGZcLwb%ce}f?{1=7htT0HH{qrK~`N2y2qY_37tdK&TvWLDQxebKYn{nWKo zgtbJA^Or2M)eu^J&xCU~@f8y(0xJGJGEq;<=&erd*BB?yNr-Kl}**Nj# zFGF6xW~{!WRo?@IG*&N&?s8Kc7HzHiGSCiecfVsIuhlw>DK4F81WPkb?pjlBN(Gb! zhK2&Sg#zga*j-qH8PZjCzj+LG(JV)Tpj`*HdlsMfcRJ!utEss4K;? zbFnFoEB-Ua>wOZ)L8FUb1dVl}(cl2%ao!=~5M3&9h!r4?cAP|rHu&<@=(9#1N*a~G zXvg7A(@RqTn&Fg!FRmTX-)Q4;82I;AtN87%2s(KV_=UBg6TaM!t%A`CklEd>oI zV54gOfHx!?g4w0udt}hC@P8fmi(k>uOne5f&&4-U!y<&_uDIxNFHq-Jn=%(zq_<3O zv7s(EZNY+IONc?cHRyVJ6TvX3o$^TKrdAD_$e^!i&?JT-`bJjG&~j6{2G3-0Hkp8K z(x9v8?T(K9JE63gG7)b3n;OopD>uD{WDqZwXpw0y;*|*R+G)SnF1=&ghA==Kc>Ut# zml=q3yGr76lt*I|FbvdD%TK@ssE^cc`NDY!Wvflu)B$QLVWMIpP^I8a|0tZMg+ZGZ z3IZf*AUU|^c~cA`a+lUWl0+9JV{7-j}U^8jj9P=AF5$0}S@af7jV59cf=#61bvzcYS;u-KAYwEc(Ly4ko z&t9d*@#Qf8e!q2+E)=6Q_W!1lW*4 zxmQD#&@ZY>sQ+42*Q=;(E+gtpS2A308Hvy`D!{fsr(sfOPNE>OZmQ>g^n^xg6!nN4 z7tn?io74{fjnZ7vmmmtNln^06{Ep#M1&DzvLRElJzPO-b+(mEjw-m*pg_4vUEtIHC z*Fp)3m4#GEC@fd;7357Sy7OBfQ$=VYcTNp95akKvj@V#C%Ebh>dcItQ{1Dz`?olJk zq-zV@99f_N0)Du_2Z+2sULAqkFarKgc49lPTERNusuAVNdFp&`9S3QELK5;$uq@*H zS>gyAajjLmS>(L(A}i55z$+q4%u+L{{rCYQkjAO=8&JqUv^QQx@pmI@l#E*)Xv zoc_*d5ark4(R;KC%kNUhwdZ3T)Z8O$h`Hk+&c7RN0M9S?;N0r;;mqFp*qn$Hx9A>s z`~&24k9>jc2IrMWQ3;O5EuwT3hl?VDx~Ggm@01j@R)h%DiX4663h5&w&@+AHcG_6{ z%g=@MRaw^@HRz0iO{dkhp%U3D8u$%CnqbGe*ANw@D-|Ao!&yKqt|yG5>v8nJVp39D z?5@YNMfZdXWWwD3XbWCe>9`ea3h5n!+hXbwy#4`#^V0wTCvD#2(M&Ft4y5*2`K_01 zj#|nYu)B>;oTK>-k5a8@NWF+{pj8S;Gfi#$*865USDQ`_lWQxC*uTFu>jIT4yg>a<8*R;gV`ZWP?hdz`_jr|za!pmAew5FH&Nx7+vIRdHkAv`Edq z>W8X~Z#$27bB%(teSjReu^z}ga-g=Ju+|qnBDg0EcgMzwjt0?j8T;~tw|&*N_hMb{ zAK-w0ymXR#LX3sG;H-{9GBWwRPsdGu+t)52i(uYYHe7JW4(FPv_-2%z?DAD3$w?bc z2lnCEDheqgL_X16SM-E7TH%r6e{>uyh{h(hLGGBE4$jH@D^xXA$gwkFGTjPDD!>B z;+tGkjb_1EY;z4XbHyg6jS}5MzUzVnK}u%q8zzZfQ@)Fra_lq++dU+Xl1aym=b$SBoGOl!HDF^Gm72io+hyg?jr zUyOLi{V~4dUAKGWuAJyUL8H;D&(W{X|HQ)zN;ZH0*W;n`%c)8)zjSY?WEAj7<~6KS z`kZNKoM-1+T)$rr=45xhn~j(96t)yd_IJeeI>UjC_lUS$0zM6+QBW81b*v5yEGdr0x;49Kt$*UE#SB&o8%GQlW}TCMc8E^j)_up zD)!{lJ|S(E$BKjAhOLlzEUpJ;iYOSTf{Eq73Cc$&XPeS2{MH8hfXA>KG}|0r;^gUq z+jzHCYNkrFUk2|VC0yyYa?L_>r?}ZCJRKLmj-K{)M#0(5@f*f6p$g?P9AHwusig_* z(O~1ecBwUk^RxZc=d~U&4)uA<3QR*qr?QrqJ9?Ya1m?4S^sX644*(c=SsRsN(AC4X3i|idfq4d!s^f)PBebnP~j7(_b zzGZ2!p^G9yZ_D}=P`D*4Qe8Qo*8e%HSnyWSQ95u=3 z7?$AY2-yXW5`Es%M=w)HP_Bb8%keo26Sc~;e)`&nZBAr`Ph=!qV1Stcerq*krSYOm#bkHK-s78J znh=!hYd-@!b&oLE9j#D=yn}*1VRh0mRJjdd2vqNwHd7#w{FbQ_0ZqO{?b`jvP<=O6ti|VDJbNz!M||GA*+(dY2TOPM4&VG?lU8|wNa>; zm;~}%(-_9wsiA5!5D2#z+y(8FyOY3JM^4(D$VOqv<}0ji^kD7XgL>XEt!4FOP1pCdZ9s{`r7)bu5jGO-oSVI@aSswB*;>Nb!JnDnGgtVt)9#O?)x zVUsqeQ`IN(V?XMUPv)p2>1)s5{E{}5#hcaM$mK|+;llttuk~YG)hld1eC<;-maTz%PA{(pzvQ0AVilc^BerQSd)!Pj1r@CsW%6nc&6VdbuZ_`9UTvj0eQ6QpX<9lPGy-XZ&!wHN^# z_5on>syLw#B~m)%&K0dO6!OT;D~?iV{AGI7hZ4r0lLL=25KeH9QKUg`ntdK`#AnyZ z9kchy^|L`Y^ERKQhmRgR>2Zm13^vLEP-`NimXefe*}vj&4AP$0hPx!CMZTc60?P+= zpJq8I)IilxpoP^%byFLtPPGkyYXmAQvk%IFxMeBwswA9d zZ)ro3TD1qe{PQXaDcSoGX=J;+dWOD+oZeyN#JDOkd`?}D&J#Re4t+C3WtFVJZC}#J zB2J|gKweWhdRN?A42dXJi&qYV(we5bjo;JwyTHF2Z;(B+dZYuM(?|P@Ut9|XuUN(9~r-zeUvK7HBuX&GgcxoV#>{`4n+@}y= zY3;1Can%6Q9%-)QpMR6t{jh~mS~^-8?%@O`t%znwDSSEMkQ)8b$c;SeqLS-pT~2Mm z)<+A8{xP!fAmINm5?i{4s`>n<==2_rEdsZwGj5;)Bo5iw`5s0(oc`_HZowSjE7B1C zs?-C8B9o62AWD$1G?%toQaFT)aVU1Cjt~1MXlCoI8Lo8-YD2>xOvkuwpq|7m=L_WKtbS8MlH&ZW* zn{DoRgUG4<;I~JiCD-A2fnIi6l3S743x5-HSMf>t^z9f2;J{JdD|e0cKnHR*xr1y?1<&)*9` z$lXK8g@Qo!CtQ24fv-49rLyeqf=pqi-Tg!ol#0@qq}fpA0Mcf*9qQ>>dyk>$RzFR0 z2!4vJq(42(l8RS}t|t@3igQJu!oMgEM>x{YO-m46(-K9qEb^9WL=%;Dz$Cg#4Lvhr z)|SQ?igQTKI1fOorUY;2lW`p3{X6gyEK98#Y#v=8wN9!h>`1W|jbY}96vBFK8rDgv zK^cUEGq0q2YE=1}Bxymql*N>Wtefu1G6{y%UFA_!iP<6_l2u8{Y{Ka-%}l+lEJKhC zsTL_eUXrOS7RQP1XK-40GEo^p{X^9#s8;Pieel}$ zKACW4S8w}C{{+_{(Orlymtt59L*CQ+wq2XSP)DnbStcja%?3jk&S;IhV}FZR?Uz3qXy_~W%0 z56l==fHDWU?wJ$69A@M>pu;1@z)9nJ{;@fJ4r!Z2*F8$$fQqt%-fF(0R*Y1fC|MKM z!oXj2{LDdZy=(eKof{O4sc*v#NF?eunw0O&LKuuI!sR9!^(NFPe@7LY7W6hZq%Zs` zK1T$;|AJDcZ+uI>Holu8@Lhqrv2T1Mzc#*_U-$NR<)i)K+qM7e_IE`DzVFQL4`1om z#y2_w-@!Tk;T!k0@$JP?HZ;FK%Iy!|$NRo+f1MHd4m$e7=l(w>i>u3O&mMk;Wt8VQo`N zhjD18-BDo76T{w#)_T_<$tzhu5Uap(>E9^qsXOhC3l^!k%5Lu9pF5Ay_ASJnqMsSZnj6|(Nf^4}&p8l|k)L}`yGWo%J4uspAcu8b`^}F?#8Z<5 zsma-cGsWL=rV!nQTn5yV$rn;D+nk->A-RO+qGe7R_{&-?K0#=nZLza`K)pev_ZZk<^# zw>{DzcV&BhCrSo;2FOp3^}yq&xRT%QWGKxZh&i4!Z`q}TN_{g40X>Qb?`zzaVLu|J zt~?Zm~$byUT_z}V27!1q0R(WZYQijSZk4L zou_baE`h3|v)h=)T@ajYJgjYNjCb?pgF%y`qZ8yXA#H*4GmyiCyWPt;^F@Bcc~UAm zE+c6&DCInoCg0^=7DY+>`b-)?N$>TUWTd2D_n8z;NzOi#22#?(K9dGfQbwOiF_d&y zpGntI(#?G)4W^`S93#HkQZM_-KRXqK;ZLRGOkkydruKdW3Cfqd(Ny=4lUQ``CxN<{H#{b&o0>=m&5!ltto<^wWB4jj0k?VOXX)5 zqzv%0E231qMcK~sfS(mJe%4TSpT#X;P0av9YY?0PkoB94qdm=u$pN_UIs;&5q;v?$ zqcn_nYr}ZAY7`EFpSg&iy&vLdGen${Dt-h^{Pq@NWh(C?e&!;6MqF&rlOgVv0q&LI zdx$Z)2!^F`uwJfpff^slwcZYKt+j*`$23@=!LY7d$!Zum6z@$^z0do@+l?C;*K#?D z1F6g_PGep;Y}MoJ=pr9BI&eEhTGwA?P&kJ5!Q^gGPrAE{e@={Q)Gn1#-Hf-6QPKAX z`(aendMIcy((5J?quL{|jUn~2_k1)N;_5~3rz)!=zXR~8yBMFM-2{B<3i#A+=czV9 z>M-x-%cFsT<}&$$;4bN5d(>&_P+%ove_&K~7HN<3Qsf%KE2u!gYO=03-t8`nqQom~y?-@v z03{yoC(%fWfA1$Tni4DfNgPOtDnI*bor5TGK|dv9DACqW;&qfbwx7hol*sp!cs(Uv zWZdSf4T`12{rw~kp~NaR(OG+gZg|ZX_zk~giL^(=-xZQ3{ObTr?eQ<7*Ja#_@!z5` zt|{7^ZZRl#H!=qHB-Y#e3qfgMU`z1^1A7iHs*~_s6($69v~mi))REtSwta*HbYzmly-B%tr+1umM1p7KLJM zq4Fteu@A(0u z;|O`Hcuz#x-6JY!Wzy;c%Yg6EPbq4Vn&?_b2RV;=ptB*vUK`n?jL#7~01p&;Hz5fZ z4!CMN$FkOExMP1p2}*|Bv`*wMYVu)*d%{gtvR?fTxS1D?2U>J8zNqJaSUi$OfHCD) z#YXw6^Q(*lp%28euPgopCW>1cIk*_X??&fpL`l_@R+EXdTJDI4{bJQk zoXzQt5xK(|=AC@SQ;=b-=JRX51HB#ef*?HJH18Uc(Y z29b5()6n2xR%uKrCk+D3WeNt?`6-9-a8Fn#xX4cMeT=tEfOuzB3`+XityNKivn!_T z`xQM!w;)GZVioBzC}C^fh2R+aHN?X~BEJu67o0NO>nhPV1h)vzPX_T7FCwpt zdBEAEi#UdvTloiDMaK>d*e+zmYJG&q4!&8G4v5dT8^{ea`Jf1SQ#5zuhTF4e$xT>v zK5q&@bzlIn_{_7XdJqQi*~f7Yis?t~uF^`Iv@=5ro{3}f6wVV>g6lcbLB5F91&!HHscwKb+qk-3{$+E0pQ@>iQ684*GQ^eQbJy&|#SiFA?%FyOo%RUreNZe4jRNMYw%Kj=L2s&ibsU>r}<~;=mGrDc{W{i{U!h` zcGnXzqHE(81TC)Vn1k7-MA4O-ZMkQ9CSP%$$l!bCzg8lduNXzF%Q0)F`R6p8d_@OR zF_#&-{NhBHv;%rl2zJcg$~SZWK0TNE?)5BE)%>Vi5;d05c6!__UP0W#73Z< zaEBy69g=$~LmKQXHp1$j44c*oZc2gn!QnVev5E`tN)Pz!g_=}{dZTOcFC z|B;*=UlPpTvYQ2W{6Marif?AVkuNddNag^UXyF>GRE)a6o8AUTZ#`%gS z5DPeZ0&^Fza=S%tHGDkara03S&sUJ%SL2{M{+Zgs93?%>c^f3QHV-6*?HRN+F=iek zu}yP9V(Hb40_RaO&Ki{L2OJbwLCLqN$x|tLv6|edCYMrj6q42b8PdN$58N!oufpdM z5e^@Hs8~R)*uJ66bg35?MJcPEJhx8S96uvp6PxOJ0<*?X8Ml$gu{qr5~T-t`y zd(KHg>N4-)-6v7A)%9SeR9Yz&zhrZ*qXj0ob9iaaThrVlu9N0$VN`vQ@<}^zB+7sX z9r7^e;KN|Vn{DtlCas664{%xELIL;mQ8+7A3B1kg4H&_*s^P3=+vDww=ISN5nZfA{ z)HGrpGHv5Jn>!JS#|C$w(ZU@9waoF5zb$+$j?iugK$g$uo)Alr76>aOG+;F`2e@CL zEi^xp7w&7UP>S&e`Qu~t^3@|g`HN*2;1G7ki;{^U*9rtZh>z~Xjt5sW3TjSX6eK?d z^28X8JWDozi5Nj(d0w)42R%MVj3w2`-9sm1FCA)sE)mb7=hE@CJU5qCH3?7~e@Gi% zchbX&aL>p|$BJOFxfB!UBHS7Ny+ zh7WssNFMH!pW9t)B3H;vQ(9_^;M!P8s9e9HWs>3gFLKZ#gW+`pK*h6+uJ3_I<0f4% zAyQEny8LrY8BHj#_n`quAtWIi7Z2+3F$c2O&1d~*{7hG9J zsU8;DvG8k~ZJGd2FkiJwzz4Yc=74*`AS5Pnjo`H)VPcvvYyO<%0UK&1#wM>$iUK)2 zjjL6YQxRAWib(q5Ts)dnTR{IDqw^-X5;bEDfJRe{gXB5K+KAG9-jT~L8R1lT#wV8) zz*&VkT}~}cJc)duccG-<61h2^u==sfa_sESDR5FAG5Z9O zx!et>A>=tn&~tM9>@&XVDOCr2^OtP?R9>BY8h$mKYf*mle%6?Q{|AlP?VDfB8siT) zW|w&{zi~gv9Bk;JQP_V>ysg>W6^$7NKk{&gz>%a{nmzno<#o;uNP4d|8keiJKPPFX001g(KVog}`4O=3ZM#WCuRS1US? z|D)aRxGGEUWC}8p2US754Kn#QxHgA(&U(@8>ubmK>37TzfkSYUkM7*L9!54YjaylbC4mOIupA|J#Vq>QoT02FucD^GptT8W){aT z^~i-iwk|xow{@zs`yK`r2i@X#B4#$Mf9O3=MOD%KUi#PO_v7CArSp{X$G@`q{Wc89 z?{8P<*YXIP-)8_)9#oSDQS#r_&e1tIyJJSdwM?#tWMPQH5 zn;!_s=ZYG9HCB5)eDg|;BpZ(Il?_LGD781w=0dcT-vx0`;ND8-RFvSnL9Db(VSL7Iwb~+cccs!91WTGaoA&Ep_?FVs?dV z?t6X7S6spDYv*ec+%mX;#k|iLSNdYe)NfuwGvz_6EUTN>*ora!l&a_rb8gY zto|V9uQaO}u9cXjf>jWP+5mj~uWbO@TtBbJ3(j{NX)?&M98>!4&(9%ZuI1hD8ENwW zN{da#{dJH@6Z4B+Z#Ny3JK^8_I4Pi(8HRj<0A#i3KjMvQaWO7#%AW~VQgD3zOw)nz zdcoe-LG&{lm~ngDbPz09K6L&hP%9KafB<~SQ0pvXMT>jtD5;sZRpWakUjMD`34^_X zXy>U>$-A9hqxcOezya-9x8RzdDTB&zR&m`wlc@CDPPAJ(#BV(W+S!S_F^%}Efg6~> z)n{-QQX3%HW^fAX@d-=v;$hfak{5d&WS|ivVlZit4n`Q^d*;RRC=N9GggJq|jGCs6 zv+s>Vs<9eWH@f%E3#jRjjGAgtd#gFHW(%mH4yN1TUXH^7q-1zq6t{=P2dG5Fby5xu z_LV&D5j+$h0a_&VgH+T>bS=~UAbL1`Cckwry48xc;zPG~;QYoR_gUt|n+!=>bl|k! zg?mcT8SeNS?3^FuSP*}=&C!|RXcAH9Z=WtO-t3!SlGqaReNBY(c1dEZ>W1r`w;yVS zl0=a7e)69^LXNtl;c#9u8U=!TDWGPRj|ZtBy5AY3oUJhkw8+O-9ES66a4!OY91pg_ zOjFk4494;0C%`w}ivvA z%+592xxIF-0qrk~zZHXWi9*By)Sk2?N{7&zL@ykCn-eM2;yfNe-TUDxT9Vkn?QxzM zBwyIyNKI}+Uz-v`-C-Rda5Es3#LBDv0@$$2=UsMYCpzb)u2l2ox1t|t!C7xxv^zG< z;x-QE4xx=;Ww@1ZA7Z#47NCKc9j|*jXXUD4qj1y0uwl)LRtIAIzY>ck7=k_`_F1*S<9aMCzVI5*({vbSvN4UN31Vux2Tzv?MmrLu8#YX zJjH9kmXU});2A$0$0*zivT#uK%YV5YDoRo)LnsAH$f=d9W;V&5N1$TL#tj1I*Z{j> zatHD#a6qw>#M;eoJ%HCq3)c?N^jgwLL)F_tg4cTEaKZH>Vn<*T@bzrtH|~X6)QRnj z183hZdo@!Cwj64`Uh)U>mrUyNoG-RC!m_*=yK-nQL<0pD4Iq zsHECyaYDmLpVELT+Wk7-^no!nigrVv5I5@__~4?4!nN0=}Ic@bvzyS7xuA`RY{uP1zqk_$cWGpI)6^hw7Q-)3hswE!EK8Y z+{*?C5)27V776=mJ--#kh5JYGTffARfkSvfk#_Dh4n9~;&>q;i&)8DB0`>>R3QkIO z;HiB1CGZWOz@6{0yQe6nVBB3pY4;FH3qGaCF?zt~56aBCL>pOHa6Dy>ATY8z16)}g zpTG-zLJXfUm`}KlPq+bR)B>IuDc#Tu?4$kp{x;gHg*N$%yp`C4iu1DDEq7v&7z%LbcH)pV2i`=m=959g+VI4)#Uf)pLZT&orPyX)WMIb7biI+!lngS z7GNJ?zp0mBxV@gz3HXKD@H7?s$md-YY@li>p%iLZOOF;)D1h5P^K%ZN07kn6Oci+EK+%xW0(CWk+3rU)}1s@yL)= zh{O>dR9#~O;fztpZb=wGw+O=7qx9K3D7&vcr#H*7EFDG~P>-*DQE!$H^jVsUhm-R= zFVb!S>-we+tqXzB;6&az8Uys0XO27&daDE}fURv0 zMAvO#nsrbozBS!l^g2Y5Q`6mxUnh3B7v4JULy38g81D0q44eQo7NK!@#yEn;4)qBm1AdZUaW}iVbi@Sq(GB@r z;1K@y%YOvR{9{pzd^OO3KN;Z${CRvebL>Y_d^N-DkaiH>0X{ax2sS$ftQSzP!PlN) zBnACw+?dWVg3^f}gH5u-6w{`IS28wBPbk3fWS786S_HGE2P^N9KL)=gwml~}+temc z33}yS!499072NF;rhrH83SN;**C35{nhW#@V3>mDgZ#z-ND?fhaBmhpcd+oGEV-*} zwW(G9ZLpf9y+LZvBOD2$6mKT3b_(o@No<8XsA#Yz9DTzlloTGJlb%=pGx5@r!b8MI zcMuj*3y=2Z&4TRn zt!jgg09s!VUPtW;@b;2%GIVdlNR>k@g1Atd%xfD{Z#p7(cpp zvjtCr_(=wAJpT2~b_$q~9PGqaOZQvZW;=(@3D#iH=gkPh1Xx<&$0$*1X#qChf}`-{ z-}t$<*)}kStZ>!B+SqYb_VT~bM51KVuz9g*gsAG@8p0ltBlLaF#v2-_4NLSkp( z!K^Smv}J1f*dA98%NdCbUdp|Ucv4|K)m-M6&Dfg})0Q8_ylizId;>`yZK>7~zhtaC zK%49Vsym<}5lYHDC;~-W7zxE17?J~TASs|BsSP8+2Ua@Ex&X#ibupSW(p`~C3ODEHmn2@0uiFnswBiscb^hU*7M*215C!mXHue4=?r&lo-7Onz6ix3@q+jiwG}|c>ZSa+TTddEkh=deJ$6x}aPV6MB^w_5eBNZb=Z_Y^ z3Jvw4Pe=~Juv{|wg4zY)I#q$lnua3D;r51dv-WCwN)%=dq6&h|@}E(AI0U1(3PXpf zn@Bwbm>R;+LHvWLx1tp=H3Y4ucOLrv=M_f)L&*7MT+!3ar_zoa>d>CrxE(gG(Z<1Z zb~l&?K22jI735`RFXXq!Mv=k`+LwAvb~OabrS(>5{Xok=hJ6hu-ldJmK_009qE^WI z8eE+U%D|d!u7UlbIe3NN4d!|4qk7>7s?}f@TmYKBhIS zi&}5F2}kJgR%Jo`8;fAvgJhEk)#2nQ49{fPCAX!xx112RsVpOdmH`Dfs zos$%kJ-Gt8Ks=IE$ow7|I|)$1dPXmMolJ%8QcMBK96rf z6%6BkNbX&rUD6ZQvBtqDaxsIeC%XtNA``GeU>6z8wz{xggx&vw2)0}PByfdUMNVYc z_xwW%Oh$DItRi6L{}ZbSw+G%l>0}hyNJf#>YKLGH0f(h7FayY!NYWcXu-a4uNB~^S zXeJ(Yn(;0Qg}*%lYsdwafnmil!w2KBUtxyCi9@UrIl)%wdm{Lt2_BtbJlNidks7g! z(9T%a$P6C`ad_!9eBgS$ki66Co{g_^GP}p`Fqa{_hvP7e9L(-P=8k4Ea>U=7;b^6R z@fM48oSJi$;t9wKY8Bs#@lM!XaEo%L+b{sbEF8kp> z&IMx!waJ8amIn67zA2)z z+m{8EjsBM2k!;kH93V$B#tw;^LB@C`fqwxs9oKGG!*`YpA9XN%+}z8-(ijJe!b!q! z2G4;>BORlHUp%5yITm-SX43dMahcQ}Kl^Hl<(mMhNoQq z8BQhr`wNJ!`x7j%e&9Uf+FXgM-45^n#D8dcTaO+bcL_;zNkyD)-bLdooU% zc2kpuzHrpmjSC(4uxfz5d~0)Cp2OD+U<;b&Hr{J@bl_`lc1Iid9&MFyriV2eB`=Zf z2@1Z|rZ&(5)DpP2Lf>AYlPt~>$kjn$Tx03LZI~aCT!7Eo+zcPS;qdK?$6%CqW4s~s zoNq~G=afd0x~m~|hrXWx+j8DygZy!x>cS8o1?vJ?B5kbGij2Sk-1jgZ2`9p$aeyTb^a&Y(FDS94H*ttvHpp}G#ttH12}(n|qLB+( z&-#QpMM2VqT-20$IC!+kLxLB5(}vkkc2LTf^iX`&Sw{MtFPTd4gm?oeUAV#b73s9X zw0}%pcutNEekl(FOB@LcLUa&^!Xc#h8-|mFFf!PVp#Z29n<*Z<=&{dE-@Q7i9Yc<> zW5_{v4EZceel}PGmRfj>9YA^@tgG>(Ed45}`KkHoj*;Ypo~9O(@|FcIl>QDkr0q!g0c;9C{voXL}2?#5E`bE!mA7$ zzBTzOEG3<+39OcL3|lwDMpdXlh7Cp?s5{QoOLAWZwvj1AN-5<6suVzBqJ72kbQ2 zk6xniV|YHC{lfr-pQLB@RCtg)J#<>P9g07EplBpm3S^K=AbsspP(>qgoxXH|jGt)w zI|Vw1R?_47OQ0KxP=NGfUx4moI=#V)Rf_#{0?|g z6R_jE)f2uCVQwc5cev2@4r3&${Ac4X+TenMyNaY=VgAThz$8fF;eG_ogbh|EYr{vh z7_QLcJQfL9kW_~{ps?5)#HPS>Wsnn&Y5EKK^}MmrSBxQj1vp0RAhHNF18N*ZP*ccF z3f6;GCePt=N?3LAMW6`+-7fics4ZTHy5eJEGl5SK_TcZNuO^?IsGy`EJHa?+7qF6D zfT=F>?XBry>j1udn?XvA($Q`7xIm9H8Z!ZR2wMmcR;{eeM8G%&@PTo_P5^sxunP){ z3wWd3uo|2pg~bR6Bv(S$K?8eBc;lgqKOJ{Sl7( zqyh!kK)v7^s(Msa$WvZY8jy<@g!Mvb7di_;(G(+SXJHWdU^;D3BMb-l1gw!-DMAfZ zd}m=d;Rz{Uj}jiR2&zlfw~_c~wGhDFKl0CK8$?Ov0f-BmE;7*1Qr3e@7 zw1mAqhpe}7Qln(T)yZ&$oz-AqMbbZYSUa5a39^BafH}a1iKCn>pU#edn^t@LU*-1f z*>B@i3>Wvm8s)>;CI&aX1M$V8!Y%I@ajP^+Fkj|3rV&%XYQR3(gxzr`4t({vO$4g} zcMKjGt^CG6^~WPNY5bur%paZtf54I}d*ABHfOdH{TW%SyzpLx4(e5_A%@$MY5iGMN zEGeufI-kK#z>ZQLT90(jVJcNrrG3J?q5~~Q31Y5hR zK89zxaCxe66kky~5NJq;TyHes=i*kIZq?jzRSp|SI)u}@=^NU33-0?!d7MIWpo{`^ z+hYyh;JXv}?Xyj{!f=4O{~T)K*`{xhJvI((b42)jYsGzG`W*w(FBhU5{6g{l+ByNB zz~SgHjAMUL<(XjZ--TN=cxRr;B-MgANQ{Hh7+e&&z$U;n(iuMm^8}p#EpS=^=M3n%Y*;iyVy_`qmCV3p;l=_EI|TwvM$irwupA6P{KN4O z*4SFo5;L(3-Zg8GN~6{RKHC%Kvm|P?FlnZd#Ee3+$HJW@QT~G}fwe%WQTZ=s13apl zq!)Zt7Fo#%#txN*jSPXpgcvAEG$0d);7X##p~WU1^=DNaI~3A9lOqdBUz%)3#xj^} znZyIFv1S8im|&wt9ETNr=eh)kb;9#0iv3q#zI>RSh|{np)EX*T&fr-cNU)4wirrk&pE25pj9v29+0bZ zf|`GePl%ay87y@I`MBW6h_+i2Ir+!8<2KGi@6qQ)Up^~$%#5Bq9yA$vB-%pni`n(a zf#bol!3_iOVMJ`dCGWxYN_g2PD>H}0Q66Xs6Y#bq2f^=2=Go;tv1o7$Wp4z+VQ20_ zCgqV~pVPY~yU~&i|5XhivdRSg0=jY*RZ6}M_R_0E!=L48H&^CJ0LU(>+KVvv}2C@ zuomNo>RF^3atk37_HMLTa6JqbcA2pmID$^L$z7J6_qP}nbllsHok7=^!Y<1Y%E1D*Uw?8P>Cgtoq&2+*a&5zL7qfEW}FS$Q64jXz?(P@fci42eE) z0Vj=5;gx|S$JHX4JTqV=ZC_&C>pL;WUS|M?(cur07aQv-&UQ0=7%g2TIR&#qu9K|(q!WTjtjL4Jfj73!bWb8!V7n zK{O0^YLR?92oDCpF~MWVe>=zk(V5E zxHS2xdhxQGepa-CeZ<6Ve3*VxoGxB^+wdWdI`P?0)k~aRqIq9mmoN7fpBbZ~@!^}g ze1#uP#GOmrw&XX^JxsbkC%COWlC>&B`t*#Q_z?%eYbmH_*|*lwM@{fi;}r?GfQc)5 zIKXutm}bH?%Y@JA5@Zd0e|#~@9naBE`ONviYW}$B2ims~EYcyV!&O>kap!O^TckN3 zu-lrrJtoaaG;RwEYhkt^Gon)T&kR`eZ++6VIuDKv3T#j&D-(SwhZ&9`g?u& zv9)V{*ANJeKp2Zt8aAk(j>ZO6u5M6&{U+O>hL6Gq*XswQZrqbd>$z_M>D7>L_57VcKwii(bW#lw&US=#g!2IA@9WS_3OU;J-b*(ZOaB$R83Kf$n_ahoFg{fYW6cOgH+Ovt zgpmbgNd6NgpHh?Os>w(2QYIl;S%ybn{$wB)_AMfOvb_&tVJwAy+Y;0j!BrRqfjB~x zam{t^e~&nh|L+jTnF!+Wew8>LRQc@Jp^vrn>Ff|+*^AMOq>oCV{`cu4{-4uF%QJn_ zhq-UQGFYXLtV9sNLP)KUMRUKw=wmxx%0wVj3i0SepZ=fu?`L1xhOblg&OZk4|Ly$4 z64J-M@{i|V#xC(b{TLk0e}X#ycqtQsOew^p@A-ew*x%yE20`+~|L=?cBP^%?g`I(% z3PL->2MFtS2HU2uv+k4V^P{14e*}Z}zc=r<{PT5xc;&w`@4UM1ts^n-BVmNXygxCV z&3iju$`l||R^ic?yrBIe8rJ`R=08%NRJKRXKL+pr?fe)28}tA7TQUDR+Wg0>^N*J@ z1;~_Dc>JUJPr|QZ_ObJ;=Zjjb{%g+{KP7(P*he3`$B{z)f>~|o11@#{%10H?_qFJ2 z!yi?=_%+WFoj?=*>e&f+1wy(#gTJ^f?q63;SK#y;U(m_$enZMv( z{M~@lE%X}u^O|rR9?=HX3Df8xd`!l#yg-dm5CaIbIWM+ff_ zxMycggr1@WiirJq_dv;Tm24XyB~q!qboq8~A&|huoy^97!DEip6Nt7@0n!j1r&3U7 zeB!Gqwl@~iq5DhqAdhcc**)n5fHr|S&*PGs{7Dc$d`CVk-u&P#l#%_G-Dp1n*$>*m z3XLEsch=k@pTiBRbHSkmX@H6V+HgAK)yMo9~PFqUIPgt(8o2gK@cSGN-J4*F_A2-Enx|7kN}y`T&-3sY4K`z*&nj( z&>)_Pi?IzN%}YW;s7Oge3VA3gZOx}NZwVI93FcLR+x&ngda3C%KP!lN_$7wz_soy0 zyQ96b!A|?f^L(^AbI+VPGjrz5nV)wq=@<34T%KHb*iFX(pQe+yeO9)6V{uD{t1LQsrP2;9Yi^b3OMu{B0LXD{Dl*ebbJ5@Y-2cd_!aWs2g0N6yZ?d?9lrNFuz49DB*o{>#z5YI z&$vGIFHYJ3$GODymE@Uj*>&W;_kP#9`z!Po;;yH0F(8iqwTGT=duqCf?rFyL)*w3o zGJYhmXy=_^O}Ht3aWQdmAydgxlifXyS%9sSlot#vxH@n|NJhtLD$|gMW_l~f{BX-*B+n5!@hsm z|D&~KQ~TFdm1Xz;XsCD!q_l=#@e+vO8QS-3JI?QXaRgHEJ8((Ok2mfu-qV34D*nBd znP;%O{0_bcJ@c`j%HcEvewIFh4G_L_Uw2=$6d0;D=L}7_z%## z13PKJhV|Drd}oaw?SrRsCG35?<6J2#s%|Tz4~fxM-z4sO%HRp(gtgl+ z5XJLi`b6^Bu?D1R^vx_z#2ODyvRKx zm2{!f`Sh;L2na^b#&r+m+5T~Owi{Sr-@bWj^4Q8}aP#J8j;D(!fAG01X;GHG8iZGL z@m1>slP_CA{oay$R#LM(wspzhl)T%ZX?NiJKi`BX_!2e-pKso_^7R`;`@xlf z&Ba)}#^u#(r$F4g>zB9^7$+@{+;}gvdiDG1DtrsNNnW)2q<`C90vF)^+E*_F312r8%h>zh2c-@M7*xLPhLPS#>bSSA zV<~~P&n`gBrv_Q@3lz+b5bdq2=D`~Pj(iKjG&*YC_4KO;kO&_uLD})?s{Lh8Ly%|O z83~+`z!?dgk-!-V*h#>lVT_`UemZ{kjy3Q!(L=G*aiykxu&yJL8VaUEgL*KnM`C^9 z5nWrOYx=Uay4D|wMkBES9pTByh~A_k)@dCgWR3K#TcfRu1XD6fC>ai>!+kQeS+5KZ zk6fl#>6IND+b>(ARadWE3GPFgs1X@TREI)ZQ`J?Wcr-H<)7P%m*HvAWjBnMN6&xoZ zG(w3Ek%*3Q+Y*%=hqi!gEZsnC*Hzz=2quF=sp{?FWZX!^5kN6LED|y$Bf)F*Cu(c* z^Rhy3u%!kkA5nAuRaf;!Mtso%NXsQ-MTA|nbIBlg2YZhUb&GME(mS{2wrZQdOR9f#3Mxx=qH9BfWk7v?} zOj_@Yhf{hio~G&~4H+KMQBBeCP&k$jrX%rKHH)sX_=#XDrNn2^z``!#a6HmSyixya z)^sqqDG^Qv)A8h*HO;O5=8Z;2Q`fa^&Hm2HP&|fAf*7iJ`4#$nk2u%H-G|^$n&R6U zjtmT@KYUVLPxZ@5VP~~>wni`4XH|HmUZdsZ3q>-h0Z4TP&zUkHd3KrG(4Lf zqT*F~IWOY;e1g%$U~rz;nl>CsrZd5)p2sJ@9-6N2Ub9AKF&GMkqr85MbTSxACE}=Z zqd%BR$A`m7gO}Yt$-H!^l^Mx!zX9%4vR>E?o2aLfPpVgzLmS?tDW`?M5sMq~{{B?h zqCiTG&1s`(qLs1CkdYpQq3nYK)m(NNZ_Z4CNnV$m4Wv*lkRcJVjs8q5WYmn*7_~K- zSTY`sB1cT9j16hbDe|k7j80Kr4Xi%$vN})YoO=C;^h~jwl&(?*%&9N3srePwkXK)) z(4JULQcKE{i>yM^+~sN1>8}Top3zGN-?mbIf9b}_D*R`qdxBD*vYjZP_eL)>gYy^r zRPECV(lt(Jd+?ps_Fx)si*Xg*g|#E3fuAylW6g=8TZd+r@sHWGEyFc3zZUDOY0n)z zxc2M!{ruC7um$wP@(qIkaC&z1HtxHoo@dJI>&P3#_<6797!X zmH7ZqaY^ZdQN?KvXOTqN@yFnU@poKM^oKVgGq`I->*>U@F#My}<#PJ~#{j1Q$8X5x zmgDZUCjcJ@d;!pdZCkB7mwO0>(*Z}Yq}B*H2z(DcBVJ&0E_WPo3a|lm6Q9lH(twEp zEZzY>N|1zumAEXR8}JaI4tOqvGm2HYC<$pxE_Wwj0`MWgaX=lyoB;FzW&v9PrvPsP zoCZt)&H&y4sNs^suK~IN9|hC_vw$AJLx7EdM*stW#{qi)-KeAl;BvrGzy`oEz;?iK zz#hN}z%*bM@J_%f!21EG0Ve=w0QUoGL%H0~0NsFZ0P28L`5wR)z(&AszyRPrg4jAc z44BxS%T-{}I*X%UJ%AIpLjHg~I12a%;0$2pS&09IT&@dn8t}`2p4*Tv;3(iU;1u9- zKpiKAD;9z;pbsz$7yxwN4mkt%0Nx2W0r(K0hNH|`zyRROfNs3q(T)1o0i%FZyCFBg zX~0qNp8=cz)b^l!ivaI|o&dVBM%)89{-@9{zzM*I0jF`R#|&Wh>!?rcu8rQ0{1AKq z`~U;vs4u{YZzFxc>|f?`i?IEs{~(vU5^&;&$OkpM9LVJk0*?Lz$_qFHSc&r? zJ^0La7vL1&DB$SxkQ1Q$5abQm_)jPYAl}c*RbaI^@eAk^!Iz=O1b>PA5}Zao0*?JU zmn(z5PyY+}1G--UJ?a1dLHPi+*C2O-aM*ib&awSk%`xJ3T(D?C8SP~d{zCkwyK=cp zh@xcyIBVK+{CxNw*aVr7v2eN9yDF|;QoglpRQvQNR$pGb{9<5P{F?x?pMkCtLO9}^ z#%~|tDZX$-vj@M?P%a16Vh+pe+BkmgXy>0NOqp zZ7aV5+ z?j~2+#`0UGm&GGq`w?%XKcD~4fJ-a5H-JkhI1P0WRdDA5H>lt$f$LLnR|40g-~zzi zq~Q91>sD|hz;!9OJ-`JN+&FOU3hr^>S{2*@;93;i%fK}%xEbKCRB*+x^$iNH0yv+7 zy8<|mf@=is3I*2%T&02=1a7&4yA?QH!BPKqp@Mq|xC#r7^ppC(;$SX!1Ioj7rFn5_ zhebaxa*}?wkbaVmP=9tfoUbF)_e}#Qb&&Mfjs9&B^pXv4pvRjRx|otv*cAv%AZ!SJ z((!%~M*Le5wjW`iqA<8h{OB7sW8nV*(&2f$Cch2+?;>OoGDYy6pzXtc*#%^uneKzO zh04<@8r+Y|%mxRN7W|$>oInzN9`u^p_dR4A?k-yA()Sjxb5-ssX>xgXmNvN>&|G~E zmj`VgVHP8SXaW3+(J;GHx!hauYn#YUz@+UV8zbz*2A9YA6_|-;mu{Mg*79}*(^8H` z|MW8lpQEKQ^t|!~G2|RCgRRUqMB%3vcc+qq&x6_QrY!-{h zY*stsQOeeco64EI4FU(R%g#u(k`V-bZp(ghlCv}ABiFJTZXW6^~!gf)QPs#k-+^%bBeUUwkuYc^ry z2s>yK_9VhI+;^eofq2q*#%B|D6k&rlVQzFj_uGVBfv_o?uvUZ>Z!MgcK7=(C2qPUH zMOYMJk3*NWwnYoIc6X_BSi5>j(QtREdx>*+#O-npC)|sj!-MWc&f%UWcad{AQ1l5j z?}b_i+D}_KwVwmvlWx!D&WB&)k%^8mWAQi##*1+U*>S{~z}RgO#et*zxX~UT1&+SR z@?rAx8;8tKE6Yze_`JXIuXENJU~fy6~$O!S99 ze*pUDw4;B7(YKOqB>H05oQLnooRNBWfL2F&=~p1E{t1Dy9wKF&CiWRR9?T#KaHiowl|mi0og&;me=R^ z1>O1OJqr3tjD6p>qaT^2uLR5G7+W91sRw$MvY`C;IhY+MS<-m@{x9Wne?-QD5nS5D+|VbXs|Z|zrojQl_1nBoXk;|QJ&A&B2 zS26jt5r2~XIPuTsa=q|VKeET!K{g#0cSrFpt_t+RZEU1Tv7{I|1=k)xhBmG+bBPsG$)BHE-hNb8XtupMEJxDxm*>6 zf0n}g7MD8TH0$FH%HNA8CttF27or0nUoYk~cqT`3_W}R#Kjv~z*yZ~cW-qTM9j1Jd zb8HW^vl-@sZ+sC`&kN){vmIf}@vFexb|LIt@jZB$?FUf(yF5rZKSpla6oEDJh4(sIyxlhVN;wxX{zVU7jCj^x8nc2wkjWnQ# zQ{^)*!pB6sDcpZXZcNZm3+UIpd z{tuE5O*zTYeh4}F+^syvmFH3Ad9U((zw-Q-;&~BIhA&3IzgRpUe1dz<7tb?CxaR`# zTp=J{^k%M1w0k*EEfe^6Kgh+&J0a4Ob3WSdCPz!p!c{(vaB%=a;ok+VcRSzg?G!knc0U5=Lz z^~6vgi01=BusQK8^Eo4)Zx-SIPdslCd`~wRe4{W&kIwOp0tN)^5ilX(sDNVvjte*; zU{=5>0jC9=5l~yEG3ORg7tkYMqksVcdjw1fI4aJ9 z3+NG$ho1cZvKdaB5Ozw&+|b;-Mz6fCHxo-|bZ>2SO|_@0rZ&T#aiEK1*s70{MG^WV zdQ+*DlADAf?n;Gd(@A^T~`(Q*GDA%JP0<8yBE(}5lq zySR7(YkiKTjuozR%8M?#;G*l!Tk3Elsp)z*Tq`fXs=UYr#s+>RXbF@>hr{tZ*J1(=dSZx+oD@P9!Rffjxx}ec z2>Vy!(9Uwo+!wv5IZBF(7G7OaQ3_A#=5tE_Yl-7Vrvu!EMnb8imWibz1F>))UbskW zI55o5!{Q9tptj+f>-1*tDj{>dYA70y4OI2TGriGpRf7kJ)L<&T1u&h7h1U**W8q{Z zgvVeqG`M!8p}q<)1XK;kzfCoImHhRG`l}Mjc$yBE$75^zGtsDC)!&l%+ zN`6S3-ihFUsbpv+PkANIxb{Z|Rz_+Y>Q{!6X*vr_BB>5(D_JI2B70Td+N#XetG2}I z)!O2ix)PI zF8j0MG8a%9DXomOKvup~UCv7RPpVX`y^Xgx^Bvh~2SWA;L#vBmN}&)v4R4*;mF*GZ zTiMUuEb!YD{6jW;o)q-jKk#1U%*EOCsG% zSuCvS!dw)}|0)~&)xZ-UIc}Hz-)%PZdj-8J{BPLM|An9rh`JTo)E*Z2Q3d~Z0-sgz z&jU|#2nab)y^{N9hCfe};|Y$13{bCBE=|l!B;F0YL;JWU$0riM68J)LYqG&_vcZ4O z20vtj|6|}OpQEijCvS-i-DN}nO&dIomxxcIhcoctncS4XkB9g(+uMZKuD;#}y#u45 zD}Tl5t^W(jr@WZxKVgoiB>ggmKVMVF6PE+8p&e%rN=gpb0)H{`uZwYWi3t8v;LirW zr-vhCKXy#uJq;ZHDM7Cz;!@~IkI)lc;Ol@VKH23QA@zI>@RV+k;3L!hs=&*6qr?}X z3nV^r{>je3fMpBAi^MEoGw@5D7ihW|2gsm{F;FEw`U;MbY8nQ;E{QDgcE;bW$@yHf z2zv#1%D-EYn->GsWf%oFiuq%iXfW3Tf3|a}CdcbC2|8y=eB}H{miKRoUJ_a0^MYQ^ z|0I1SEIO5Ie1a3RJu!HBA<5zeoVnvVI0r+_!t|Fo%4b>2K+LF$n^_3kJ-cMIj!~o_XHo$ zBOJ^2GU2@t_3qX%D^m2~GKRlUYgF)cLQmX6Fl>Jxe0DIrHAU|KmY|=wpFj5slAi!i zau`$e{8MO9h3a=b!-M{tg0JAS)rS5a8~nEgA30BEXC=Vvs11Du=2fg+T*DKTdY)l; zkqHaX_$^m2Vi=I0sUbxCA_bl+W} zdVI$Q?{;%LDfP`Q?87C%lb%c{da@JvLj1pGgFj${|24zI?p@6@zf+Xq81R(txRUN? zP>F@ojWfJkYuv&auro67J}vR0{U!zff3@NBfuNV`TT;$9AW_PXT-TEJ?fU|sRq}Zr z2988Oqu7&r;7On5x|-zw0K<#SS-{@|Pj*|b`&EdrUkd)?3jd4HkrDq1A)M zP8VWhtR=P|*@eF5gvY@!7dGcy)BWgrEKb$FuWk@D4IO$65b>h44;vR=O37vv$D4bj)#6 zg1<)!4D%CmtFh2o$c}{>9{L&KbZ-kg{6ic1pW5JmEBNT5zn1;}GBh+Q@6WiH@2mR-Pr+Sh57ShgqQ{d%3 zjKn)H;`vM@S#laXmk4hF_y=KU+%p`H;U9Ci3wpVqff{1&u)xc8e|8oXAuT%R@7~1e z*tuVLcK}cMOep&EJfr8d*8d|b3d^BtSz-KtDJRu|Qi%LZUX6?%cDS2k*cn@Rw*s$A zA`841c$!DZ^+Xx;fZ)@(ieqjQ0Y7JWNi_@K0lrYZoO>~skK9L+?Q0wGL@)PMWP9A~ z6!iCV-n}A;ZwPviBF}e$e-PuG?zcIitjEeva=Rh-z0i`GTMs;|mj?bU_2gmTnLQEq zmc3&Hl1YZ=wATN>67+JPN6r_TF%M*Xg#4xc{2B19yh^$ysQ0DR09&(T3y+$)JlI7Tb}bS;awP!2S7g}B^t!!v6$`g*3G5P*;W{EV z)?eGy;Wzx(v>1kFhzi`aetoCE%jjxa*X{?lW%D&n9c|4Z#>R)?ZxwM`TduV6q_AbTBo7Z%hP6aX_!tlA6QE`vL7 z=pr9>VaSI{%i9nO(j`Z-IFw3327S@M@K9rOcX!jeHp5$8TfG|U($&!{HLI1^Yg+W(&C23nS5^dbj-nOpUsKKUi-Fyjq2f{-`#vDnKY2n+O_WhZK zZLJTZgLEVj-6pI#@07ju;Y7sn@ott3e2p=J24wzmliyiu@IM&caQGxMUQ?HU(U|2T zFo=d48!j9ZYz-qc5~L~!MsdBHT>?tIp}(&uHgnpnO|!Dx2xA>%6sjlFe9Y5sMXxaT zawtnWerDTE#Zz>uriN8*e>}N0nCvq`nZBSg7>xGM%~8Zsld9eFn&UEJFgNaeQB9?u zRpWvTx;pL3;MZjLIFw0;M=)TcTlzjk4W53(X!HfsK?D2a$u#U)Zz?7Iv$bl9Y)-SX ztyO;Imo&%2O0nlFbi?dGN}f-c9ZJpiRBFARE-Ol<$GwT6!7ce_$7fND8c!kXAaZFM z1GCIGxH5j@j6G?_GohPJeX(e+_6}p%k(Fb`w#AXDuwqdV#75q5BkuuWT z5@FcK{wW4LGx48>#II?+hps$aeF zWZJ3=(#qm2Pam7$m?gxiGE``1+84s*qIRu8>3Mv11E@kbnsLilNJuxvf3(vOiA;s=Z!P#LjqKe>9UC%nw9}w}vKMnC#t>31`BA;ZDQ4 zqM=$#pS?++qqPN!k|*A+L5l)$@yhm6Z}ylUy64Cn&Rg4Hji;I2=G&1sWAHPQ&Uh{E zi}i<8Dbu(ICcA)Mo~AjRsxZEj9)%jiHFFkLtQJS%Tvumn+U8VT-X)3!YSCiXlvG-3 zxIFFUR%gF-G&dg>+gd0amWG_fy{^HO+EKJgeJW6a)KQqhN-?~v=8PEUrcZ;#)vWE1 zBIk`jPks&0nu)ocnJP{tn=M|FwlOd=V&K-{RGfzO>21ca$7<1xFpblYj%8-1@L_%C z?9fxm2fHK>c48GRZ{^K)t*uLyFlG%OMdeuQohxh2xQ?C6G_DP&GDDcTG&f`HgEcI( z3gsoN0>Mb$tVXF^s;1M}BJGPYw?(lz7+W))Ta@{-Ok*dqFNMcx^BR%-0h+oJ zZ_y^T7HU2DwZK*dwnox};>P5;C9pD6WR3!S_+fco&q)i#!P%7!>N{yw&K#?V(U^6N z(17J&Hbn4K_wMyD^<F_F#1%LixmW*4C`qHkjYLir3qZ#xxxy^(n!<~!Sj<X`UlqtXbxv<)m z6xcQ>+?UT|N>4B@{Jr$(^Vkd4=N7u1-g>h?A8ncTuBNuoY>ax5VShsepUMcvLg9I; z4IRonL%zAAfO%%uc7jG`PRHtlS>%s$I>0IcrKhT`3H4%C%a}7`oOh#wFYd)tY0QO( zV325^q%8!jH#8W;267T}nH1I({MQ0k3+ojL8G^bE+nyf7-&t=2Hey3Tj6e7)B=)o{ zxk)u-8d7E`hUc6?LPQsR`|f$vy4-{V=017019a)|bCIrijMt%ij~h2!Q%0 z>qbZ-i7{M1E!UducK*dwD4aS}FEGX_xXeCJk?H#u^N&xIM4*)F5Jf5csamtV!yhNx zOuJ?WhDvtkn`fWC{c3ML9c>6E*>)^mC=90p8-tx51Ct-Kw_~e6v$cJKmGjfJFJLx* zwRa~AW>gSZ=N^v>IZ~vLdINx;YCyMnfurWt0M~GJstL+BY=%qw!$c$F>l${DQ$< zNGvYX(FZf%*s#&6Q`?5{Dk)$($st~TdeXo+DbT z!|_UXSAvY+P{cze^jL}Em*+kux*i@nb1vi4ecj|F6tC-XB>u~Dml8q8)2;F4`@9nB zl2AZ3{tdwYU)-r6)1MIOODOYSD&nj84+4)PzGnXAy%ZAC_btd#9!YjGegz@n=-4+U zDbtttTu3PIwIEt@s{C&Oj_!<*@#Q@c65hyS&EA_K=_LFD!s!kQ`7H0DkWeN_`KNr! z{0p|)UU;a8W_-L~MNZz2A@9Xd3ATgv;TDRI?t)~Psf+7@pYY(n@|=^aFOSeCH#&Ozwu&DD50bi(Q)#ZI|(Q7XiYyL z;!Eg0LHbW9@#XvJ65cG{hY`t|>C5N;p~RQ(g-R%@QDI{b z%Xkv~DMTim%)fk}QNEWJP||Od5e4*!^gRk*#=k_tOLh{j7V#A_ZN139%2&pd^duh@ zM`6PvepnI;sM2=X#D7V|Z$Cl&UYq!riuW`^Cx{;t@zrwPCExQ -#include -#include -#include -#include - -#include "matmul_mkl.h" -#include "../avx/matmul_avx.h" - -using namespace std; - -// Function to fill matrices with random int8 values -void fill_matrix_int8(int8_t *data, int rows, int cols) { - for (int i = 0; i < rows * cols; ++i) { - // data[i] = static_cast(rand() % 128); // Random values from -128 to 127 - // data[i] = static_cast(rand() % 256 - 128); // Random values from -128 to 127 - - // std::random_device rd; - // std::mt19937 gen(rd()); - // std::normal_distribution<> d(0, 6); // mean 0, standard deviation 64 - - for (int i = 0; i < rows * cols; ++i) { - // int value = std::round(d(gen)); - int value = rand() % 7 - 3; - data[i] = static_cast(std::max(-128, std::min(127, value))); - } - } -} - -// Function to compare two int32 matrices -bool compare_matrices(const int8_t *mat1, const int8_t *mat2, int rows, int cols) { - for (int i = 0; i < rows * cols; ++i) { - if (mat1[i] != mat2[i]) { - std::cout << "Mismatch at index " << i << ": " << mat1[i] << " != " << mat2[i] << std::endl; - return false; - } - } - return true; -} - -void test_mat_mul_int8() { - // Matrix dimensions - int M = 64; // Rows in A and C - int N = 64; // Columns in B and C - int K = 64; // Columns in A and rows in B - - // Allocate matrices - int8_t *A_data = new int8_t[M * K]; - int8_t *B_data = new int8_t[N * K]; - int8_t *C_avx_int8 = new int8_t[M * N]; - int8_t *C_mkl_int8 = new int8_t[M * N]; - - // Fill matrices with random data - srand(static_cast(time(0))); - std::cout<<"filling matrix A"<(A_data[i * K + j]) << " "; - } - std::cout << std::endl; - } - - std::cout << "Matrix B:" << std::endl; - for (int i = 0; i < K; ++i) { - for (int j = 0; j < N; ++j) { - std::cout << static_cast(B_data[i * N + j]) << " "; - } - std::cout << std::endl; - } - - // Create bias matrix - int8_t bias_data[N] = {0}; - fill_matrix_int8(bias_data, 1, N); - matrix Bias = { - .row = 1, - .column = N, - .int8_data_ptr = bias_data, - .qparams = { - .scale = 1.0f, - .zero_point = 0, - .q_min = -128, - .q_max = 127 - } - }; - std::cout<<"Bias matrix created"<(bias_data[i])<<" "; - } - std::cout< Date: Thu, 5 Dec 2024 16:16:05 -0500 Subject: [PATCH 3/3] tested and fixed neon kernels --- kernels/matmul_factory.cc | 4 ++-- kernels/neon/matmul_neon.h | 5 +++++ 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/kernels/matmul_factory.cc b/kernels/matmul_factory.cc index 31c3b5ba..228456ca 100644 --- a/kernels/matmul_factory.cc +++ b/kernels/matmul_factory.cc @@ -11,7 +11,7 @@ namespace matmul { MatmulOperator& CreateMatmulOperatorMKL(); MatmulOperator& CreateMatmulOperatorAVX(); MatmulOperator& CreateMatmulOperatorCUDA(); -MatmulOperator& CreateMatmulOperatorNEON(); +MatmulOperator& CreateMatmulOperatorNeon(); MatmulOperator& CreateMatmulOperatorRef(); MatmulOperator& CreateMatmulOperator() { @@ -20,7 +20,7 @@ MatmulOperator& CreateMatmulOperator() { #elif defined(QM_MKL) return CreateMatmulOperatorMKL(); #elif defined(QM_ARM) - return CreateMatmulOperatorNEON(); + return CreateMatmulOperatorNeon(); #elif defined(QM_x86) return CreateMatmulOperatorAVX(); // Default to AVX #else diff --git a/kernels/neon/matmul_neon.h b/kernels/neon/matmul_neon.h index d5beddae..4b17330d 100644 --- a/kernels/neon/matmul_neon.h +++ b/kernels/neon/matmul_neon.h @@ -25,6 +25,11 @@ class MatmulOperatorNeon : public MatmulOperator { void mat_mul_accelerator_int4_fast(const struct matmul_params* params) override; void mat_mul_accelerator_int4_fast_no_offset(const struct matmul_params* params) override; + void gemm_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) override; + void gemv_accelerator_int8_int4_fast_no_offset(struct matmul_params* params) override; + void gemm_accelerator_int8_int4_fast_no_offset_v2(struct matmul_params* params) override; + void cblas_gemm_accelerator_no_offset(struct matmul_params* params) override; + void mat_mul_accelerator_untransposed_fastover_column(const struct matmul_params* params) override; }; inline MatmulOperator& CreateMatmulOperatorNeon() {