From 7e6f37d51fdcd4c5130b3261b0de513af3d1d50a Mon Sep 17 00:00:00 2001
From: David Goz <david.goz@inaf.it>
Date: Fri, 2 Aug 2024 10:57:53 +0200
Subject: [PATCH] omp block matrix
---
cuda-omp/omp/7/mat_mult.c | 185 ++++++++++++++++
cuda-omp/omp/7/mat_mult_block.c | 208 ++++++++++++++++++
cuda-omp/omp/miscellaneous/structure.c | 59 ++---
.../not_opt/jacobi_serial_not_opt_len14 | Bin 0 -> 17416 bytes
4 files changed, 427 insertions(+), 25 deletions(-)
create mode 100644 cuda-omp/omp/7/mat_mult.c
create mode 100644 cuda-omp/omp/7/mat_mult_block.c
create mode 100755 jacobi/serial/not_opt/jacobi_serial_not_opt_len14
diff --git a/cuda-omp/omp/7/mat_mult.c b/cuda-omp/omp/7/mat_mult.c
new file mode 100644
index 0000000..8ad7fcd
--- /dev/null
+++ b/cuda-omp/omp/7/mat_mult.c
@@ -0,0 +1,185 @@
+////////////////////////////////////////////////////////////////////////////////////////////////
+// - Naive matrix multiplication algorithm
+// for (size_t i=0 ; i<N ; i++)
+// for (size_t j=0 ; j<N ; j++)
+// for (size_t k=0 ; k<_N ; k++)
+// C[(i * N) + j] += A[(i * N) + k] * B[(k * N) + j];
+//
+////////////////////////////////////////////////////////////////////////////////////////////////
+
+//////////////////////////////////////////////////////////////////////////////////////////////////
+// Author: David Goz
+// mail : david.goz@inaf.it
+// date : 31.07.2024
+// code tested using nvhpc
+//
+// - Compile the code:
+// $ nvc -mp=gpu -gpu=ccnative,debug,lineinfo -target=gpu -Minfo=all -v classwork.c -o classwork_omp
+// - Run the code:
+// $ ./classwork_omp
+//////////////////////////////////////////////////////////////////////////////////////////////////
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <time.h>
+#include <assert.h>
+#include <omp.h>
+#include <string.h>
+
+#define N 512
+#define SIZE (N * N) // matrix size
+typedef double MyData; // do not change
+
+#define LOOP 100
+#define NDEBUG
+
+double wall_time()
+{
+ struct timespec ts;
+ clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &ts);
+ const double ret = (double) (ts.tv_sec) + (double) ts.tv_nsec * 1.0e-9;
+
+ return ret;
+}
+
+void CPU_mat_mult(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ for (size_t i=0 ; i<size ; i++)
+ for (size_t j=0 ; j<size ; j++)
+ for (size_t k=0 ; k<size ; k++)
+ C[(i * size) + j] += (A[(i * size) + k] * B[(k * size) + j]);
+
+ return;
+}
+
+void GPU_mat_mult(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ #pragma omp target
+ {
+ #pragma omp teams distribute num_teams(size)
+ for (size_t i=0 ; i<size ; i++)
+ {
+ #pragma omp parallel for num_threads(size)
+ for (size_t j=0 ; j<size ; j++)
+ {
+ MyData value = (MyData)0;
+ for (size_t k=0 ; k<size ; k++)
+ value += (A[(i * size) + k] * B[(k * size) + j]);
+
+ C[(i * size) + j] = value;
+ } // omp thread
+ } // omp teams
+ } // omp target
+
+ return;
+}
+
+void GPU_mat_mult_no_loops(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ #pragma omp target
+ {
+ #pragma omp teams num_teams(size)
+ {
+ const size_t team_size = (size * omp_get_team_num());
+
+ #pragma omp parallel firstprivate(team_size) num_threads(size)
+ {
+ const size_t tid = omp_get_thread_num();
+ MyData value = (MyData)0;
+ for (size_t k=0 ; k<size ; k++)
+ value += (A[team_size + k] * B[(k * size) + tid]);
+
+ C[team_size + tid] = value;
+ } // omp threads
+ } // omp teams
+ } // omp target
+
+ return;
+}
+
+void check(const MyData *const __restrict__ cpu_matrix,
+ const MyData *const __restrict__ gpu_matrix)
+{
+ int flag;
+ for (size_t i=0 ; i<SIZE ; i++)
+ flag = ((cpu_matrix[i] != gpu_matrix[i]) ? 1 : 0);
+
+ if (!flag)
+ printf("\n\t Result OK");
+ else
+ printf("\n\t Result wrong");
+
+ return;
+}
+
+int main()
+{
+ double time;
+ MyData *buffer = (MyData *)calloc(4 * SIZE, sizeof(MyData));
+ assert(buffer != NULL);
+
+ // host reference matrix A
+ MyData *const restrict A = buffer;
+ MyData *const restrict B = A + SIZE;
+ MyData *const restrict C_CPU = B + SIZE;
+ MyData *const restrict C_GPU = C_CPU + SIZE;
+ for (size_t i=0 ; i<SIZE ; i++)
+ {
+ A[i] = drand48();
+ B[i] = drand48();
+ }
+
+ ////////////////////////// CPU naive algorithm //////////////////////////////////////////
+ CPU_mat_mult(A, B, C_CPU, N);
+ /////////////////////////////////////////////////////////////////////////////////////////
+
+ // copy/alloc data to the GPU
+ #pragma omp target enter data map(to: A[0:SIZE], B[0:SIZE]) map(alloc: C_GPU[0:SIZE])
+
+ /////////////////////////// GPU naive algorithm ////////////////////////////////////////
+ time = 0.0;
+ for (unsigned short int loop=0 ; loop<LOOP ; loop++)
+ {
+ const double start = wall_time();
+ GPU_mat_mult(A, B, C_GPU, N);
+ time += (wall_time() - start);
+ }
+
+ #pragma omp target update from(C_GPU[0:SIZE])
+ check(C_CPU, C_GPU);
+ printf("\n\t GPU naive time %lg [s]\n", (time / LOOP));
+ ////////////////////////////////////////////////////////////////////////////////
+
+ /////////////////////////// GPU naive no loops algorithm ////////////////////////////
+ time = 0.0;
+ for (unsigned short int loop=0 ; loop<LOOP ; loop++)
+ {
+ const double start = wall_time();
+ GPU_mat_mult_no_loops(A, B, C_GPU, N);
+ time += (wall_time() - start);
+ }
+
+ #pragma omp target update from(C_GPU[0:SIZE])
+ check(C_CPU, C_GPU);
+ printf("\n\t GPU naive no loops time %lg [s]\n", (time / LOOP));
+ ////////////////////////////////////////////////////////////////////////////////
+
+ // free CPU memory
+ free(buffer);
+ // free GPU memory
+ #pragma omp target exit data map(delete: A[0:SIZE], B[0:SIZE], C_CPU[0:SIZE])
+
+ printf("\n");
+
+ return EXIT_SUCCESS;
+}
diff --git a/cuda-omp/omp/7/mat_mult_block.c b/cuda-omp/omp/7/mat_mult_block.c
new file mode 100644
index 0000000..cd84ad1
--- /dev/null
+++ b/cuda-omp/omp/7/mat_mult_block.c
@@ -0,0 +1,208 @@
+////////////////////////////////////////////////////////////////////////////////////////////////
+// - Block matrix multiplication algorithm
+//
+// const size_t Nblocks = (N / Bsize);
+//
+// // loop over blocks of matrix C
+// for (size_t ib=0 ; ib<Nblocks ; ib++)
+// {
+// for (size_t jb=0 ; jb<Nblocks ; jb++)
+// {
+//
+// // loop over blocks of rows of A
+// for (size_t kb=0 ; kb<Nblocks ; kb++)
+//
+//
+//
+// for (size_t i=0 ; i<N ; i++)
+// for (size_t j=0 ; j<N ; j++)
+// for (size_t k=0 ; k<_N ; k++)
+// C[(i * N) + j] += A[(i * N) + k] * B[(k * N) + j];
+//
+// } // jb
+// } // ib
+// - Exploit shared-memory.
+////////////////////////////////////////////////////////////////////////////////////////////////
+
+//////////////////////////////////////////////////////////////////////////////////////////////////
+// Author: David Goz
+// mail : david.goz@inaf.it
+// date : 31.07.2024
+// code tested using nvhpc
+//
+// - Compile the code:
+// $ nvc -mp=gpu -gpu=ccnative,debug,lineinfo -target=gpu -Minfo=all -v classwork.c -o classwork_omp
+// - Run the code:
+// $ ./classwork_omp
+//////////////////////////////////////////////////////////////////////////////////////////////////
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <unistd.h>
+#include <time.h>
+#include <assert.h>
+#include <omp.h>
+#include <string.h>
+
+#define N 512
+#define SIZE (N * N) // matrix size
+typedef double MyData; // do not change
+#define BLOCKSIZE 32 // number of threads per block
+
+// sanity check
+#if BLOCKSIZE > 1024
+#error BLOCKSIZE cannot be larger than 1024
+#endif
+
+#define LOOP 100
+#define NDEBUG
+
+double wall_time()
+{
+ struct timespec ts;
+ clock_gettime(CLOCK_PROCESS_CPUTIME_ID, &ts);
+ const double ret = (double) (ts.tv_sec) + (double) ts.tv_nsec * 1.0e-9;
+
+ return ret;
+}
+
+void CPU_mat_mult(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ for (size_t i=0 ; i<size ; i++)
+ for (size_t j=0 ; j<size ; j++)
+ for (size_t k=0 ; k<size ; k++)
+ C[(i * size) + j] += (A[(i * size) + k] * B[(k * size) + j]);
+
+ return;
+}
+
+void GPU_mat_mult(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ #pragma omp target
+ {
+ #pragma omp teams distribute num_teams(size)
+ for (size_t i=0 ; i<size ; i++)
+ {
+ #pragma omp parallel for num_threads(size)
+ for (size_t j=0 ; j<size ; j++)
+ {
+ MyData value = (MyData)0;
+ for (size_t k=0 ; k<size ; k++)
+ value += (A[(i * size) + k] * B[(k * size) + j]);
+
+ C[(i * size) + j] = value;
+ } // omp thread
+ } // omp teams
+ } // omp target
+
+ return;
+}
+
+void GPU_mat_mult_no_loops(const MyData *const restrict A,
+ const MyData *const restrict B,
+ MyData *const restrict C,
+ const size_t size)
+{
+ #pragma omp target
+ {
+ #pragma omp teams num_teams(size)
+ {
+ const size_t team_size = (size * omp_get_team_num());
+
+ #pragma omp parallel firstprivate(team_size) num_threads(size)
+ {
+ const size_t tid = omp_get_thread_num();
+ MyData value = (MyData)0;
+ for (size_t k=0 ; k<size ; k++)
+ value += (A[team_size + k] * B[(k * size) + tid]);
+
+ C[team_size + tid] = value;
+ } // omp threads
+ } // omp teams
+ } // omp target
+
+ return;
+}
+
+void check(const MyData *const __restrict__ cpu_matrix,
+ const MyData *const __restrict__ gpu_matrix)
+{
+ int flag;
+ for (size_t i=0 ; i<SIZE ; i++)
+ flag = ((cpu_matrix[i] != gpu_matrix[i]) ? 1 : 0);
+
+ if (!flag)
+ printf("\n\t Result OK");
+ else
+ printf("\n\t Result wrong");
+
+ return;
+}
+
+int main()
+{
+ double time;
+ MyData *buffer = (MyData *)calloc(4 * SIZE, sizeof(MyData));
+ assert(buffer != NULL);
+
+ // host reference matrix A
+ MyData *const restrict A = buffer;
+ MyData *const restrict B = A + SIZE;
+ MyData *const restrict C_CPU = B + SIZE;
+ MyData *const restrict C_GPU = C_CPU + SIZE;
+ for (size_t i=0 ; i<SIZE ; i++)
+ {
+ A[i] = drand48();
+ B[i] = drand48();
+ }
+
+ ////////////////////////// CPU naive algorithm //////////////////////////////////////////
+ CPU_mat_mult(A, B, C_CPU, N);
+ /////////////////////////////////////////////////////////////////////////////////////////
+
+ // copy/alloc data to the GPU
+ #pragma omp target enter data map(to: A[0:SIZE], B[0:SIZE]) map(alloc: C_GPU[0:SIZE])
+
+ /////////////////////////// GPU naive algorithm ////////////////////////////////////////
+ time = 0.0;
+ for (unsigned short int loop=0 ; loop<LOOP ; loop++)
+ {
+ const double start = wall_time();
+ GPU_mat_mult(A, B, C_GPU, N);
+ time += (wall_time() - start);
+ }
+
+ #pragma omp target update from(C_GPU[0:SIZE])
+ check(C_CPU, C_GPU);
+ printf("\n\t GPU naive time %lg [s]\n", (time / LOOP));
+ ////////////////////////////////////////////////////////////////////////////////
+
+ /////////////////////////// GPU naive no loops algorithm ////////////////////////////
+ time = 0.0;
+ for (unsigned short int loop=0 ; loop<LOOP ; loop++)
+ {
+ const double start = wall_time();
+ GPU_mat_mult_no_loops(A, B, C_GPU, N);
+ time += (wall_time() - start);
+ }
+
+ #pragma omp target update from(C_GPU[0:SIZE])
+ check(C_CPU, C_GPU);
+ printf("\n\t GPU naive no loops time %lg [s]\n", (time / LOOP));
+ ////////////////////////////////////////////////////////////////////////////////
+
+ // free CPU memory
+ free(buffer);
+ // free GPU memory
+ #pragma omp target exit data map(delete: A[0:SIZE], B[0:SIZE], C_CPU[0:SIZE])
+
+ printf("\n");
+
+ return EXIT_SUCCESS;
+}
diff --git a/cuda-omp/omp/miscellaneous/structure.c b/cuda-omp/omp/miscellaneous/structure.c
index c12b212..fc0041c 100644
--- a/cuda-omp/omp/miscellaneous/structure.c
+++ b/cuda-omp/omp/miscellaneous/structure.c
@@ -17,7 +17,8 @@
#include <assert.h>
#include <omp.h>
-#define N 8
+#define SIZE 8
+#define SIZE_2 (SIZE / 2)
typedef double MyData;
typedef struct my_span
@@ -29,26 +30,20 @@ typedef struct my_span
void allocate( span *my_struct,
const size_t size)
{
- span tmp;
/* allocate the buffer on the host memory */
- tmp.ptr = (MyData *)malloc(size * sizeof(MyData));
- assert(tmp.ptr != NULL);
- tmp.N = size;
-
- /* declare how the object 'span' has to be mapped on the device */
- #pragma omp declare mapper(span tmp) map(from: tmp, tmp.ptr[0: tmp.N])
-
- my_struct->ptr = tmp.ptr;
- my_struct->N = tmp.N;
+ my_struct->ptr = (MyData *)calloc(size, sizeof(MyData));
+ assert(my_struct->ptr != NULL);
+ my_struct->N = size;
return;
}
-void print(const span *const A)
+void print(const span *const A,
+ const char *const string)
{
printf("\n");
- for (size_t i=0 ; i<A->N ; i++)
- printf("\n\t array[%i] = %lg", i, A->ptr[i]);
+ for (int i=0 ; i<A->N ; i++)
+ printf("\n\t %s[%d] = %lg", string, i, A->ptr[i]);
printf("\n\n");
return;
@@ -56,27 +51,41 @@ void print(const span *const A)
int main()
{
- span A, B;
+ span A, B, C;
- allocate(&A, N);
- allocate(&B, N);
+ allocate(&A, SIZE);
+ allocate(&B, SIZE);
+ allocate(&C, SIZE);
+
+ /* declare how the object 'span' has to be mapped on the device */
+ #pragma omp declare mapper(all : span S) map(to: S) map(from: S.ptr[0 : S.N])
+ #pragma omp declare mapper(left : span S) map(to: S) map(from: S.ptr[0 : S.N/2])
+ #pragma omp declare mapper(right: span S) map(to: S) map(from: S.ptr[S.N/2: S.N/2])
/* init on the GPU */
- #pragma omp target
+#pragma omp target map(mapper(all): A) map(mapper(left): B) map(mapper(right): C)
{
- for (size_t i=0 ; i<N ; i++)
- {
- A.ptr[i] = (MyData)(i);
- B.ptr[i] = (MyData)(2 * i);
- }
+ #pragma omp loop
+ for (size_t i=0 ; i<SIZE ; i++)
+ A.ptr[i] = (MyData)(i);
+
+ #pragma omp loop
+ for (size_t ii=0 ; ii<SIZE_2 ; ii++)
+ B.ptr[ii] = (MyData)(ii);
+
+ #pragma omp loop
+ for (size_t iii=SIZE_2 ; iii<SIZE ; iii++)
+ C.ptr[iii] = (MyData)(iii);
}
- print(&A);
- print(&B);
+ print(&A, "A");
+ print(&B, "B");
+ print(&C, "C");
/* free the host's memory */
free(A.ptr);
free(B.ptr);
+ free(C.ptr);
return 0;
}
diff --git a/jacobi/serial/not_opt/jacobi_serial_not_opt_len14 b/jacobi/serial/not_opt/jacobi_serial_not_opt_len14
new file mode 100755
index 0000000000000000000000000000000000000000..f546329f4a6d492f2c2a4495e5a65dbeb9ea452d
GIT binary patch
literal 17416
zcmb<-^>JfjWMqH=CI&kO5KoTZ0W1U|85lTh!CWxmz+l0^$>6{s!ywJT#=yY9%D}(?
zQ|AC>!RQ|#!x$JCU^EBV1O^6X1_lNe1_lNTCWwFq6T}1<Edvo|fYDH6z-|NC2bD&#
z86@T-2O=35U^D}R0$2c~ALP~_x?nED58VVX14eIvx?=~FhUo)wLHaoKz{(jo^dS0R
z^aqf^3=9k~8s<Nc-#|D7WB>yLLkQG=FdE`2m@Sa70ND`mv?K+juR$8bV_;~IhJ*`@
z&VdLsz-W+KkWk>$k`$0TL2P2M08~*B)IMC{!U1(ZjD~ubK|d!m$;?DQCq*|WGq1El
zx5C0q*UUt(IA6~Q9IqhrL2BLoLcunHi~+e1=5J8^3o<Y;fYk9bFfj0eRUrtFUJwtY
z52O~P7ZmU8&~(Gaz`(%Cz`(%Gz`(%Ez`y_scOH-i9N30|fk6n1oH-70P%L9puZ2T>
zDGqTL9O98U#Jg~agZv2#J7jMJ;ZUEBL%a!xxFHU4O&sDXIK=aCh%@A5Cgp;XhZ#eB
ze0pwvUVL#$Vo^zaJV+uLB9U5=S;A0UP?VWhlE#phUyz!|P*7S@3?h;<5{nq(;}eUE
zQ;SOC(-Jdt7>Y}ZO7e3Ua#M4QQ%e}ql5_HlQyFp-b8_;N8Pdv&GD}h!(o;)HGILY&
z(^3*EK?-y8%TgKQ<8v~TlA+e;CT8X_=;s%d=qKlt7MG+J=@%CyCTD|1AS_7oPEO4$
z$uBn7%}6XRPtD2EPft$PHPN%sGX=8@4fTxl40J0CD^n5+D$3H0bFvc)O-+;X&5Dwf
zE6U6Al2Ws>^Rgg$hQZy()5$sBNY6;m6cm$845*L+gAY>0$iT=T1<fxYeyLO@CrBkX
zG^6~`2c=jBaCr>M7a%ntEMWi<2jv^6Qjjo%0+Ki=kAVdl7#K8=#6kHECT@Tv4l8p&
zau!J9p!^9E17QawaZug`iGi>Ok~qjsATbaQKoSS#Wsn#MM<9tK`zrxS98?y7)PisZ
zk~k>OfW$z!07)FVoTxw&2bCirwIJMpB#vA*bs&j@$}Nyu5T1Y}4l2_?Vj%3%{D#A$
zo3&R%fx)BoKnd%M*$fN}9?eHM4ufPF82+0!X(%xKS8dTyVBnW`VEC^J;%9&)Uq1N%
z|Nnp0Dh&mO3{cd)ya47Gf%u>(dwBrN&jRs5QT1{In4bjVgQDo=0x&-c#0N#q%L!n9
z5Qq<ol9vr&z88oOii(#7V7?QG4~l}931Gezhz|<-mjPhD5r_{8@|O-^z7~iN3hI{z
zV7?NF4+`R!3Shnzhz|<dmjYnE5Qq;7(w7WiJ{O1&3d)xs{{8>&(;K3~;nDo$f#d(8
zAE6$dPaVS@!yH2$LmY#GJsRIEU|?Wq{?EkUXUfLF;L%;j;nDed7sxS&-(Cd%|Nq~k
zyIR2T8(74%^Ju5*f*02R|NlR7urlxI{!S?ehNHjgcq@e%7+iV_m^?aNS2%VaaqX>T
zI`TMD=>X%=pBkcT1Q{3%4;=mZ@8Hp2&rTeD8F%r><4Uo^j7NXgxh^^SQuNi)--4kX
zM}G_EHXQw_8F~OjOn6cD@Bjbq&;>4?Klg!RYX1aD2L{6fFXI3G|L@ZI(WUbdM0f{8
z*yG>-|DMVJJq|u+ez6A3{s&@vbhcLf|NsBR%zyv?do=F_u^Bu%TS2kH-xA5n07^jo
zEdi_y3@+WiOOF232%T{BM);GXzXfVLj{X+NZE$Tp$=_kf$iU#({F9Bpj}xS_^Or|A
z*r=m_^G)>y85oWv)k>Z}_=@@H-+Wnn5Z}=<{J{TH9-Tit4nAV@=nWBfIQlpLm`C#g
zhofizeLni{*@+{MYb_5j9yzpk#>u1qo*e)&4j%n0=yBlaUqOcl9+nqM&$@K}YJT*i
z^Ma%0b^exh3=9k|o$oZS_4+aWKi2$@sVwm5&%7H)9#^uSVmNv-{2R<kKMM{s9R2x^
zb-DlpgGaB(^`pO@-8lNQ;J~{hk510GdGyz_10cqUqdx_uj{X$%66ieD{OAwk1<Q*i
zen)R>9XR?i{A8~Oqet=usDT~_A2A<!^l|3#qrU}N#2qb<)~Fr5t#z>Z0255qMV!OY
z@@x%1NIOXE$I;({HR1xTxA{9v85tOwf3xxTonQtz|C2}eTu{)1!tFahB;1Zd!;M{l
zf#K-E@c*D-y8#V0VVuEsbjArtupONN3Nd)F9Rda0Cup!;<8Rr*z`)R5x&kfOet?4Q
z67$iA;kRKf`dM%s9Bc{v;9$E9PN|S!dpP41B-kF#xN!8RAlK2Kf<Z{Z76S4&DA>SZ
zR=WafA}HK?Js3d&augJF!j{KtG?79MB<Lc{VR^nr_~>n|2S;xU9D%8I5f*Ui{O8j7
z{pgSK7wf?x2jY2jwt`BF7ju~z7(6?Vdvv?5XgyG8<J$VR&iwxg&)#?r-`<kr0*<||
z|9yIOj)IDtG=6y(P@ZP?>3n7Q&7=7+hevM(#|!V@|Nr|~elC6XVj>GjI(Efz*A7te
zc=WemtH=NU|2uDV-tUZEaqtC`OK<ppm)?y3oi|)Mr)K>B|G)XcAJ2nNS$hkZdLtM+
zTXR4{2bg`5Z}=pC>nxq%k$l_p;A3XbgU^^dU6*(^AJy>boErfuMGih@_UYUTD*t@C
zw}L#9#xLIjvXr6QcY#OuR8Ycj?R?<V`QM{^D=4GBP+?(Ua13z)N%w~C$YONd0V)VQ
zdVLT0bpG?{{O)nwbqC0R7uKu{3=vVTtxx#-Hi061FUX-j-BUr9flCRO?p}~XeY(L;
z@$CHJ(R`c<Zog+Y&v(ymk#8QICp<JScr?G@@M!)az~A<Tfq`KMC>R}^e=za4J!N2E
z0Eg$%zXCc;NB;`yFnTl}Q1G#QQTo=!(sxCP-qGIzEK;DzmUOXBolqlpWMQ=ci_npU
z)q*U7F4n0_YIwjRQeY8DkN?L!EWejF9{o{%^tXWPj_%Tq{T5OV3}9a#{VnL)@S^DF
z|NsA8K_!uE!_lAmtWvxT4E*h9KqX`B0*~&!pkQ?AeAIdG#Sacp27A<bZ-0aoxSTZq
z`TzfmO-u|7`)47{i~+~)3+bQ#|9kYNUNAfm?HJ`4;~480=NKRB*_-s^FsQx+JAhxF
z0py)SAnzRX_<z8|@_6aHm!QO9c**d<%jJLm{|7VDfB*l#KSI)h;l<J)|Nrj?#n6lF
zpa1{wTqotguq%*(fnnbRNe71g98wMpFAn|q|KISfXXnK@kIqLP%}+iAgm@l5&>-Q!
z;Bow*0GK*Hfy05}#m*o9|9^p{j~y!v5*R#eL1c+AIL1NMKIeXr!2btn6Apu8B~1^e
zejiA}qwyV>>ipnw@SVNK!5_>XjK@4K{wooLo4^h=0bPHC?L!3x28L2`P+j?=?*IS)
z$5_uRC@?U_LR3OXgnh635NuHRIfi&7{}1)({0pjqf;&IIkpA)iKgh{`o$oz6kMOVm
z?E$T+n11~K@7c+s0+RG<{>xaW;*os1xAdPUxW+#CklC}h<coj@;|Y)ELmVERPCq;>
zua$oA>3k3!>lkx*r+}mb!!8j}5HFE%VA#J&!hzvM{P+L=Arj(niK8G1x9|V|L&^hi
z<==VLF$C<!&hsEAzWDb4|Na*sb@#sg|KEDMF0R-4zfW(;Q4UCXqJmVO7~TezIa=`i
z<I(NPu>ZD%131A*fKv~ssClvF+yDQbmUl|;?+3LeJerSkyeRwzDtW$^?uiCl77Iyy
zFSEY=|BnbmkLEWLhTl9opTB7O_W%FU58)?|Jp8QXqQY_XcO2_+b_NCzmq*2;`G~~P
z-w_87J9eIQ?Db$g`dfh4k&^+Ee>^)+xwamta|gx8(a$j~hkUwOj(hw+aO7aMAcKeH
ziPBfT$+y5x1r=A2Z~>=xsJf%SBM%&DtQKUE@UT2s`V`a*=-b1<z~I>H_}`;j99$r*
z1=q<Zj{eo)J;}kqaPS3Fug8CIf#GrRA*)C8K@F(6Ut?I<j{X&7VFHyGAj{vnbe?Q}
z_@~#A(b4i`N&L|t<*uD495qikHveEOQT%`K=uZvSG7biYqd#L<&K><J$a2c1^I-GC
ze~c#_El-q)c+5EZTYxnbq~Lc9%P~+EJ>;8w$))p@NAocbuWp;`pxVcy`8dbR4PXEN
zKl-Em^#gu+hQo#@!9IRH2^1=b@CGGwSd%aq)K2s`4r&WCcy#)xoN(zp>C)+=^2aCH
zSHPpQM&*V_XNk%MkIoR44=$Z9DnDGTeI@w&r!p}xcqV^1`m4@14(u!e(QE9Wn$+g2
z2B>k~S)%g5r`P5ysJK2_slmGwELr{oWa5gWl^U!k*g?M4_zp6CgGX=q1&`+U3P=BH
z==OtDc9y8@=sf4rdGP384b{2q3=A%s$3R*Sbe5>R`G2hQLuZM~i{@v44!&Sw{McEd
zvc#pg;D4u&$^y^-$2<={V}-;+ca6#)%Nms*{QcJ$LDim%$_1b96qN&B-6AKvx=T*{
z@anEP!Qs*EqH@ElS4PdNo96_G^WkW<fB=Yb!mpRb^5|~?j-#I=J;3UA`1Fb}fuuHg
zFrM(VJOpB{aO9tS*n{zcr{w|u$p>6GZ@6^FsQhuUj8XaF$UpV4i{cI6&cmLXw>_JG
z3pnymKkU+SgTEz{k%7UlS4Q2Z^MH@$4M_D93TiynsJvm~Z<`2~)=~26m9anas8;f(
z;Q_y16+6FP87qEyh9d`S1z4Dl9ITCSVFhs_SXhq!jbLFu`Zx0Ek%P5@EFi%c7G_ZD
zJNh@|$kD%%hmUGUTsT@O`rydpTFC>9M;x^yu7HJpiXH&bo|eZ-S9H6mEU|P^dBNW@
z7nJfl4|KbzEO6-#QF-C2dC0}mMdeLNgb%Ze%7Y^dD+OW>{W!8PF~a2#$I;(HH;*i=
zj9@tiCXXLESQ&Zn$ihlN77!W3a>%vwf@AZ8e=dv%jvTBEIdJr6<l&>MBW@f8*_|A5
z!PW9mi9R@qyBzv)^f$!F2v9PQU^#yDcO;}rI{G_?<<QZeA;*p+ecX8g6!0I8RtpMv
z{y*qq=_|qC%l{9Y@I87{Z$v{%8OL~N`3+M1z@^(o<qsrX9Q`c-c6J2BKCCu*{6BE?
zXUK7|Wsa5yYEQrJ>2^{1;c9r`<*R@H|MSZ;;LA^)H7Y+keN?_!p5pJ9Wn^G*wLHP!
zFTx0_vQKz)e)dTI<e7ZIMe~49ugq0Y76E0i3%wPLoh~XLJbS%4I)8N5s61(Y{>PE=
z$H5m&pp5y#r8nTeW3K~aXN}4O-{coQ$samhRPH$TIxu<uKh)`>a>MiBBUYc}AD#!F
zf}3NW2cI)T{0s8`507r1&pzEOA3&TJ-n}9c9=$e<9-S9FG=D&nUmK`fSflcR$*b2!
z(vg4q5yJznEl)~{(k2}JE5O5k^sgWfxIFW;JXHF&J4EGyWr)g?5+iU6M-yChYPeYY
zsNAWM2MZa3g$zJKH=ykxOR$iI$NyuNB`RO|dsF^ni?Xx{ASYgcwnq+v3UHPK9{&$G
zS{|%@`}z>SJcDDH;Q`Oie~zJ!kV-B2-97_J2ZrF?pvH(t^BaY3QBg?+22h1^@WcQA
z9-YtjS3nv`LLdJB-`6JP0G6Bj0aO_u0o84);1-2vH_JVr&W9eIul9jjfgb-4gImvY
zKm7miXnD8xfluf2|E#+u6c}FYdJpPe9#H^Q%b?cOVMJm8xi8$Y^S@)5XXiJ^P*0E(
zeR@?dh$=9Idi3go=wKhsr-rvZ4*q2ZjcmM-1sT@*wnPrrp7H5?<<t4irSreze~~*L
z&4(F1EMJ$t0u6Y;>~oAiJON?l=fj}kkT)D470@9OP@;lkMg}GZJ^kXMWc`x-{G4LF
zWQL^DGzCRl1r3d|{LB;uElmSWkWjQiEV`(wg`QCwsBr``OcyTzseoWd(D-3uNou^2
zOG;9XV^L9JB}9UOfuW&A&oL`8IX@}W$R)lwwJ0+&Cq6H~BtE~OBwp3PC`B(R6Eyq@
zR?o$$5L%p=o@%9FlUk9QTw0Qtl#^<wU{evFUX+;<Uz}M5;#DH^Kn)KDE>2E`%%s$G
z9fi!)yc7jn1=SQC1=SR5g}jRRlKc{|0LTK6x-77|EHrhMaCKZ<4A8KHxH=xg*?P$!
zvp`s*A{{OKG!@j;6f`O^gh8qq7`QkUJWEoG5=%1k^T3W$&{aq+DgrTc(z(F$ei8BR
zKK@P$6$*Zl5EjTc3Ka^fDO?P83gJbWC7F5Y3aZ6i4CP561Ko=<Qxr1uQZkcMi>*Kc
zA(^?UU=!06qKjj>xER30wGaoOkRR+JTwL@8s5~YO8iM-&|9?F$8e}dgYl0dXpg}g!
zz#3>U4K#oT8Z-lCHhdUnF321Z2AKuIA2S|zZ2j@Xj!~M~8Z?&!8Zcvc^Z!4n!7uRk
z|9=H028MZ`|NjSdbsN6?|Nnr2fkENx|NjDv3=9Xp{{J7q$iVR7>;M0tkwJs+|Nl2I
zGB7m!`u~3eNc{Kz{|^`$7$*Gr|33rd@W226!{)a@EpV_qs)868D+Cy&dDuB7FtQ7P
z#6fdye4qdS2hDMS)POKZ97HpssApgRjsJtxJG_RtiGfeRjZeahpSzr+fx%wNTFY2P
z38YU4)b3<pV3_gw|NmJa0Vh6z9wrw)i8f{rK80o$rZ6r(2}eEwCq52G&>$jco~8se
zpY!Ga|7lS5eN0Y#61~i>d<s1*j(i$ztnPdU&1_61d>U?i3Ql|yPGIdI#S9D}_kspe
zGQRx(ZvYZ<;ZsQG;{c1F0p;N@|Nk3-A^;@Ez`y`<-xCG~h817`|91vOup^&98<R6H
zn{PjJA4@N54_h-k^Gzl$9(E2lko!RDL6zZ&AOHX7g3Q6D-i8sR-WjAG<R2GC1_q8_
z|Nl=0+3Cb5(8uJC>K{*#e?0jDm@@ec9QiaHk$eT3KWbrQV0iKB|NlNvqI2cpV(<}Q
zV=80t;%6?4U=!eHDytCSV=5~U;AJYy5a3}dOYjh2W-1F{{>{q8!^Pms!_MK#4Kn8p
zBLjoNumAs3q3-Bsa^;ig1BGiZiz}Z-532{CK^vO~pG7mfFP}pWpM@jXIZjB<0f(aq
z69a?CumAr+Neg7uC>{-g(GVC7fzc2c4S~@R7!85Z5Eu=C(GVC7fzc2c4S~@R7@Q&S
z0lL2X0F(v|Wq@3YPJ`AQg2eej1Oo%Z94HMQuVP?eNMMJoPX;aU0!g(<K=_g%4rtvs
zbe%PHIWu^jE=USyE(3HuHE1jeBn}!&1JR&yHV|C_T{jIHa|7`+SRm$uhC)F63s8CR
z;3)$GgE-U!ps^Q_yab3~U|<l1(l81%)B|Qi*X_Z^GC<>KAaT$*5{L$k@quX21T~0$
zzzVS+CVm3Shq)VOA85fGNIi`I@BjaN5Pt<!{eLJQ>Rg5oP<{hc9;B8Y7`Bcb-L7fS
zcwGjiw?XM+Q2H8_eg>t#L1{J)Ncf0BX*DQq2BqDgbQqLQgVJSCx(!OByT{$x*-Amf
z-7i#A0le;3!N|bS)WFQhSV1EwHATTAu>>NnYiy~>08tzaTFI;6?(D3oU<h8Ii*6=5
zpP7M?0k;1DS&o^33BG<GRh*dtmOfF%Sr}mH5>=d)0hWGI#n~9J#XB<tI|D4epsMF!
zfTcTBaZUzU`a~7yg6C6Iac+hNXuP3{^Dux`wjhgv*vt&P46ytS;=?dA10TZzXugMu
zfoNt1eg@dS1rQ&GnHdBaK0x<1z{Ef_GlL*}Uj~Q|!_1(S`_OU%CI+IJ8HC~cF+hA6
zW@Zp!fUSpyiGgTlP$>Y-?;t)5Gc$-W%)lxx&H!4M15*d1nHgZY8YaR3-j~9{@bN!F
z1i}K9Ym5xy3~Qj~!^$<#_&Fm39|LIR5lAJhe6s@YZ@@Iil>xNcn};C)T3&<nfUrAM
zya7!-3MvjO-$80YxPSrVR)oJ`<xdS*T!;Zy9)r|^@M5Sru=)Zd2EyyX>M`T%91izi
z1FPp_04)Oo>4%lCxcvJDtX_b@Lk?1&fb_sHws8IhmK4NrKWKeCBbIam8pCJ6lHSz8
z0fL!str$V!Cd$A7jYqI`3=9mQc^DxE9tP05FpwCm9!X%tzTc!2EY8EQKmg(|@EivN
z1A`kAXfF_EI&6iiFA#^Q2hSleFfdF8i-T-L#Y-7Mduqfm((M|sdL9M`L5Mlvxe*2i
zhAm)mR5L)l6JT*31_r1*!E-kZ3=C0Vam;c+h8eqm?|{{VOhm=saEJ>sVK+w=hqwz9
z#J}kIHxL}2sCI&Q2{_CtVPaqqWD;XofEK@vVD&r<8-yU?2A(5gU|^UCQjb*b!SeM2
zsQLnEIK$?eHbcch3kN~&0MEfNFfbeji-SZ^F&j7@#ThoBxyK8vo`(UpKM$tsIyf8<
z`5Gqv7^)sto`dI*7#J8{K*brvApQc+Wic=?d<KidbR%dE9PulSL)-v|xDyWXFlI>j
zqnEdP!S2CQf8m+Z&P&V8N!2r8h>s5lat#TIjCTw5a}M$J_lu9$Gh|3EDk&~0O-s{D
zW=KgbN=?rMZTpEY$&F9W$<IqIW{8hZ$&XLZ$xlkmiBBoXFDi~tEUjQj&d)8#Ni9iD
z0h^0OMSNOjUS@nEXu&{yYF<fEB|}<KVs2`DN@;FxB}05X5-&b8FS7(iDQs;6HW!2Z
z;TGiR>l*Ls=Mo>!5Ra-og(2Q0($CS?(-~C;yuFGcKHl9gG~U%C9%6!rOArHiD&Em2
z-rvnF*fk_R#L>ye6=W)SS5$IwDcB>hJy#A6$h(z-GxE#hkrqaPrtnjP^K(i;s{lY+
zic!_ar=}#9Br+7ICg<m+6hrrDp(??)-Xa}o+ZL*7lpSEGQjmRMsA3pfz0jpW&V%d(
zL*JDK*<FUAzBoTVBQY-}CzZhqyb#1OCq2I?vm_&zA>PwJzPKbMGcUfhI5maADZex?
zC9$Z|IX^EY6BIn4gar>`Py&GNRYSKTskj*GBJg5|<op82B9UU)1~ya;;B9iKLeL#_
zsDhxqb*LhcJ#wgG@$n(P&d_wmpjTX(TauW>pa%*IFr5KoW#*Nn78Nk)<>i;8>Nz@j
z>Xsy?Gw5X`7H2T%rBvn>SLQ<Kk|GAZ^t@8NvecsD%=|nQPJ9uAUQudJB1i+2RghD{
zpa%{cy^_?55(Yg`o?_4|$_K?0gI;O|G#6#06d`yJ9gyq<(E($pROTh-W+pS}rRSG`
z2|cg{5W|v+iy8Ei^K)}k^FZqj5G5UXEKvUz)O3K=|FC`pY=1MT9f)i)NG&p7nSp@;
z)Te~$hwT%F?H@+p`w6RWU}k{!Z-Uy8F#WK8Nd;6O0;&+E57u9W@j?4wL49VZa(Msc
z0#slI)WIM-Kunl^SW6hRzy{Q(hAM~mb26Y7!20D-XMo$wFgAz=wPQhjVwiqdf2aXu
zAp--05=araZ3$t()WPTw1_lOD9~;Jp^_x~e^~1ss#6@;L#54v5&_XDXQ(^i~K>Oh*
zKpGer;4OZT8jv1z|K~#8<^+;tV1W0-9zgZO#uY&NL1Pi<`pd!XUC91TkP)zPl>|_t
zU|<06{RD}_FvuJj8$=hO>4){}8ld_?>xMw;U>LSX5k}Xd>4&u!CO{qf2{r%(H2|g$
z%4KLq(+}$>wm=gICrlxfMt6TFR6i{IVg1PkQ2h^J5>WlH_=Iv9CZg$w^*dqv=s|0&
zU}`}$y8ZLf^uzk6A3zzKfq@|g&3<UGfQ$m;mC!xxpgaH)fSoG<+gFdS9>#~!TR|Ik
z7$D~&fb_%rVX$)tz=8+@rVmCRWPqd{nEkME9oV@Au=5F^$q_yd2aCs3Q2)d9!}`xZ
z^dWY_?1Atg`|Uw&nD}`#{qTMgw4*Ttq!<aq^ug#GX!>F8`593CFn7X?fa!zr?}65g
zA%!2T9k~LkA6-3q{6B`;4^<8`mI0<8)c=MHLe+x!pt1x;qa_iLFvAC^fuKDDAOX<0
zCn!o`?tvz622lQhg*`|e4DYaju#(UuVfI2~K|NWdv;z|al^dW1XbcPtp#C(HFh=<c
E02<39ApigX
literal 0
HcmV?d00001
--
GitLab