Date: Tue, 23 Aug 2011 21:11:24 +0000
Subject: [PATCH 008/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6764
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/Section_commands.html | 21 +++++++++++----------
doc/Section_commands.txt | 1 +
doc/pair_coul.html | 26 ++++++++++++++++++++++++++
doc/pair_coul.txt | 26 +++++++++++++++++++++++++-
4 files changed, 63 insertions(+), 11 deletions(-)
diff --git a/doc/Section_commands.html b/doc/Section_commands.html
index 0618a5d250..3671b787b1 100644
--- a/doc/Section_commands.html
+++ b/doc/Section_commands.html
@@ -442,16 +442,17 @@ package.
diff --git a/doc/Section_commands.txt b/doc/Section_commands.txt
index f9b9b1a189..8bc22c3a49 100644
--- a/doc/Section_commands.txt
+++ b/doc/Section_commands.txt
@@ -686,6 +686,7 @@ package"_Section_accelerate.html.
"cg/cmm/coul/long/gpu"_pair_cmm.html,
"cg/cmm/cuda"_pair_cmm.html,
"cg/cmm/gpu"_pair_cmm.html,
+"coul/long/gpu"_pair_coul.html,
"eam/alloy/cuda"_pair_eam.html,
"eam/alloy/opt"_pair_eam.html,
"eam/cuda"_pair_eam.html,
diff --git a/doc/pair_coul.html b/doc/pair_coul.html
index 3cdc20760b..2d3b2f623d 100644
--- a/doc/pair_coul.html
+++ b/doc/pair_coul.html
@@ -15,12 +15,16 @@
pair_style coul/long command
+pair_style coul/long/gpu command
+
Syntax:
pair_style coul/cut cutoff
pair_style coul/debye kappa cutoff
pair_style coul/long cutoff
+pair_style coul/long/gpu cutoff
+
- cutoff = global cutoff for Coulombic interactions
- kappa = Debye length (inverse distance units)
@@ -89,6 +93,28 @@ Coulombic cutoff specified in the pair_style command.
+Styles with a cuda, gpu, or opt suffix are functionally the same
+as the corresponding style without the suffix. They have been
+optimized to run faster, depending on your available hardware, as
+discussed in this section of the manual.
+The accelerated styles take the same arguments and should produce the
+same results, except for round-off and precision issues.
+
+These accelerated styles are part of the "user-cuda", "gpu", and "opt"
+packages respectively. They are only enabled if LAMMPS was built with
+those packages. See the Making LAMMPS
+section for more info.
+
+You can specify the accelerated styles explicitly in your input script
+by including their suffix, or you can use the -suffix command-line
+switch when you invoke LAMMPS, or you can use
+the suffix command in your input script.
+
+See this section of the manual for more
+instructions on how to use the accelerated styles effectively.
+
+
+
Mixing, shift, table, tail correction, restart, rRESPA info:
For atom type pairs I,J and I != J, the cutoff distance for the
diff --git a/doc/pair_coul.txt b/doc/pair_coul.txt
index d906461547..a51a3e15cd 100644
--- a/doc/pair_coul.txt
+++ b/doc/pair_coul.txt
@@ -9,12 +9,14 @@
pair_style coul/cut command :h3
pair_style coul/debye command :h3
pair_style coul/long command :h3
+pair_style coul/long/gpu command :h3
[Syntax:]
pair_style coul/cut cutoff
pair_style coul/debye kappa cutoff
-pair_style coul/long cutoff :pre
+pair_style coul/long cutoff
+pair_style coul/long/gpu cutoff :pre
cutoff = global cutoff for Coulombic interactions
kappa = Debye length (inverse distance units) :ul
@@ -84,6 +86,28 @@ Coulombic cutoff specified in the pair_style command.
:line
+Styles with a {cuda}, {gpu}, or {opt} suffix are functionally the same
+as the corresponding style without the suffix. They have been
+optimized to run faster, depending on your available hardware, as
+discussed in "this section"_Section_accelerate.html of the manual.
+The accelerated styles take the same arguments and should produce the
+same results, except for round-off and precision issues.
+
+These accelerated styles are part of the "user-cuda", "gpu", and "opt"
+packages respectively. They are only enabled if LAMMPS was built with
+those packages. See the "Making LAMMPS"_Section_start.html#2_3
+section for more info.
+
+You can specify the accelerated styles explicitly in your input script
+by including their suffix, or you can use the "-suffix command-line
+switch"_Section_start.html#2_6 when you invoke LAMMPS, or you can use
+the "suffix"_suffix.html command in your input script.
+
+See "this section"_Section_accelerate.html of the manual for more
+instructions on how to use the accelerated styles effectively.
+
+:line
+
[Mixing, shift, table, tail correction, restart, rRESPA info]:
For atom type pairs I,J and I != J, the cutoff distance for the
From 7a1f6fc1ee186a10d4464ab13451218471a1fc54 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 21:11:35 +0000
Subject: [PATCH 009/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6765
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
lib/gpu/Nvidia.makefile | 14 ++++++++++++++
lib/gpu/README | 17 ++++++++++-------
2 files changed, 24 insertions(+), 7 deletions(-)
diff --git a/lib/gpu/Nvidia.makefile b/lib/gpu/Nvidia.makefile
index 1776cb44a9..b126f3433a 100644
--- a/lib/gpu/Nvidia.makefile
+++ b/lib/gpu/Nvidia.makefile
@@ -54,6 +54,7 @@ OBJS = $(OBJ_DIR)/pair_gpu_atom.o $(OBJ_DIR)/pair_gpu_ans.o \
$(OBJ_DIR)/ljc_cut_gpu_memory.o $(OBJ_DIR)/ljc_cut_gpu.o \
$(OBJ_DIR)/ljcl_cut_gpu_memory.o $(OBJ_DIR)/ljcl_cut_gpu.o \
$(OBJ_DIR)/lj_class2_long.o $(OBJ_DIR)/lj_class2_long_ext.o \
+ $(OBJ_DIR)/coul_long_gpu_memory.o $(OBJ_DIR)/coul_long_gpu.o \
$(OBJ_DIR)/morse_gpu_memory.o $(OBJ_DIR)/morse_gpu.o \
$(OBJ_DIR)/crml_gpu_memory.o $(OBJ_DIR)/crml_gpu.o \
$(OBJ_DIR)/cmm_cut_gpu_memory.o $(OBJ_DIR)/cmm_cut_gpu.o \
@@ -76,6 +77,7 @@ PTXS = $(OBJ_DIR)/pair_gpu_dev_kernel.ptx \
$(OBJ_DIR)/ljc_cut_gpu_kernel.ptx $(OBJ_DIR)/ljc_cut_gpu_ptx.h \
$(OBJ_DIR)/ljcl_cut_gpu_kernel.ptx $(OBJ_DIR)/ljcl_cut_gpu_ptx.h \
$(OBJ_DIR)/lj_class2_long.ptx $(OBJ_DIR)/lj_class2_long_ptx.h \
+ $(OBJ_DIR)/coul_long_gpu_kernel.ptx $(OBJ_DIR)/coul_long_gpu_ptx.h \
$(OBJ_DIR)/morse_gpu_kernel.ptx $(OBJ_DIR)/morse_gpu_ptx.h \
$(OBJ_DIR)/crml_gpu_kernel.ptx $(OBJ_DIR)/crml_gpu_ptx.h \
$(OBJ_DIR)/cmm_cut_gpu_kernel.ptx $(OBJ_DIR)/cmm_cut_gpu_ptx.h \
@@ -252,6 +254,18 @@ $(OBJ_DIR)/ljcl_cut_gpu_memory.o: $(ALL_H) ljcl_cut_gpu_memory.h ljcl_cut_gpu_me
$(OBJ_DIR)/ljcl_cut_gpu.o: $(ALL_H) ljcl_cut_gpu_memory.h ljcl_cut_gpu.cpp charge_gpu_memory.h
$(CUDR) -o $@ -c ljcl_cut_gpu.cpp -I$(OBJ_DIR)
+$(OBJ_DIR)/coul_long_gpu_kernel.ptx: coul_long_gpu_kernel.cu pair_gpu_precision.h
+ $(CUDA) --ptx -DNV_KERNEL -o $@ coul_long_gpu_kernel.cu
+
+$(OBJ_DIR)/coul_long_gpu_ptx.h: $(OBJ_DIR)/coul_long_gpu_kernel.ptx $(OBJ_DIR)/coul_long_gpu_kernel.ptx
+ $(BSH) ./geryon/file_to_cstr.sh $(OBJ_DIR)/coul_long_gpu_kernel.ptx $(OBJ_DIR)/coul_long_gpu_ptx.h
+
+$(OBJ_DIR)/coul_long_gpu_memory.o: $(ALL_H) coul_long_gpu_memory.h coul_long_gpu_memory.cpp $(OBJ_DIR)/coul_long_gpu_ptx.h $(OBJ_DIR)/charge_gpu_memory.o
+ $(CUDR) -o $@ -c coul_long_gpu_memory.cpp -I$(OBJ_DIR)
+
+$(OBJ_DIR)/coul_long_gpu.o: $(ALL_H) coul_long_gpu_memory.h coul_long_gpu.cpp charge_gpu_memory.h
+ $(CUDR) -o $@ -c coul_long_gpu.cpp -I$(OBJ_DIR)
+
$(OBJ_DIR)/morse_gpu_kernel.ptx: morse_gpu_kernel.cu pair_gpu_precision.h
$(CUDA) --ptx -DNV_KERNEL -o $@ morse_gpu_kernel.cu
diff --git a/lib/gpu/README b/lib/gpu/README
index 5f6637a9c4..61486b0d15 100644
--- a/lib/gpu/README
+++ b/lib/gpu/README
@@ -53,8 +53,9 @@ Current pair styles supporting GPU acceleration:
7. morse
8. cg/cmm
9. cg/cmm/coul/long
- 10. gayberne
- 11. pppm
+ 10. coul/long
+ 11. gayberne
+ 12. pppm
MULTIPLE LAMMPS PROCESSES
@@ -68,7 +69,7 @@ LAMMPS user manual for details on running with GPU acceleration.
BUILDING AND PRECISION MODES
To build, edit the CUDA_ARCH, CUDA_PRECISION, CUDA_HOME variables in one of
-the Makefiles. CUDA_ARCH should be set based on the compute capability of
+the Makefiles. CUDA_ARCH should be set based on the compute capability of
your GPU. This can be verified by running the nvc_get_devices executable after
the build is complete. Additionally, the GPU package must be installed and
compiled for LAMMPS. This may require editing the gpu_SYSPATH variable in the
@@ -100,10 +101,11 @@ NOTE: Double precision is only supported on certain GPUs (with
with -DFFT_SINGLE. For details on configuring FFT support in
LAMMPS, see http://lammps.sandia.gov/doc/Section_start.html#2_2_4
-NOTE: For Tesla and other graphics cards with compute capability>=1.3,
+NOTE: For graphics cards with compute capability>=1.3 (e.g. Tesla C1060),
make sure that -arch=sm_13 is set on the CUDA_ARCH line.
-NOTE: For Fermi, make sure that -arch=sm_20 is set on the CUDA_ARCH line.
+NOTE: For newer graphics card (a.k.a. "Fermi", e.g. Tesla C2050), make
+ sure that -arch=sm_20 is set on the CUDA_ARCH line.
NOTE: The gayberne/gpu pair style will only be installed if the ASPHERE
package has been installed.
@@ -111,8 +113,9 @@ NOTE: The gayberne/gpu pair style will only be installed if the ASPHERE
NOTE: The cg/cmm/gpu and cg/cmm/coul/long/gpu pair styles will only be
installed if the USER-CG-CMM package has been installed.
-NOTE: The lj/cut/coul/long/gpu, cg/cmm/coul/long/gpu, and pppm/gpu styles
- will only be installed if the KSPACE package has been installed.
+NOTE: The lj/cut/coul/long/gpu, cg/cmm/coul/long/gpu, coul/long/gpu,
+ and pppm/gpu styles will only be installed if the KSPACE package
+ has been installed.
NOTE: The lj/charmm/coul/long will only be installed if the MOLECULE package
has been installed.
From 541c47f4d45df954467efc37aab211ff65174e79 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 21:12:16 +0000
Subject: [PATCH 010/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6766
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
lib/gpu/coul_long_gpu.cpp | 124 ++++++++++
lib/gpu/coul_long_gpu_kernel.cu | 411 +++++++++++++++++++++++++++++++
lib/gpu/coul_long_gpu_memory.cpp | 158 ++++++++++++
lib/gpu/coul_long_gpu_memory.h | 79 ++++++
4 files changed, 772 insertions(+)
create mode 100644 lib/gpu/coul_long_gpu.cpp
create mode 100644 lib/gpu/coul_long_gpu_kernel.cu
create mode 100644 lib/gpu/coul_long_gpu_memory.cpp
create mode 100644 lib/gpu/coul_long_gpu_memory.h
diff --git a/lib/gpu/coul_long_gpu.cpp b/lib/gpu/coul_long_gpu.cpp
new file mode 100644
index 0000000000..60c0d35d7d
--- /dev/null
+++ b/lib/gpu/coul_long_gpu.cpp
@@ -0,0 +1,124 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+/* ----------------------------------------------------------------------
+ Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
+------------------------------------------------------------------------- */
+
+#include
+#include
+#include
+
+#include "coul_long_gpu_memory.h"
+
+using namespace std;
+
+static CL_GPU_Memory CLMF;
+
+// ---------------------------------------------------------------------------
+// Allocate memory on host and device and copy constants to device
+// ---------------------------------------------------------------------------
+int cl_gpu_init(const int inum, const int nall, const int max_nbors,
+ const int maxspecial, const double cell_size, int &gpu_mode,
+ FILE *screen, double host_cut_coulsq, double *host_special_coul,
+ const double qqrd2e, const double g_ewald) {
+ CLMF.clear();
+ gpu_mode=CLMF.device->gpu_mode();
+ double gpu_split=CLMF.device->particle_split();
+ int first_gpu=CLMF.device->first_device();
+ int last_gpu=CLMF.device->last_device();
+ int world_me=CLMF.device->world_me();
+ int gpu_rank=CLMF.device->gpu_rank();
+ int procs_per_gpu=CLMF.device->procs_per_gpu();
+
+ CLMF.device->init_message(screen,"coul/long",first_gpu,last_gpu);
+
+ bool message=false;
+ if (CLMF.device->replica_me()==0 && screen)
+ message=true;
+
+ if (message) {
+ fprintf(screen,"Initializing GPU and compiling on process 0...");
+ fflush(screen);
+ }
+
+ int init_ok=0;
+ if (world_me==0)
+ init_ok=CLMF.init(inum, nall, 300, maxspecial, cell_size, gpu_split,
+ screen, host_cut_coulsq, host_special_coul, qqrd2e,
+ g_ewald);
+
+ CLMF.device->world_barrier();
+ if (message)
+ fprintf(screen,"Done.\n");
+
+ for (int i=0; igpu_barrier();
+ if (message)
+ fprintf(screen,"Done.\n");
+ }
+ if (message)
+ fprintf(screen,"\n");
+
+ if (init_ok==0)
+ CLMF.estimate_gpu_overhead();
+ return init_ok;
+}
+
+void cl_gpu_clear() {
+ CLMF.clear();
+}
+
+int** cl_gpu_compute_n(const int ago, const int inum_full,
+ const int nall, double **host_x, int *host_type,
+ double *sublo, double *subhi, int *tag, int **nspecial,
+ int **special, const bool eflag, const bool vflag,
+ const bool eatom, const bool vatom, int &host_start,
+ int **ilist, int **jnum, const double cpu_time,
+ bool &success, double *host_q, double *boxlo,
+ double *prd) {
+ return CLMF.compute(ago, inum_full, nall, host_x, host_type, sublo,
+ subhi, tag, nspecial, special, eflag, vflag, eatom,
+ vatom, host_start, ilist, jnum, cpu_time, success,
+ host_q, boxlo, prd);
+}
+
+void cl_gpu_compute(const int ago, const int inum_full, const int nall,
+ double **host_x, int *host_type, int *ilist, int *numj,
+ int **firstneigh, const bool eflag, const bool vflag,
+ const bool eatom, const bool vatom, int &host_start,
+ const double cpu_time, bool &success, double *host_q,
+ const int nlocal, double *boxlo, double *prd) {
+ CLMF.compute(ago,inum_full,nall,host_x,host_type,ilist,numj,
+ firstneigh,eflag,vflag,eatom,vatom,host_start,cpu_time,success,
+ host_q,nlocal,boxlo,prd);
+}
+
+double cl_gpu_bytes() {
+ return CLMF.host_memory_usage();
+}
+
+
diff --git a/lib/gpu/coul_long_gpu_kernel.cu b/lib/gpu/coul_long_gpu_kernel.cu
new file mode 100644
index 0000000000..bc3747a7e3
--- /dev/null
+++ b/lib/gpu/coul_long_gpu_kernel.cu
@@ -0,0 +1,411 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+/* ----------------------------------------------------------------------
+ Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
+------------------------------------------------------------------------- */
+
+#ifndef CL_GPU_KERNEL
+#define CL_GPU_KERNEL
+
+#ifdef NV_KERNEL
+
+#include "nv_kernel_def.h"
+texture pos_tex;
+texture q_tex;
+
+#ifdef _DOUBLE_DOUBLE
+__inline double4 fetch_pos(const int& i, const double4 *pos)
+{
+ return pos[i];
+}
+__inline double fetch_q(const int& i, const double *q)
+{
+ return q[i];
+}
+#else
+__inline float4 fetch_pos(const int& i, const float4 *pos)
+{
+ return tex1Dfetch(pos_tex, i);
+}
+__inline float fetch_q(const int& i, const float *q)
+{
+ return tex1Dfetch(q_tex, i);
+}
+#endif
+
+#else
+
+#pragma OPENCL EXTENSION cl_khr_fp64: enable
+#define GLOBAL_ID_X get_global_id(0)
+#define THREAD_ID_X get_local_id(0)
+#define BLOCK_ID_X get_group_id(0)
+#define BLOCK_SIZE_X get_local_size(0)
+#define __syncthreads() barrier(CLK_LOCAL_MEM_FENCE)
+#define __inline inline
+
+#define fetch_pos(i,y) x_[i]
+#define fetch_q(i,y) q_[i]
+#define BLOCK_PAIR 64
+#define MAX_SHARED_TYPES 8
+
+#endif
+
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
+#define EWALD_F (numtyp)1.12837917
+#define EWALD_P (numtyp)0.3275911
+#define A1 (numtyp)0.254829592
+#define A2 (numtyp)-0.284496736
+#define A3 (numtyp)1.421413741
+#define A4 (numtyp)-1.453152027
+#define A5 (numtyp)1.061405429
+
+#define SBBITS 30
+#define NEIGHMASK 0x3FFFFFFF
+__inline int sbmask(int j) { return j >> SBBITS & 3; }
+
+__kernel void kernel_pair(__global numtyp4 *x_, __global numtyp4 *lj1,
+ __global numtyp4* lj3, const int lj_types,
+ __global numtyp *sp_cl_in, __global int *dev_nbor,
+ __global int *dev_packed, __global acctyp4 *ans,
+ __global acctyp *engv, const int eflag,
+ const int vflag, const int inum,
+ const int nbor_pitch, __global numtyp *q_,
+ const numtyp cut_coulsq, const numtyp qqrd2e,
+ const numtyp g_ewald, const int t_per_atom) {
+ int tid=THREAD_ID_X;
+ int ii=mul24((int)BLOCK_ID_X,(int)(BLOCK_SIZE_X)/t_per_atom);
+ ii+=tid/t_per_atom;
+ int offset=tid%t_per_atom;
+
+ __local numtyp sp_cl[4];
+ sp_cl[0]=sp_cl_in[0];
+ sp_cl[1]=sp_cl_in[1];
+ sp_cl[2]=sp_cl_in[2];
+ sp_cl[3]=sp_cl_in[3];
+
+ acctyp e_coul=(acctyp)0;
+ acctyp4 f;
+ f.x=(acctyp)0;
+ f.y=(acctyp)0;
+ f.z=(acctyp)0;
+ acctyp virial[6];
+ for (int i=0; i<6; i++)
+ virial[i]=(acctyp)0;
+
+ if (ii0) {
+ e_coul += prefactor*(_erfc-factor_coul);
+ }
+ if (vflag>0) {
+ virial[0] += delx*delx*force;
+ virial[1] += dely*dely*force;
+ virial[2] += delz*delz*force;
+ virial[3] += delx*dely*force;
+ virial[4] += delx*delz*force;
+ virial[5] += dely*delz*force;
+ }
+ }
+
+ } // for nbor
+ } // if ii
+
+ // Reduce answers
+ if (t_per_atom>1) {
+ __local acctyp red_acc[6][BLOCK_PAIR];
+
+ red_acc[0][tid]=f.x;
+ red_acc[1][tid]=f.y;
+ red_acc[2][tid]=f.z;
+ red_acc[3][tid]=e_coul;
+
+ for (unsigned int s=t_per_atom/2; s>0; s>>=1) {
+ if (offset < s) {
+ for (int r=0; r<4; r++)
+ red_acc[r][tid] += red_acc[r][tid+s];
+ }
+ }
+
+ f.x=red_acc[0][tid];
+ f.y=red_acc[1][tid];
+ f.z=red_acc[2][tid];
+ e_coul=red_acc[3][tid];
+
+ if (vflag>0) {
+ for (int r=0; r<6; r++)
+ red_acc[r][tid]=virial[r];
+
+ for (unsigned int s=t_per_atom/2; s>0; s>>=1) {
+ if (offset < s) {
+ for (int r=0; r<6; r++)
+ red_acc[r][tid] += red_acc[r][tid+s];
+ }
+ }
+
+ for (int r=0; r<6; r++)
+ virial[r]=red_acc[r][tid];
+ }
+ }
+
+ // Store answers
+ if (ii0) {
+ *ap1=(acctyp)0;
+ ap1+=inum;
+ *ap1=e_coul;
+ ap1+=inum;
+ }
+ if (vflag>0) {
+ for (int i=0; i<6; i++) {
+ *ap1=virial[i];
+ ap1+=inum;
+ }
+ }
+ ans[ii]=f;
+ } // if ii
+}
+
+__kernel void kernel_pair_fast(__global numtyp4 *x_, __global numtyp4 *lj1_in,
+ __global numtyp4* lj3_in,
+ __global numtyp* sp_cl_in,
+ __global int *dev_nbor, __global int *dev_packed,
+ __global acctyp4 *ans, __global acctyp *engv,
+ const int eflag, const int vflag, const int inum,
+ const int nbor_pitch, __global numtyp *q_,
+ const numtyp cut_coulsq, const numtyp qqrd2e,
+ const numtyp g_ewald, const int t_per_atom) {
+ int tid=THREAD_ID_X;
+ int ii=mul24((int)BLOCK_ID_X,(int)(BLOCK_SIZE_X)/t_per_atom);
+ ii+=tid/t_per_atom;
+ int offset=tid%t_per_atom;
+
+ __local numtyp sp_cl[4];
+ if (tid<4)
+ sp_cl[tid]=sp_cl_in[tid];
+
+ acctyp e_coul=(acctyp)0;
+ acctyp4 f;
+ f.x=(acctyp)0;
+ f.y=(acctyp)0;
+ f.z=(acctyp)0;
+ acctyp virial[6];
+ for (int i=0; i<6; i++)
+ virial[i]=(acctyp)0;
+
+ __syncthreads();
+
+ if (ii0) {
+ e_coul += prefactor*(_erfc-factor_coul);
+ }
+ if (vflag>0) {
+ virial[0] += delx*delx*force;
+ virial[1] += dely*dely*force;
+ virial[2] += delz*delz*force;
+ virial[3] += delx*dely*force;
+ virial[4] += delx*delz*force;
+ virial[5] += dely*delz*force;
+ }
+ }
+
+ } // for nbor
+ } // if ii
+
+ // Reduce answers
+ if (t_per_atom>1) {
+ __local acctyp red_acc[6][BLOCK_PAIR];
+
+ red_acc[0][tid]=f.x;
+ red_acc[1][tid]=f.y;
+ red_acc[2][tid]=f.z;
+ red_acc[3][tid]=e_coul;
+
+ for (unsigned int s=t_per_atom/2; s>0; s>>=1) {
+ if (offset < s) {
+ for (int r=0; r<4; r++)
+ red_acc[r][tid] += red_acc[r][tid+s];
+ }
+ }
+
+ f.x=red_acc[0][tid];
+ f.y=red_acc[1][tid];
+ f.z=red_acc[2][tid];
+ e_coul=red_acc[3][tid];
+
+ if (vflag>0) {
+ for (int r=0; r<6; r++)
+ red_acc[r][tid]=virial[r];
+
+ for (unsigned int s=t_per_atom/2; s>0; s>>=1) {
+ if (offset < s) {
+ for (int r=0; r<6; r++)
+ red_acc[r][tid] += red_acc[r][tid+s];
+ }
+ }
+
+ for (int r=0; r<6; r++)
+ virial[r]=red_acc[r][tid];
+ }
+ }
+
+ // Store answers
+ if (ii0) {
+ *ap1=(acctyp)0;
+ ap1+=inum;
+ *ap1=e_coul;
+ ap1+=inum;
+ }
+ if (vflag>0) {
+ for (int i=0; i<6; i++) {
+ *ap1=virial[i];
+ ap1+=inum;
+ }
+ }
+ ans[ii]=f;
+ } // if ii*/
+}
+
+#endif
+
diff --git a/lib/gpu/coul_long_gpu_memory.cpp b/lib/gpu/coul_long_gpu_memory.cpp
new file mode 100644
index 0000000000..7ea4077954
--- /dev/null
+++ b/lib/gpu/coul_long_gpu_memory.cpp
@@ -0,0 +1,158 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+/* ----------------------------------------------------------------------
+ Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
+------------------------------------------------------------------------- */
+
+#ifdef USE_OPENCL
+#include "coul_long_gpu_cl.h"
+#else
+#include "coul_long_gpu_ptx.h"
+#endif
+
+#include "coul_long_gpu_memory.h"
+#include
+#define CL_GPU_MemoryT CL_GPU_Memory
+
+extern PairGPUDevice pair_gpu_device;
+
+template
+CL_GPU_MemoryT::CL_GPU_Memory() : ChargeGPUMemory(),
+ _allocated(false) {
+}
+
+template
+CL_GPU_MemoryT::~CL_GPU_Memory() {
+ clear();
+}
+
+template
+int CL_GPU_MemoryT::bytes_per_atom(const int max_nbors) const {
+ return this->bytes_per_atom_atomic(max_nbors);
+}
+
+template
+int CL_GPU_MemoryT::init(const int nlocal, const int nall, const int max_nbors,
+ const int maxspecial, const double cell_size,
+ const double gpu_split, FILE *_screen,
+ const double host_cut_coulsq, double *host_special_coul,
+ const double qqrd2e, const double g_ewald) {
+ int success;
+ success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,
+ gpu_split,_screen,coul_long_gpu_kernel);
+ if (success!=0)
+ return success;
+
+ // we don't have atom types for coulomb only,
+ // but go with the minimum so that we can use
+ // the same infrastructure as lj/cut/coul/long/gpu.
+ int lj_types=1;
+ shared_types=false;
+ int max_shared_types=this->device->max_shared_types();
+ if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) {
+ lj_types=max_shared_types;
+ shared_types=true;
+ }
+ _lj_types=lj_types;
+
+ // Allocate a host write buffer for data initialization
+ UCL_H_Vec host_write(lj_types*lj_types*32,*(this->ucl_device),
+ UCL_WRITE_OPTIMIZED);
+
+ for (int i=0; iucl_device),UCL_READ_ONLY);
+ lj3.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
+
+ sp_cl.alloc(4,*(this->ucl_device),UCL_READ_ONLY);
+ for (int i=0; i<4; i++) {
+ host_write[i]=host_special_coul[i];
+ }
+ ucl_copy(sp_cl,host_write,4,false);
+
+ _cut_coulsq=host_cut_coulsq;
+ _qqrd2e=qqrd2e;
+ _g_ewald=g_ewald;
+
+ _allocated=true;
+ this->_max_bytes=lj1.row_bytes()+lj3.row_bytes()+sp_cl.row_bytes();
+ return 0;
+}
+
+template
+void CL_GPU_MemoryT::clear() {
+ if (!_allocated)
+ return;
+ _allocated=false;
+
+ lj1.clear();
+ lj3.clear();
+ sp_cl.clear();
+ this->clear_atomic();
+}
+
+template
+double CL_GPU_MemoryT::host_memory_usage() const {
+ return this->host_memory_usage_atomic()+sizeof(CL_GPU_Memory);
+}
+
+// ---------------------------------------------------------------------------
+// Calculate energies, forces, and torques
+// ---------------------------------------------------------------------------
+template
+void CL_GPU_MemoryT::loop(const bool _eflag, const bool _vflag) {
+ // Compute the block size and grid size to keep all cores busy
+ const int BX=this->block_size();
+ int eflag, vflag;
+ if (_eflag)
+ eflag=1;
+ else
+ eflag=0;
+
+ if (_vflag)
+ vflag=1;
+ else
+ vflag=0;
+
+ int GX=static_cast(ceil(static_cast(this->ans->inum())/
+ (BX/this->_threads_per_atom)));
+
+ int ainum=this->ans->inum();
+ int nbor_pitch=this->nbor->nbor_pitch();
+ this->time_pair.start();
+ if (shared_types) {
+ this->k_pair_fast.set_size(GX,BX);
+ this->k_pair_fast.run(&this->atom->dev_x.begin(), &lj1.begin(),
+ &lj3.begin(), &sp_cl.begin(),
+ &this->nbor->dev_nbor.begin(),
+ &this->_nbor_data->begin(),
+ &this->ans->dev_ans.begin(),
+ &this->ans->dev_engv.begin(), &eflag, &vflag,
+ &ainum, &nbor_pitch, &this->atom->dev_q.begin(),
+ &_cut_coulsq, &_qqrd2e, &_g_ewald,
+ &this->_threads_per_atom);
+ } else {
+ this->k_pair.set_size(GX,BX);
+ this->k_pair.run(&this->atom->dev_x.begin(), &lj1.begin(), &lj3.begin(),
+ &_lj_types, &sp_cl.begin(), &this->nbor->dev_nbor.begin(),
+ &this->_nbor_data->begin(), &this->ans->dev_ans.begin(),
+ &this->ans->dev_engv.begin(), &eflag, &vflag, &ainum,
+ &nbor_pitch, &this->atom->dev_q.begin(), &_cut_coulsq,
+ &_qqrd2e, &_g_ewald, &this->_threads_per_atom);
+ }
+ this->time_pair.stop();
+}
+
+template class CL_GPU_Memory;
diff --git a/lib/gpu/coul_long_gpu_memory.h b/lib/gpu/coul_long_gpu_memory.h
new file mode 100644
index 0000000000..04914a2514
--- /dev/null
+++ b/lib/gpu/coul_long_gpu_memory.h
@@ -0,0 +1,79 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+/* ----------------------------------------------------------------------
+ Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
+------------------------------------------------------------------------- */
+
+#ifndef CL_GPU_MEMORY_H
+#define CL_GPU_MEMORY_H
+
+#include "charge_gpu_memory.h"
+
+template
+class CL_GPU_Memory : public ChargeGPUMemory {
+ public:
+ CL_GPU_Memory();
+ ~CL_GPU_Memory();
+
+ /// Clear any previous data and set up for a new LAMMPS run
+ /** \param max_nbors initial number of rows in the neighbor matrix
+ * \param cell_size cutoff + skin
+ * \param gpu_split fraction of particles handled by device
+ *
+ * Returns:
+ * - 0 if successfull
+ * - -1 if fix gpu not found
+ * - -3 if there is an out of memory error
+ * - -4 if the GPU library was not compiled for GPU
+ * - -5 Double precision is not supported on card **/
+ int init(const int nlocal, const int nall, const int max_nbors,
+ const int maxspecial, const double cell_size,
+ const double gpu_split, FILE *screen,
+ const double host_cut_coulsq, double *host_special_coul,
+ const double qqrd2e, const double g_ewald);
+
+ /// Clear all host and device data
+ /** \note This is called at the beginning of the init() routine **/
+ void clear();
+
+ /// Returns memory usage on device per atom
+ int bytes_per_atom(const int max_nbors) const;
+
+ /// Total host memory used by library for pair style
+ double host_memory_usage() const;
+
+ // --------------------------- TYPE DATA --------------------------
+
+ /// lj1 dummy
+ UCL_D_Vec lj1;
+ /// lj3 dummy
+ UCL_D_Vec lj3;
+ /// Special Coul values [0-3]
+ UCL_D_Vec sp_cl;
+
+ /// If atom type constants fit in shared memory, use fast kernels
+ bool shared_types;
+
+ /// Number of atom types
+ int _lj_types;
+
+ numtyp _cut_coulsq, _qqrd2e, _g_ewald;
+
+ private:
+ bool _allocated;
+ void loop(const bool _eflag, const bool _vflag);
+};
+
+#endif
+
From e77ec4dea6c9c91fde9e1cb6cdbb63ce83dad07e Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 21:12:28 +0000
Subject: [PATCH 011/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6767
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/GPU/Install.sh | 7 +++++++
src/GPU/Package.sh | 6 ++++++
src/GPU/pair_cg_cmm_coul_long_gpu.cpp | 4 ++--
src/GPU/pair_cg_cmm_coul_msm_gpu.cpp | 2 +-
src/GPU/pair_gayberne_gpu.cpp | 6 +++---
src/GPU/pair_lj96_cut_gpu.cpp | 2 +-
src/GPU/pair_lj_charmm_coul_long_gpu.cpp | 4 ++--
src/GPU/pair_lj_class2_coul_long_gpu.cpp | 4 ++--
src/GPU/pair_lj_class2_gpu.cpp | 2 +-
src/GPU/pair_lj_cut_coul_cut_gpu.cpp | 4 ++--
src/GPU/pair_lj_cut_coul_long_gpu.cpp | 4 ++--
src/GPU/pair_lj_cut_gpu.cpp | 2 +-
src/GPU/pair_lj_cut_tgpu.cpp | 2 +-
src/GPU/pair_lj_expand_gpu.cpp | 2 +-
src/GPU/pair_morse_gpu.cpp | 2 +-
src/GPU/pair_resquared_gpu.cpp | 6 +++---
16 files changed, 36 insertions(+), 23 deletions(-)
diff --git a/src/GPU/Install.sh b/src/GPU/Install.sh
index 32ee689bed..4f16599890 100644
--- a/src/GPU/Install.sh
+++ b/src/GPU/Install.sh
@@ -45,6 +45,11 @@ if (test $1 = 1) then
cp pair_lj_charmm_coul_long_gpu.h ..
fi
+ if (test -e ../pair_coul_long.cpp) then
+ cp pair_coul_long_gpu.cpp ..
+ cp pair_coul_long_gpu.h ..
+ fi
+
if (test -e ../pair_cg_cmm.cpp) then
cp pair_cg_cmm_gpu.cpp ..
cp pair_cg_cmm_gpu.h ..
@@ -112,6 +117,7 @@ elif (test $1 = 0) then
rm -f ../pair_lj_class2_coul_long_gpu.cpp
rm -f ../pair_lj_charmm_coul_long_gpu.cpp
rm -f ../pair_lj_cut_tgpu.cpp
+ rm -f ../pair_coul_long_gpu.cpp
rm -f ../pair_cg_cmm_gpu.cpp
rm -f ../pair_cg_cmm_coul_long_gpu.cpp
rm -f ../pair_cg_cmm_coul_msm.cpp
@@ -134,6 +140,7 @@ elif (test $1 = 0) then
rm -f ../pair_lj_class2_coul_long_gpu.h
rm -f ../pair_lj_charmm_coul_long_gpu.h
rm -f ../pair_lj_cut_tgpu.h
+ rm -f ../pair_coul_long_gpu.h
rm -f ../pair_cg_cmm_gpu.h
rm -f ../pair_cg_cmm_coul_long_gpu.h
rm -f ../pair_cg_cmm_coul_msm.h
diff --git a/src/GPU/Package.sh b/src/GPU/Package.sh
index 60aa6820ca..2eb20755e6 100644
--- a/src/GPU/Package.sh
+++ b/src/GPU/Package.sh
@@ -15,6 +15,12 @@ for file in *.cpp *.h; do
if (test $file = pair_lj_cut_coul_long_gpu.h -a ! -e ../pair_lj_cut_coul_long.cpp) then
continue
fi
+ if (test $file = pair_coul_long_gpu.cpp -a ! -e ../pair_coul_long.cpp) then
+ continue
+ fi
+ if (test $file = pair_coul_long_gpu.h -a ! -e ../pair_coul_long.cpp) then
+ continue
+ fi
if (test $file = pair_cg_cmm_gpu.cpp -a ! -e ../pair_cg_cmm.cpp) then
continue
fi
diff --git a/src/GPU/pair_cg_cmm_coul_long_gpu.cpp b/src/GPU/pair_cg_cmm_coul_long_gpu.cpp
index 153cb98a9e..d19b2d6512 100644
--- a/src/GPU/pair_cg_cmm_coul_long_gpu.cpp
+++ b/src/GPU/pair_cg_cmm_coul_long_gpu.cpp
@@ -144,9 +144,9 @@ void PairCGCMMCoulLongGPU::init_style()
cut_respa = NULL;
if (!atom->q_flag)
- error->all("Pair style cg/cmm/coul/long requires atom attribute q");
+ error->all("Pair style cg/cmm/coul/long/gpu requires atom attribute q");
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU cg/cmm pair style");
+ error->all("Cannot use newton pair with cg/cmm/coul/long/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_cg_cmm_coul_msm_gpu.cpp b/src/GPU/pair_cg_cmm_coul_msm_gpu.cpp
index e88144bea1..d946b739d5 100644
--- a/src/GPU/pair_cg_cmm_coul_msm_gpu.cpp
+++ b/src/GPU/pair_cg_cmm_coul_msm_gpu.cpp
@@ -136,7 +136,7 @@ void PairCGCMMCoulMSMGPU::init_style()
{
PairCGCMMCoulMSM::init_style();
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU cg/cmm pair style");
+ error->all("Cannot use newton pair with cg/cmm/coul/msm/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_gayberne_gpu.cpp b/src/GPU/pair_gayberne_gpu.cpp
index 6713242885..c137f0f67c 100644
--- a/src/GPU/pair_gayberne_gpu.cpp
+++ b/src/GPU/pair_gayberne_gpu.cpp
@@ -156,9 +156,9 @@ void PairGayBerneGPU::compute(int eflag, int vflag)
void PairGayBerneGPU::init_style()
{
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU Gay-Berne pair style");
+ error->all("Cannot use newton pair with gayberne/gpu pair style");
if (!atom->ellipsoid_flag)
- error->all("Pair gayberne requires atom style ellipsoid");
+ error->all("Pair gayberne/gpu requires atom style ellipsoid");
// per-type shape precalculations
// require that atom shapes are identical within each type
@@ -166,7 +166,7 @@ void PairGayBerneGPU::init_style()
for (int i = 1; i <= atom->ntypes; i++) {
if (!atom->shape_consistency(i,shape1[i][0],shape1[i][1],shape1[i][2]))
- error->all("Pair gayberne requires atoms with same type have same shape");
+ error->all("Pair gayberne/gpu requires atoms with same type have same shape");
if (shape1[i][0] == 0.0)
shape1[i][0] = shape1[i][1] = shape1[i][2] = 1.0;
shape2[i][0] = shape1[i][0]*shape1[i][0];
diff --git a/src/GPU/pair_lj96_cut_gpu.cpp b/src/GPU/pair_lj96_cut_gpu.cpp
index 108b108125..ce51b12a7b 100644
--- a/src/GPU/pair_lj96_cut_gpu.cpp
+++ b/src/GPU/pair_lj96_cut_gpu.cpp
@@ -128,7 +128,7 @@ void PairLJ96CutGPU::init_style()
cut_respa = NULL;
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ96 pair style");
+ error->all("Cannot use newton pair with lj96/cut/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_charmm_coul_long_gpu.cpp b/src/GPU/pair_lj_charmm_coul_long_gpu.cpp
index f66319d66d..fa48db0d9a 100644
--- a/src/GPU/pair_lj_charmm_coul_long_gpu.cpp
+++ b/src/GPU/pair_lj_charmm_coul_long_gpu.cpp
@@ -147,9 +147,9 @@ void PairLJCharmmCoulLongGPU::init_style()
cut_respa = NULL;
if (!atom->q_flag)
- error->all("Pair style lj/charmm/coul/long requires atom attribute q");
+ error->all("Pair style lj/charmm/coul/long/gpu requires atom attribute q");
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU CHARMM pair style");
+ error->all("Cannot use newton pair with lj/charmm/coul/long/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double cut;
diff --git a/src/GPU/pair_lj_class2_coul_long_gpu.cpp b/src/GPU/pair_lj_class2_coul_long_gpu.cpp
index 6441f05419..2c7ea41531 100644
--- a/src/GPU/pair_lj_class2_coul_long_gpu.cpp
+++ b/src/GPU/pair_lj_class2_coul_long_gpu.cpp
@@ -142,9 +142,9 @@ void PairLJClass2CoulLongGPU::compute(int eflag, int vflag)
void PairLJClass2CoulLongGPU::init_style()
{
if (!atom->q_flag)
- error->all("Pair style lj/class2/coul/long requires atom attribute q");
+ error->all("Pair style lj/class2/coul/long/gpu requires atom attribute q");
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/class2/coul/long/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_class2_gpu.cpp b/src/GPU/pair_lj_class2_gpu.cpp
index f351c7ab97..c45fb5bfdb 100644
--- a/src/GPU/pair_lj_class2_gpu.cpp
+++ b/src/GPU/pair_lj_class2_gpu.cpp
@@ -125,7 +125,7 @@ void PairLJClass2GPU::compute(int eflag, int vflag)
void PairLJClass2GPU::init_style()
{
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ96 pair style");
+ error->all("Cannot use newton pair with lj/class2/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_cut_coul_cut_gpu.cpp b/src/GPU/pair_lj_cut_coul_cut_gpu.cpp
index 791e47410e..cd19e81ca2 100644
--- a/src/GPU/pair_lj_cut_coul_cut_gpu.cpp
+++ b/src/GPU/pair_lj_cut_coul_cut_gpu.cpp
@@ -133,10 +133,10 @@ void PairLJCutCoulCutGPU::compute(int eflag, int vflag)
void PairLJCutCoulCutGPU::init_style()
{
if (!atom->q_flag)
- error->all("Pair style lj/cut/coul/cut requires atom attribute q");
+ error->all("Pair style lj/cut/coul/cut/gpu requires atom attribute q");
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/cut/coul/cut/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_cut_coul_long_gpu.cpp b/src/GPU/pair_lj_cut_coul_long_gpu.cpp
index ad0414e6bb..2b3a915f0e 100644
--- a/src/GPU/pair_lj_cut_coul_long_gpu.cpp
+++ b/src/GPU/pair_lj_cut_coul_long_gpu.cpp
@@ -145,9 +145,9 @@ void PairLJCutCoulLongGPU::init_style()
cut_respa = NULL;
if (!atom->q_flag)
- error->all("Pair style lj/cut/coul/cut requires atom attribute q");
+ error->all("Pair style lj/cut/coul/long/gpu requires atom attribute q");
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/cut/could/cut/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_cut_gpu.cpp b/src/GPU/pair_lj_cut_gpu.cpp
index c37ce7bd56..63a908d3e7 100644
--- a/src/GPU/pair_lj_cut_gpu.cpp
+++ b/src/GPU/pair_lj_cut_gpu.cpp
@@ -128,7 +128,7 @@ void PairLJCutGPU::init_style()
cut_respa = NULL;
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/cut/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_cut_tgpu.cpp b/src/GPU/pair_lj_cut_tgpu.cpp
index 2807fdaf60..043cf0e5bb 100644
--- a/src/GPU/pair_lj_cut_tgpu.cpp
+++ b/src/GPU/pair_lj_cut_tgpu.cpp
@@ -140,7 +140,7 @@ void PairLJCutTGPU::init_style()
cut_respa = NULL;
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/cut/tgpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_lj_expand_gpu.cpp b/src/GPU/pair_lj_expand_gpu.cpp
index fc0c8812a3..f45429a9db 100644
--- a/src/GPU/pair_lj_expand_gpu.cpp
+++ b/src/GPU/pair_lj_expand_gpu.cpp
@@ -127,7 +127,7 @@ void PairLJExpandGPU::compute(int eflag, int vflag)
void PairLJExpandGPU::init_style()
{
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU LJ pair style");
+ error->all("Cannot use newton pair with lj/expand/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_morse_gpu.cpp b/src/GPU/pair_morse_gpu.cpp
index a0d6be917b..ae5fcdfe64 100644
--- a/src/GPU/pair_morse_gpu.cpp
+++ b/src/GPU/pair_morse_gpu.cpp
@@ -125,7 +125,7 @@ void PairMorseGPU::compute(int eflag, int vflag)
void PairMorseGPU::init_style()
{
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU Morse pair style");
+ error->all("Cannot use newton pair with morse/gpu pair style");
// Repeat cutsq calculation because done after call to init_style
double maxcut = -1.0;
diff --git a/src/GPU/pair_resquared_gpu.cpp b/src/GPU/pair_resquared_gpu.cpp
index 3d9a5b9f9a..69b6060394 100644
--- a/src/GPU/pair_resquared_gpu.cpp
+++ b/src/GPU/pair_resquared_gpu.cpp
@@ -155,9 +155,9 @@ void PairRESquaredGPU::compute(int eflag, int vflag)
void PairRESquaredGPU::init_style()
{
if (force->newton_pair)
- error->all("Cannot use newton pair with GPU RESquared pair style");
+ error->all("Cannot use newton pair with resquared/gpu pair style");
if (!atom->ellipsoid_flag)
- error->all("Pair resquared requires atom style ellipsoid");
+ error->all("Pair resquared/gpu requires atom style ellipsoid");
// per-type shape precalculations
// require that atom shapes are identical within each type
@@ -165,7 +165,7 @@ void PairRESquaredGPU::init_style()
for (int i = 1; i <= atom->ntypes; i++) {
if (!atom->shape_consistency(i,shape1[i][0],shape1[i][1],shape1[i][2]))
- error->all("Pair gayberne requires atoms with same type have same shape");
+ error->all("Pair resquared/gpu requires atoms with same type have same shape");
if (setwell[i]) {
shape2[i][0] = shape1[i][0]*shape1[i][0];
shape2[i][1] = shape1[i][1]*shape1[i][1];
From 5414a66be0cd7905aa182233753e0931f561dbff Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 21:12:41 +0000
Subject: [PATCH 012/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6768
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/GPU/pair_coul_long_gpu.cpp | 282 +++++++++++++++++++++++++++++++++
src/GPU/pair_coul_long_gpu.h | 47 ++++++
2 files changed, 329 insertions(+)
create mode 100644 src/GPU/pair_coul_long_gpu.cpp
create mode 100644 src/GPU/pair_coul_long_gpu.h
diff --git a/src/GPU/pair_coul_long_gpu.cpp b/src/GPU/pair_coul_long_gpu.cpp
new file mode 100644
index 0000000000..468821dcec
--- /dev/null
+++ b/src/GPU/pair_coul_long_gpu.cpp
@@ -0,0 +1,282 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+/* ----------------------------------------------------------------------
+ Contributing author: Mike Brown (SNL)
+------------------------------------------------------------------------- */
+
+#include "lmptype.h"
+#include "math.h"
+#include "stdio.h"
+#include "stdlib.h"
+#include "pair_coul_long_gpu.h"
+#include "atom.h"
+#include "atom_vec.h"
+#include "comm.h"
+#include "force.h"
+#include "neighbor.h"
+#include "neigh_list.h"
+#include "integrate.h"
+#include "memory.h"
+#include "error.h"
+#include "neigh_request.h"
+#include "universe.h"
+#include "update.h"
+#include "domain.h"
+#include "string.h"
+#include "kspace.h"
+#include "gpu_extra.h"
+
+#define MIN(a,b) ((a) < (b) ? (a) : (b))
+#define MAX(a,b) ((a) > (b) ? (a) : (b))
+
+#define EWALD_F 1.12837917
+#define EWALD_P 0.3275911
+#define A1 0.254829592
+#define A2 -0.284496736
+#define A3 1.421413741
+#define A4 -1.453152027
+#define A5 1.061405429
+
+// External functions from cuda library for atom decomposition
+
+int cl_gpu_init(const int nlocal, const int nall, const int max_nbors,
+ const int maxspecial, const double cell_size, int &gpu_mode,
+ FILE *screen, double host_cut_coulsq, double *host_special_coul,
+ const double qqrd2e, const double g_ewald);
+void cl_gpu_clear();
+int ** cl_gpu_compute_n(const int ago, const int inum,
+ const int nall, double **host_x, int *host_type,
+ double *sublo, double *subhi, int *tag,
+ int **nspecial, int **special, const bool eflag,
+ const bool vflag, const bool eatom, const bool vatom,
+ int &host_start, int **ilist, int **jnum,
+ const double cpu_time, bool &success, double *host_q,
+ double *boxlo, double *prd);
+void cl_gpu_compute(const int ago, const int inum, const int nall,
+ double **host_x, int *host_type, int *ilist, int *numj,
+ int **firstneigh, const bool eflag, const bool vflag,
+ const bool eatom, const bool vatom, int &host_start,
+ const double cpu_time, bool &success, double *host_q,
+ const int nlocal, double *boxlo, double *prd);
+double cl_gpu_bytes();
+
+using namespace LAMMPS_NS;
+
+/* ---------------------------------------------------------------------- */
+
+PairCoulLongGPU::PairCoulLongGPU(LAMMPS *lmp) :
+ PairCoulLong(lmp), gpu_mode(GPU_PAIR)
+{
+ respa_enable = 0;
+ cpu_time = 0.0;
+}
+
+/* ----------------------------------------------------------------------
+ free all arrays
+------------------------------------------------------------------------- */
+
+PairCoulLongGPU::~PairCoulLongGPU()
+{
+ cl_gpu_clear();
+}
+
+/* ---------------------------------------------------------------------- */
+
+void PairCoulLongGPU::compute(int eflag, int vflag)
+{
+ if (eflag || vflag) ev_setup(eflag,vflag);
+ else evflag = vflag_fdotr = 0;
+
+ int nall = atom->nlocal + atom->nghost;
+ int inum, host_start;
+
+ bool success = true;
+ int *ilist, *numneigh, **firstneigh;
+ if (gpu_mode == GPU_NEIGH) {
+ inum = atom->nlocal;
+ firstneigh = cl_gpu_compute_n(neighbor->ago, inum, nall, atom->x,
+ atom->type, domain->sublo, domain->subhi,
+ atom->tag, atom->nspecial, atom->special,
+ eflag, vflag, eflag_atom, vflag_atom,
+ host_start, &ilist, &numneigh, cpu_time,
+ success, atom->q, domain->boxlo,
+ domain->prd);
+ } else {
+ inum = list->inum;
+ ilist = list->ilist;
+ numneigh = list->numneigh;
+ firstneigh = list->firstneigh;
+ cl_gpu_compute(neighbor->ago, inum, nall, atom->x, atom->type,
+ ilist, numneigh, firstneigh, eflag, vflag, eflag_atom,
+ vflag_atom, host_start, cpu_time, success, atom->q,
+ atom->nlocal, domain->boxlo, domain->prd);
+ }
+ if (!success)
+ error->one("Out of memory on GPGPU");
+
+ if (host_startq_flag)
+ error->all("Pair style coul/long/gpu requires atom attribute q");
+ if (force->newton_pair)
+ error->all("Cannot use newton pair with coul/long/gpu pair style");
+
+ // Repeat cutsq calculation because done after call to init_style
+ double cell_size = sqrt(cut_coul) + neighbor->skin;
+
+ cut_coulsq = cut_coul * cut_coul;
+
+ // insure use of KSpace long-range solver, set g_ewald
+
+ if (force->kspace == NULL)
+ error->all("Pair style is incompatible with KSpace style");
+ g_ewald = force->kspace->g_ewald;
+
+ // setup force tables
+
+ if (ncoultablebits) init_tables();
+
+ int maxspecial=0;
+ if (atom->molecular)
+ maxspecial=atom->maxspecial;
+ int success = cl_gpu_init(atom->nlocal, atom->nlocal+atom->nghost, 300,
+ maxspecial, cell_size, gpu_mode, screen, cut_coulsq,
+ force->special_coul, force->qqrd2e, g_ewald);
+
+ GPU_EXTRA::check_flag(success,error,world);
+
+ if (gpu_mode != GPU_NEIGH) {
+ int irequest = neighbor->request(this);
+ neighbor->requests[irequest]->half = 0;
+ neighbor->requests[irequest]->full = 1;
+ }
+}
+
+/* ---------------------------------------------------------------------- */
+
+double PairCoulLongGPU::memory_usage()
+{
+ double bytes = Pair::memory_usage();
+ return bytes + cl_gpu_bytes();
+}
+
+/* ---------------------------------------------------------------------- */
+
+void PairCoulLongGPU::cpu_compute(int start, int inum, int eflag,
+ int vflag, int *ilist, int *numneigh,
+ int **firstneigh)
+{
+ int i,j,ii,jj,jnum,itype,jtype,itable;
+ double qtmp,xtmp,ytmp,ztmp,delx,dely,delz,evdwl,ecoul,fpair;
+ double fraction,table;
+ double r,r2inv,r6inv,forcecoul,forcelj,factor_coul,factor_lj;
+ double grij,expm2,prefactor,t,erfc;
+ int *jlist;
+ double rsq;
+
+ evdwl = ecoul = 0.0;
+
+ double **x = atom->x;
+ double **f = atom->f;
+ double *q = atom->q;
+ int *type = atom->type;
+ double *special_coul = force->special_coul;
+ double qqrd2e = force->qqrd2e;
+
+ // loop over neighbors of my atoms
+
+ for (ii = start; ii < inum; ii++) {
+ i = ilist[ii];
+ qtmp = q[i];
+ xtmp = x[i][0];
+ ytmp = x[i][1];
+ ztmp = x[i][2];
+ itype = type[i];
+ jlist = firstneigh[i];
+ jnum = numneigh[i];
+
+ for (jj = 0; jj < jnum; jj++) {
+ j = jlist[jj];
+ factor_coul = special_coul[sbmask(j)];
+ j &= NEIGHMASK;
+
+ delx = xtmp - x[j][0];
+ dely = ytmp - x[j][1];
+ delz = ztmp - x[j][2];
+ rsq = delx*delx + dely*dely + delz*delz;
+ jtype = type[j];
+
+ r2inv = 1.0/rsq;
+
+ if (rsq < cut_coulsq) {
+ if (!ncoultablebits || rsq <= tabinnersq) {
+ r = sqrt(rsq);
+ grij = g_ewald * r;
+ expm2 = exp(-grij*grij);
+ t = 1.0 / (1.0 + EWALD_P*grij);
+ erfc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * expm2;
+ prefactor = qqrd2e * qtmp*q[j]/r;
+ forcecoul = prefactor * (erfc + EWALD_F*grij*expm2);
+ if (factor_coul < 1.0) forcecoul -= (1.0-factor_coul)*prefactor;
+ } else {
+ union_int_float_t rsq_lookup;
+ rsq_lookup.f = rsq;
+ itable = rsq_lookup.i & ncoulmask;
+ itable >>= ncoulshiftbits;
+ fraction = (rsq_lookup.f - rtable[itable]) * drtable[itable];
+ table = ftable[itable] + fraction*dftable[itable];
+ forcecoul = qtmp*q[j] * table;
+ if (factor_coul < 1.0) {
+ table = ctable[itable] + fraction*dctable[itable];
+ prefactor = qtmp*q[j] * table;
+ forcecoul -= (1.0-factor_coul)*prefactor;
+ }
+ }
+
+ fpair = forcecoul * r2inv;
+
+ f[i][0] += delx*fpair;
+ f[i][1] += dely*fpair;
+ f[i][2] += delz*fpair;
+
+ if (eflag) {
+ if (rsq < cut_coulsq) {
+ if (!ncoultablebits || rsq <= tabinnersq)
+ ecoul = prefactor*erfc;
+ else {
+ table = etable[itable] + fraction*detable[itable];
+ ecoul = qtmp*q[j] * table;
+ }
+ if (factor_coul < 1.0) ecoul -= (1.0-factor_coul)*prefactor;
+ } else ecoul = 0.0;
+ }
+
+ if (evflag) ev_tally_full(i,0.0,ecoul,fpair,delx,dely,delz);
+ }
+ }
+ }
+}
diff --git a/src/GPU/pair_coul_long_gpu.h b/src/GPU/pair_coul_long_gpu.h
new file mode 100644
index 0000000000..a84b281447
--- /dev/null
+++ b/src/GPU/pair_coul_long_gpu.h
@@ -0,0 +1,47 @@
+/* ----------------------------------------------------------------------
+ LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
+ http://lammps.sandia.gov, Sandia National Laboratories
+ Steve Plimpton, sjplimp@sandia.gov
+
+ Copyright (2003) Sandia Corporation. Under the terms of Contract
+ DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
+ certain rights in this software. This software is distributed under
+ the GNU General Public License.
+
+ See the README file in the top-level LAMMPS directory.
+------------------------------------------------------------------------- */
+
+#ifdef PAIR_CLASS
+
+PairStyle(coul/long/gpu,PairCoulLongGPU)
+
+#else
+
+#ifndef LMP_PAIR_COUL_LONG_GPU_H
+#define LMP_PAIR_COUL_LONG_GPU_H
+
+#include "pair_coul_long.h"
+
+namespace LAMMPS_NS {
+
+class PairCoulLongGPU : public PairCoulLong {
+ public:
+ PairCoulLongGPU(LAMMPS *lmp);
+ ~PairCoulLongGPU();
+ void cpu_compute(int, int, int, int, int *, int *, int **);
+ void compute(int, int);
+ void init_style();
+ double memory_usage();
+
+ enum { GPU_PAIR, GPU_NEIGH };
+
+ private:
+ int gpu_mode;
+ double cpu_time;
+ int *gpulist;
+};
+
+}
+#endif
+#endif
+
From a3c3214428cf26bf2cfe7cfa3b360464119dd526 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 21:12:47 +0000
Subject: [PATCH 013/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6769
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/KSPACE/pair_coul_long.h | 14 +++++++-------
1 file changed, 7 insertions(+), 7 deletions(-)
diff --git a/src/KSPACE/pair_coul_long.h b/src/KSPACE/pair_coul_long.h
index 2de7fa9d14..59e2c7609d 100644
--- a/src/KSPACE/pair_coul_long.h
+++ b/src/KSPACE/pair_coul_long.h
@@ -28,19 +28,19 @@ class PairCoulLong : public Pair {
public:
PairCoulLong(class LAMMPS *);
~PairCoulLong();
- void compute(int, int);
- void settings(int, char **);
+ virtual void compute(int, int);
+ virtual void settings(int, char **);
void coeff(int, char **);
- void init_style();
+ virtual void init_style();
double init_one(int, int);
void write_restart(FILE *);
void read_restart(FILE *);
- void write_restart_settings(FILE *);
- void read_restart_settings(FILE *);
- double single(int, int, int, int, double, double, double, double &);
+ virtual void write_restart_settings(FILE *);
+ virtual void read_restart_settings(FILE *);
+ virtual double single(int, int, int, int, double, double, double, double &);
void *extract(char *, int &);
- private:
+ protected:
double cut_coul,cut_coulsq;
double *cut_respa;
double g_ewald;
From 6ac01bf0d4e08e7361a34f149445b43ef6022258 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 23 Aug 2011 23:10:45 +0000
Subject: [PATCH 014/246] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6772
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/Developers.pdf | Bin 0 -> 86207 bytes
doc/Manual.html | 3 +++
doc/Manual.txt | 5 ++++-
doc/pair_coul.html | 5 ++---
4 files changed, 9 insertions(+), 4 deletions(-)
create mode 100644 doc/Developers.pdf
diff --git a/doc/Developers.pdf b/doc/Developers.pdf
new file mode 100644
index 0000000000000000000000000000000000000000..2f45f9b99e0dba14c245815483cf9f3df82430ee
GIT binary patch
literal 86207
zcmbTd1CV9ivNc+^ZQHhO+qP}nwyiF^x@_Bab=j`!*Wba7^S^WAy>tH+u_IQ@h|IP2
zj+wc~$T23Vf`}L`BOMDAGXVpEy^$3G4-XW*jH#WuivK30FDQC3OB)wcCjxpg
z8$%aU5mRG(6H_QYJ}74wCsRXPD31+W6>0k%285oQx;->vP;`C|+{MQqVo@k1I2(96
z1kh0Nv}|1yy*Dx7Y>`Qm8!_f5-X1^kC%41tiC{xcq{;=~5XLlaYZL^tOmwP9C@%Cb
zPnV~sZ7J4tF)C>c^9LmAPcW#)v9_KaSaA1XikQLZ2n;A+E*6$$U&}uM88j7ciBuz!
zBx_v}MTBmN2oWj^h}18kxM{rK0oa179!?J=#zM%cB3#^R={DgITTECRA&E1atT1_c
z8B;)=T{=xMbwH{G!Az21mI#>V92&^3(9c3YztXcSMp~=^Yyuyem=Q3`dMbJRBc9Ck
z1oQT+S%lR!mw-N5>6`sZzYVBJ?lH(HnZ0LEnPOaM)+=*qfKEWWdV5D@<4%AW&(2qn
z>7msQErzKXfz4pSCx2khr_S7dg;0p&iB(3dzp!G%WA8R)b@iG6vC$aS%htOAV>VzA
z=k4+kN!*YQ*9~c~)29l(
zdb=23se;DL>ET8cU(&bM+xTluYQ=4bUd8YPYQ6dOL4@`z0AK;yAn{mpW{GO!FX=hZ~S2#qB!r!$;{>~=S^Phh&5mSC6t(=
z(Ux{B<;47}j-D@|qjb-@y(=^C;s(KFJH|HS8R7-@moW=XkAU{Bd@}ITm@^y;CcKseOrE$>Gf%3
z*sVbaTxx!3)UY3+<`A@_Vu>WqIFaI-RxnBf2hmjwM;YqZ#1D4wb*|chDR0ft%2mih
zE3ax%PB&lX%At==@+N!k@v1CClaW#4FYN4@3A%T2B9yM&UafMILov2Q+j?L9_UsJEp34-_^455{VquE#A}98a5`hDh!D3b~;)
z5XE`6`KLggW0gz>O4r%Li$XH}YmJg_X~eeXNJBJ*5XGC2d&4b{Qjak3c!HoOLQq(z
zcShK99BXQ_Qg0WwALtkTQn-f87_|{{`q8(~C
zXuBF9fZF4iorz*0yd2myDKwL6RFzI3=~ud7npfhV&Ntei`_uaUCe~1kHh7yPx^>N?
z)wjTTrJ7B@B6@tFVQ<<#Fuc7HHESni9O0eL`3>Mis;K<^qXOW-qbl=Vhe)zP$|^=z1J7Ij4R~?iS@RoqdUYe
zYHA3#P?)|hb1A2I2?!L=x|Vxg0R^lqMba5@v5uEoE>vP5uMWcfxgu~8nhTzH1em9N
zz*$>NmVHA;wTm64>6`-+NM?Y-Y8fYZ30Thrm=4YbeFh8SJIEeUH(}w|x
z5K0j^b1p$=jsSI|pI|lQHSGXFdi^>+9oM(7J50;X7@9WGS)$FiZFQUn^}
zjI4fBH?eSk^*U>vpa8KY7F)TH9X{JGA78|0|NEt*?L?M@=~#mj_*>`&SFyh1aAWl4
zQeq`gu>n9d*$@%yNYOfc(o^VXUxo4J29;b&FXz;|pg`R7L(AMi2z8?S?y
zdb5Zm1+R@pu(WhrTN2G7vwE+dG~J={miF>%>=Dzf{Hl^zDX;VG=CY(sMhmr;*`Uho
z>EbhxOty$q3F#-&Zc7{w^{;r)C}g*ssR_BD&iX^As;ou0|t~IxGd%kt1F_`aaYGLn2}y4eypN71py$q?K?KSHRLTbr=B)aB#7a)
ztdmYJ=+m{W)3t1tHL@c~99MH5aPm_?i&0<9?trtkyrrP&Mdai#MT4iP*?iAiRK&Z+
zR@}Da9+sont|UafL+V$7%)!C}4r8xy`{I$_i0lNNWB}S`vH)c*)z%qbzp3L6=BmHU
zd2A)BD)qmR$%Qu+iW+0A*PPe0C>;G*F8=&*eQ#|#hnz0$q#aS^iIDj2+Dm2`d?g8{
zh-Rj5p)_3fhipkW<9s_b1opg!M9K}hJ1gge)DE5)_)L8kD)H5
z%YZKi>Q|>2>*OqANh%Syq8|
zOgwR_53!-z@Dlq4&{XX0xrp}K?x5$w6zXkHe5O^$>vmx`i;Ys=Jy`7dy!x?F53pGJ
zL4TR2wj+5T8kSyW2F>i3wg%fjE)Jwgs-C?{1p)_WjqMB!=ua;&pvQNmld5)?!uF{-
zd=?Cvj@gLuU=I4(DV8>}iTwMH9II)w`=@t@BBeMkKN+aqv^qYvnhZw+aLE~o1}Ou@
zX#14zPV^F&Cr8$e0f4LZGNys@)2bf3zg3WMQm)PEGk>0Ms2A&j*Pte=uqO*2(?Ryv
zF{8H?t}H*#N8Y{z@3?R7(4;HIX%EzJ+?-$e8@hj;{>8r`C`Q(Qfgk@7Lj3_b|8ES%
z^7obhHHi8v?w?TSKccAr9du&)3%mU@-#?%eGwUDl^|$N)pV|Dtnshp3yT7iRdlxAc
zV-XCf>*hrEwpPlVvbij2wQ_^QiEPy<6bUUc=GO210)mKYRMut&&-P%!iWTcEys@99
zt5xk(AX_5S-P)l|SuspIiDnre&>)ZAoGvHOt4)bY-n&u)m4*EM?$@!PiDYAarow>j
zVj7q1#Z{o{&@u{+a}gYh4*e>0+SagaqU^NH8>Ww
zc>#7epdy;j&XAP}0M9w*2SF%HQe3;YcVU_%4FAo`jHNSoNO@->6)HO5^9c-D1-_;`
zswtqo)Fo`}`3g4b%!M`61`jN^{9!c+{DjwQ>7cv&5CG-a%h4_r27oGoEbF6LivXDP
zXrg{InQOD9>aM!_U%HgMI#Hnd*@LUjH
z+N(Ilc4h&uen2^le9kp*KS3#$`eZ1q%L!QIvf|wJ8@G+P!laBvXDr~56iVDpmEk1J
z2F=@^cprxj9I5+zp?7^bmP8M=W+R=cf6r$Ww|fbQcjpjLnS{jYX?32I5hzk?oue}f
z;}1bC$Zp6UXvJ26uDiUfs7w&EVg|Dy$QGP5e-PbR=ny%g6eBVCbc|Yf4{UVp3`H2_
zm)GSu$0a*NtI+7Ib9s_3D-SV^sA=wY)PcRalI(T*{bu29RUY|EtV|h#3hurFcVK#9
z7~3Yttgfsa+R`S%@L0+qXCr0)ln>TIu2Vxco(}Uf3lA0sP<_4nIXDt23OEw
zYkvpXEpFtsv?u|z*kyv3^3uYkeR@i3&BMDaqg|ok!Xoa6ls39O)apQg&P}XhRaazq
zZFImzfM}veilM7qlnpGuRdlb3dE4G(6wt3+4fZ<$0_6A#8_e=>yaGcrO<=7r^N<08
z#+ogPKev-@rbiPBU~v6P79P`vLKq97r>I%*y#fr{JEg(U_1@*A;f-Gel!jq4C=S;v
z;Pnna1X+p&gn>zfWObT*hu*)+GnJ#O=rTr+yGFps1~#V01d`ugwgK!?_Cvn|Y>}v<
zgrUaO#5}8igvk@SMi}6fioZ!1Y5Ty0a
z*x*ZlIWL=J^!93)&Un*3bVZCiJ_5j++tp^q;5vett1%7sgoj_$VkEoVwvFOq4Jk>#
z5yF4P_iOl^jf2v0b^D<>2g?WX%GPNe8+RW%u3TZP#vF|5>fTS=j^Yng2iN9>KwkGm
zvWklW6yW4>aqiyRh_)`-rVPu_297L*f=ld>Oj{5k!I$wB35%pNnHIO`f|Z8Hf(SH3
zR{@C?>803T6Nfy7k*H(@Rq{j>YXI{~h;wg3p@c_9o}>xguMZ9%k4e}F5u*nD8ihe*
z_jvw=tn~Gk{2oIiZET=Pdkr?X>HQvo5fOddf-)jdkOFUmC1U8*DVEDSq)_n7;!rA
zsq75ZOJPW*#QtGxn#_&J=jHs@lwlC`*Dy>llLQ0oca^m*#$y}XV8y%{h7qCOOd2mm
zXKeTgGRIH$`{Wdy6eibt44QZk4T+Gw?p%iK-31K5Z#^RcX-TLB22;2)
zL_6Tf^y?pgItE9yuKK||(brY*bP+&t6cJdmw+J?=GIAnWdeCFNbiZe(7PHweav-W!
zK%cygEapcni92P|ZxK3BocQSWdSGTKrh{P+~vmIw|>2ai`(9X1C+}Sh(HA$wV}F_B4xP8xj~ejRq3@}
zMLVDkw3A|(M4R-XWiYLMLkFk3we|U>jmNN&=Qqf1#~g1i<387aifsn0A+o4mW@ZQH+(3&++Up630p2v_
zM#f1oh`uR92ZRleQblKQA=vL>IEsFl%LhAl4;t{=UHcMm#8jNe-sO!E&S`uOnD%RJ
zOieeXrzntWdAi$qh6X6KOG2oM80Bw5;))q&(NEWx@lEpi#lt5wTAvDz``wB{-6Dva45YxPpES&E7X=u_t?9`Fe2M(vY#
z=>{ijY}6w}c4TxbZ$dJL@lT*OIPLkOIav(xQ>&CQds?1jW>~;r%H<XhjeEHo3%h`ezU0XK1sqv$jCoWSJ4Qz!5}!^p>^-D>0&m6Sg-G}
zxz;AYgoSP^WR)$kl%Glb34;b0NaAcr)TW3VaxSio8PYsXicWh#4C=GsIFqq$2PnY9
zkwQK1QjG>H)bMIEm(FNn76xqHrL4->*jWQ#u1)QY-aQB%aJQo$mT5E*BP>f!l0Z{_
z5@>cn@^r;(Wtq%qS;eM)`t6~qCp3D#!aMuD>7OoI&-=+0G5;21c)ePm>__BCB4a@KE
zhi*6f=J^oNwA+HMDRprF?vs_
zL6j?lv(rk&u7W+*Eb{ev#0xojmkUqMNu2P`%l4V|<7eBV`z(a`X*lnPgKGi|;;Otz
z^v-)G1O{Vhf)
z0m#e~P=TQgp~-~5d1x(utcpq`0I-+n4bocY5cJQl+iwk{SURCwF$xLAvR)|z(B+9i
zSDLz!Awx5#wrz^KnF;|DP;Hrv4@ApepO;CQ%_N>Fl?H_lBHC9fTVx$-_;s00B~&wm
z5ZbX6FR1XbY2Q8(Pc!U;c8Jt(wk0i{Iz{LP=C$}sRxAIG_f>8px07}il&sh61gZ+X
zVijuKtO}3Vg$wuJne!PZhDhYQVG{jMRH>Oz$ff$lNI1Y&M|Eu!pz6%3z<2Y%*bCJa
zy2WZG@5F8C8Bzh2;b2N2_)Xk-o8KZBM^_036WYE#2)YRx&{A)y_IEYy$t
zl#s}kL^IAmOZ90mBc#F`?Jj4dV?I?RorfKi7o+<}cro@4K}BN{NaFZo8H(i>@7KU!
zN%?EeOv)E}9GYc?`>&0E&%R-+ijYtj+gvevKjlUVtMt9+jjgvltHW))&eEieU9tC2
z6K-I=fRJ}ItXzW=kqS)~_MdJ~GIF>H_tLkIu0g{~+rOS+ez-t-xT(bCcsF*
z^fw{lFMWW8fa&jA!GEO+F#YRy`ll+u_TQ@lKQJegHrej^^b3hqS&Kx0d~io4b&hMg
zk4?zDC7l9DbcbT~GJd=fl+^bye*X%jegAsb@@d
zNiM18qB&JIu5G_r-)*x|E_QGHeD11p06gnS1G>V!x)m)^zZN#)FPdY{{rXT_V%=gf
zSF%a4Fak!NIBriYb4B8R-dT8M#JDx%`e*~5U7Mn>yxVn4Mq7wuV3@B$Z_P1oy&8z&
z;U>l?Gi&`Ustn$A1a*r(Ywj2CM`n{mQpfCVVRR`WhFZv}W2TjBJmXaq7>k1V#oKc0
z9P8}Sv`)AJ4t8}EVaVdb%=f~}1cpbY@bZ(FuTO3IaI}sOg-v!5q|2=4CFZQ>0yN;X
zTY=5r02lw@Nz59y4zMaeArK0p!qY+I=a}6laR+VPHJ28DmK%IDBeie^8#bOT(rcN&
zv~??&*Y-hZd7^m*l=1(;=Nepe&^}Q4$up6&vYho!9z#kZ0J$M)wqpo1$
zh=r^-*x~L762V&=%q6y~1ym&?+u%xO<-PTit1&P=C46PSv=B=d&~30i5=~^Spki3N
zflapo@3BgROKA*`LVNN1Bn%bw0Nx11doSW#fTg24URH%BdSxZpfgjNE4Kl~c2_s)q
z#ln6R39I=&ehQ7)NzYSr@FPNs3zY*XQ@dk)7eEBk(;i?Q1Ka_vmOE2*iVI%>M|9%H
z+JaJA)me<`$N~4i9O3T8HBVs#D@dxBnD+ayfg~1J3q9P-ug!(e7463(dn4J3dDocU
z&2z%EPZH2Y`vKYumjPB|paTA6c+E=;G<*~i5IY(qpVio#9B`)bZgv%di-`zCw`hmJ
zERlxA*kho~a8_=Y(z#E1c#i{z^@m{XR651z_fAobT=^K{kO)UynrN{hJd30CrU>c0
z>Bl%p@ozMNZ*}#vXw2)QEilrr=wqed#C-a=a8@}O5+Gi~xacNToKqwIw!h&WknoGQ
z6W++5P8i#?{86MnM8?f?%Tf=TqxE77ZbXz3K1Yi2gaE7(Qxg#>67fo_&{EVX2n$#q=(;4G9bX*A>vZ#gfVQDd%Z7&
z;XTw+ijyT_cs1`=**iS`^+O`1`i#wEWF{+2T~}m+b+gctZ)~Y+Uq{+$u0sOoRIK55
z?Jqd_pCaGf=IY@2#;Z%*=~ygfr_1}629Ef+c;5ruEXDcK6qpp}l_XM-M_I_k%Qu`)
zme$@L0tcZ}d{gy%*k<&o71+^~W~m;$ObO}k_o)<=>9RoK1Txr3s7T_B)FAxSk*HSv
zN_h()P=!T3Y>(&F_}8^a0b8|6s!GQVR8LZ;1k$RtZll%wwSHf|T^stg>il}P#;U2+
z@(ep8lBz%A)Gc*VQ>U`3H>>k`s$vIQAQQTWwm9Z6IMAwzo?Zx3y`4x;w3gV$)afE*
zTjCap{LzeXYx#-UmJ@nU*D)J-&67nO`StWO)b;t<+`!|^<|`&RCch&`l{;hX6q;_;
zRH#a;n_l5R@~Udc{&US;Ah!kX_9(@@o)oV?8Hn*k@H%B0^#zSk-i4#bIDaZ
zLpn|^pG)jyCNmS*d*_fy^N|1>SJ4(o)KcciC+Ef=2@23Jg{KWIK-gHlF`i~S8YDt4U
zTOBUgY@7vW*Zot-jr2;1_mB>hDX9BeL!7xnU}YexGGcoG0urM_t%RQFFancT=C$Vl
zo>9j`ZAgt~0ggX^J+_hv9EBC~zz)5@7yfDWR7WcA_s|H*39fym`IevHOvP}Cv%F||
z&G*~=t}2Cc{a)XB;=%McjBpo0Q+G|xU(L6&AA0cBOs4SbB{smbzOdJmN
zijq(8h-$jfeq9fx0@3)^(^1YXU=KVnr)G9RrQADLS!I>YyDtdrYZ??-R-JSs!=VG5
zEkf@efi}GMm9aydo^8DMhvG6OFC{MXGsXJ`GuNe5lEdc6gq0Sv-^GgWvp{N7{_
zll81H>AJ*Hhxd(
z6S4`)4q{5i=QNp*#pp6mog#=rx+HC$13wkW*Kv^dP1I5)K#bT@D0Qa8IBA?f=Y?=HUkSANd~H^O+6
zmU*7YuDkeoY}uO|W%MNSP3t1h(`V7sgTL$H9mE4w?}ZhvLx(Tj)OY^CeIv?>XH>SnkgkJYPRXkim8wge)}ni
z8MSvA1n*YC%W*|iUn}Y}-I2ge!PtE)_sDOr2te0IT~__G!ec_%nQdCK&)}Qj&Otl4
zurQp2?_cZqB(v*_B1w8b3aN;V`H5ZhPCAyaDFU7LT;vw$lSU)~Zh%98p8d{2Ky2?S
z$E$r#tqN9D`4V5_xZGRwxAs)@>6Tsu!;2O!PpN5t84WXbTC>-lw!t6OP@!@%jN`E7
z=Rm^7@h)HbrYN4#a%I&^$Dc{GI~+{Iy0Pr^bR8duh4`F>jJP!@&gYcvu&Mpd@~TrJ
zQKlaL@(*^4sb3;dkiZj#;5jhPE{6h$vGLl4I^G8njp=}tRw4!1%|sRzKqwZ<@9
zLrn;K@{Z?4m?tW^YF1>t1;tq>@3-LatWSe*YVegni6J%jV}v;`F)Q|aPsHz?Z_HA~
zB-857y|tTbf%9uLK6(MxRG>p5(%=kuVt)aENz=e9V1-1ikMJh@8G=29i8Q6G7
zn8IR^@1>y7q=9D==ixBA#UOB!5EiMOy&aQH(m%;(NHM)i#~oW|rP92X$Y)4M8bse#
zHX^m1)u6;~2Ztl?g4U~OMz1GZ!-aO_=K2yeTvpKxkNBAj!NVfk=RN0lWE?Ln!beqC
zSqTQt0S*rzwFMvJbp|!($p}DG*vawD;=z~Ce3DnJ{06Vmih460V!I3yncj+UAn;J`?
z`HD>ylt|*49R?quM(@<#m;kjb`cP=iEX6xCf<@FqKS^NeF7NT}>uMrFm9KP$V8&0s
zNVv7PGh+BmEW-GOjil)n8_)I`-(;C_Pk6WjQ@`fJQ1bc$Z##98L}yMA%I!507XJY`
zU%K@AvP!L^-PYd6r%O3)W7)Bgx7;or_zEmVbgVQ}pZcMIoHKGZxDoM=(UK1af^O=;G7f%!HT4Hq8!e^7dX#SSxNSmqO@FQ=N-HB_d@G)H3&feh8u~@u5wpW%%7EslMdh^W_R1W^u(egNgU9CEK6Dd>x^%|!<
zrkH$C)l)&NT~#HdmMRulneMuphDU8xh5om;#r&`KzyH{_{?;4*j|l_wU*hzimH(|s
z{-Q9eT&~N
zWyLScB%4I;OJOwoV;euOXPXLjWM|htb#>=_&bw#(oaGtn8)i-ie_tQ;@FyOWu9~za`#W#RT5y3MbIMw-kAoZ~2XsEGu
z7M^&4+Xm|BxeK=0jsnZp&cK4XRIA=^u-|-+lXLg(
zJc66GLC04EE$iIZ>VKSune5?c>f&J(W)Vy@VM7KL$!d?x-Ace
zxVfG$i$f2YY+d=Y!-uAN%e)7sE<7Hi!gvnU^g4J6@f_d8P30Tz!vwOX%8%COj&I1+
zWQfCsFZO1?Y7455**3Tr0)5LpuYIfJPeho2&}Vr0e&^W>{gJI+X+3hC=f`oKawQK#
z@;Ds9r!!gFa2>1B%LZz-+6gz9YX(gA(bCO}llR{VMkRo2F~GKe_Vy|=Fs)U@N-61|
z{vw-k=8wCV_5Atf^)Q;>Z|06{LqgcaAd-dtRwLgAc5P7~d}
zq6b0=Ni8OZpW1;+AJWRN%m51t>xH<5$k@C^+4Q0ZgO#F7bu!;tpI43g%^(
zt`GEeIp3g*S3Q~gupcg2X_H-UObUs4
z!rOyAT0P-K=0LJtb<%Wl1HmF%c0PP4W>r%Fm?S4IWAET>N8E*4uFJU62UTB4tmei7
z3kV)lime^ZKsw->fY)UoqJ21-2E)H@LnX*s5Zu8?LiJF_
zTjRHs9a39GPgc7bjr5MbV6zPwWxYX|_0_E7P7OPY1Gg#dQBM-;z+fw4FIBb_xb3Dd
z;WlA`>tmEm$_)3pb7%GfQ%upp@y>apJc98rn%tNy5^!-$>Ql8eLEi#_!v|So7})dp
zVgZIy#pi7>gz*(D>pigAmsNmQ*Y$-J8!ke~_DaqyMK{ORt%c(;OsQVP&Lzcm5^xZ7
zdqEro7eV%pXh`BRaX*3G3u;(wB=^*kS?>AZ8c7_}Y$x5^2pSk-=dX&5NfeF-hOU)^
zh4WVIG=5n2KJ{G#qz9a{@(0wOBiT}pb^?ACt^Yx+J)B9p%6}|8-eQL3JM`f1k3^SzE_vqp!%KqM
z!5B<18)RTo7w~T$SRY3!0=@4^EFvRV~1=h`7YW?;*m@
z@RxkwN<)i<#XXe%2gGc0WgRr{y{0ptJo#lpci3)fm=M4(ukNW3G6lM}uEv
zYk?)Z+y}uDxHY;gacH6IG<000{3dHDaERz0fTMp#NOBQ2KA4g@JDym>DmoaHA^d@Y
z9`J_+&OYZE)ePJxZHO#~lusY2A-fp|w;^ksHkbmnIj8C_*kiq=*d3&dcMSt4lB3XR
z_N1|eacrNE@yynJvsa)FW9u@4Yh3qwD%&V3(G&`{I%;_*ztikYtu|J?5=#EOWB|lPL-1@&=z6B|P&2Ef(frb`rT$!NJxFCF
zMIrKCC@v^F5M2W9>+sI{C6HPE)OD#SY;!)?`X9$MKJ^z@G*xpeTtlcREE4~?a@S70
z4b`PDYP4aBQV6yPAWuNJZ^8#VVJpf@mRWe5V0wTIBmuV@w}mdp>g*-LPKBz
zl5KiOVWdAp?_6N$JM=|X>CU>=XSrl2vjcx|z*x*|pgtQ8D<`wvewW2LZg3^Z+zz-3
z>VX3q^S=G?oz2&7?iQukd7~03ru;K}GXlU3Sqgj#pf@8_dXAho4_ZnYivT<&-!dpXB0`>LernuIu8Q21JT-Wc
zgX~4X$V9h-Lb&0Ad=9uJCHGi=S&0||g!PZI0_Y78K?GLwba$o&qW)X@LJXZ^{38Pa
zfB+KDSE@ln3tL>H^e|!b0=G=vYjHQ>YS*qp$`mqQxWTRXAWj!7ZyES+d_w!2>D$4D
zPcXl%6YSe>zp?@KEd8<;*9h8a`ri8t!W2U*vXngs1h*a;I3w@#cj{zaLKBozIDnz7
zX{mmaVr+r*v~!7EJuf>V6i|J#l@^}U0m#NTQO)@#ZD$0gkyg00#JOw%g3#_zkj?Z0
zxC1<%4}ocYacJxi0Olf=0a8b2N2bJX3{(>36toRIM|HGAtiG9|Bar3aEb@{WszzW^
z=i?(J^E+MSm{3>2_AjWkmykh#eHYFHd15(5%+y#o8
zh~5YflZ+&NQln$gpIX3p4oMngGW!#|H=EuQJThMl)5kMr#(-frmy^z~y{pRra~z4BtdW
zdaBDzfAl{-r|bE$$e`B+B`HP&lyF3*bF^R{p|}UY+GT}aI1;q-y2}_riOa+;*E%RQ
z=xaRxfNQ}&T_2#NKMRak=s
z$PhzYvg70sBYSZ=hcnT?a|=G+ZNCr;hPB5>+-VIp2R<*
zz#|oNN(n6hVy!7)Ej^$k9U?8t#o!~$5~Sm0QRQjQ4R_**UOfxOWqhoH=@7|!-^51_
z(NFXEAZm+<7~SLLC*XW)Ps=$g+I0=+As^
zVC)qd{t+bb?HGhbrM!w*H&|??#h}*ubr1zvSpahv-!i6~`=gT<-l*ZFGO6DUb#A5j
zr%NGb%JiH(qY1Y<5>gYs+fU!F6B4FAw0$wakpPa5RJkC_Kcxjto;Y}NX;n=Zv_opM
z^1ip{IIU=ZMbw^)`qPO=kB3=Yjs|%e22+LfI%hBYRbx-D0?@^B%_3_74SBHTH_!kK
zlE}Yvv41rY|7R}t7vcFox!B*2bN%nR*k5t~&c*)a*w=rgWq&9dCjs-{6z2c#3>eG5
z^2>ixH1{JQde
z+4}|Kbm#GG^v~(4w+s?JP98D6Mg2YhA*wc!T>lWW5#l!)JSDdu9c{rB`S65}ItXdWMDb4H@
z9yW^+eeBnnVhfF$-wp&{bYoTaZmx`Vuk^sXaI81IkGy)d>4f66nr;@q_0wT8f;#iP
zKM}NM1X8!L;|1^6Jm#QU=jYRGJRiiVzm<{ILFfg_QDt>Kg69tht-V%%2_Vd?t*laG;Q
zblk8SdX4&}g3oyOUH{+6KfhtR)KGRHKUN0}K(vk84u-zqJu-slC>vU?z!yRfX&qBA
z%%Io65N+PC=jp)QQOIdbQ;P4gr*JH!mOOndj5=}1_+T)^#1)kAk7K7HS7%+Bs4M)?
zEN`~Q4yn6XKMFm|V$r7xG%L2SW-h3|k+qwH8^~xoGZwqyV5Jc8#Ya)Z*2Pl_!;bJ%
z;aJ2kLdMgHmg$-(87DQsd4>6ab(+vn2c$NjD>yJ|3(&I?*z_MQht?NtII)pAZWk44
zeQZsNyd9)+CA99w;}b>a?Dl#PUACEj%%5p^ri2b0T|BgCB-hR9qeis?WGmpR;5FFc
zchH3u%ieE!i$f%cLu)msw#3w$!IY4?RYY?PB+znyz{%{t$?+hS9`W?*W&aw+1=L<;
z{YKNqUJ0ZzPr6Ey0lY`EXMzt)yFut`fk!m(Ekows{cVXKP{;B3N&@i5_dR{)xa@2!{T1nRc@M>#B~c0PHM_!00gnLeHShXLxEiO3~lUD@S!
zHPnrKM{Rz`4ihYT(u0G3b3PGaXB7}W=QE?+&@JMSDvJw)g-oj%738w4B-M_3AoAan
z^}LCp;HAjHlnBn%W(3!w74DG}K`g8G%FKnw+(T+sX{CaC6x_RC#vQ6nHERI&fv9&?
z1#8RrVE8e*FIT=?hRnEyOup>CZU1KAz=53$O$ZHnRD1y
z1u0A3Xdqq2)?~g)TxJNqlY!fvrY3wyEiNhZ=492Ty<9*e)Pnm7cCQs7AyZFta96~b
z({Trxvtl-+4A-0%xmw(sR^&T2Wuv)6Z2E^$oE8Tw<^u*-6AK?-hp?Lc)Qpa&r5
zdhX)2(=#o|Q;2GwU&UT0OYcs%njX!7Jk|#hkAGOmAzDpT^2cIR&GzL4H{bkjXK`yO
zA?*kz54GkgIOfraH9EqxWZN>aInf|-2k2ygyGFH4>+6TDp}+;T1Y*?F1L3sz);1WP
zyT>(Zt1oJXQ?e}|oL+lzp&$nr3J6Uev$k^!bz47%ozO>qDMyZ0#t9Y78;;H$>61no
zlP^wzZx?taWTcJxhSTYeNkI_ea^@gOp9@VWiNGl=LGSwj3&jZ)3kP^-iGW>DXys|`
zWvo1!o)lRP1D=yBYY?X>>#PR%6!vZ*ROvehHD9JOjK|m)!;z*L{oJd7e1|B3I~;i;3d6&N@&jjf1z{frZXKDc4MuKAK_y
zPjiGr4h(F}N$&g7>@!ZRkc`f(X$7sJE~`Z5sb);D!Hwa7tq5mHg21;&7HHn9r#b@y`M24#`
zR?CFz6UP&B$`bGW&dN6)N%N*Cl0Jo40I0h3w5ofuBUhM{#b}_A(Hdm6Lqsh>E_Ff>
z<}HU=^c2Nbrd+MTMMTGG8-pq9@4=4OD+x@e0PkyhN{bdM*`Ggi;+Z=o{S>A>4NTW&
zdq&pa^JBNTvOZ^;^qOn9~O1Vo0QhQQW>H&~?EM~>E0X2x*a0K{S%cOrPK
z-DK-ic01W|HZTN&-I56h6{ks3Nl?^hI0#zepeSDjzBnAwo%|fQoo}m6bYo%YZJDZe
zB`|dLW0Hc&b@uggNiNhJb7}DEy?Rm8Gl#MHJ1@=jb7k3LaJ26L%GAt&Oqr`2xym1F
z5IG{S#^ZI&j&R;<=&=KSg?aB&r9oW1LKhN)Ak~^e;s(I9AlJC_xXBT$jsS*gr2s#p
zpP~M=c1|{cCPJEX1lu;F>yRxno#Tdnm*kP#;+Mx&*K(9F-MOE-o_L
zYwa0yu|H$XDa0SXNok9RN`tA<$N;Ot%?#Zf`Tat@(UTMsE!nGKWkG!M+7Du
z=|r#}h?dT_l^031FK8pIT$)mjVF$bC(C4Dvd$G;ETwHuB2vcfr*v^Gj*nso}MDonz
zwh|Bm>%eP#oarQM`_QBZIB;M8a7IZI1t^+KH!N58oqB7|Q^5ps7`oGg*bahrm#}pp
z7UUV>Kz1?IA=t&Mx#pqUFl7+u*8q%dr;n5)N!8K-gAIxN*k^|Gp}MpC5Sgk>MmC;S
z?%{=*h&UU!;yAUgS6?y{Mg^d{U)ytG&DzC)6bLdbk_^x;4mdIQ+DHd9g;`T?F}n})
zxgl{Zuz7wY06&F6N)D)X_qe9Wj(Ou6G3(>EZoW>O5rRoEb|o}b@>3y#2Z?FOFdThH
z{ic!EVKn2$ag){_?)vsC;+SnDkl=Qpp_BvTJ&niDqQT3a&Pss9aL4CR{TSqJ?w8mQ
zpCe@P1F?wYyhvabXN5@Xd|Gck8x7$S8zd^R@2wE1?hZy2&bMqhcjax50tYdSL#5kf
zk~Yf*5$ySa1CHzl&hQC!in()>UZkfi{VOG)-i?M`+<4P^xur12wa&8*ZatIQn*OmL
z7n#+tiz52rE?X0KDnH@GYX~QniSpg6u~i0|NkV*gCyr9DX58Vg#i2Wpu?K%TK`!Kf
z!Gdqr0y---(ZMIB97RJnJgc(!xRe6L-6`%cxVuxJcyTE1?hZwYYk}eprMSEM9ccTW_mlIS@7(A9b9vaace0a}
zm6c2;YbE(DIBwZ>7+j=4L&hN?tGV*@@q6Z1Zl6fpoiG$zZL_MWOmpjeg-Tk|b%2sZ
z=!d=GP2hK*eh#WGA?z+R0UAr44!(_~z-!20%DBdl{tzj?hCH{>o+`oqbbTU;+2-D`
z$oY}!&0_b90@vd@T2jFQ=b%NLc}~cm%E76JGIm(18l8n6EipO#c*QC$w=br(%8DzM
zP7JsyU;By6w($iC{*?T>uRr%Z?q94@iDj5HT2PWMyn~ueB69Sc1u#
zyx}HtFgHFO=Zq=kTAd?ZpEIv%gc4;fsM39F%qv|C;YIR93UC6dOMM&fff?5uTa-1Q
z{aYrPUWbH1CdBC2SbPGlNJ;Q5NP9iB#blsoli}uNUtJ)?Z6T{!Ri$^m%aA*{DxE>z
zqI^YXOP$YFlE-gY6yCbJAjNCKhbaizSpdU`k5g)zXSrsEnn%)r8
zvG1eioic{vqD4eJom4)$)@-A0u=SemJ%#CN;lYLU0EVI7^0ftoqA%fUxct0yG`eUh
zB(k|}RC*YNB0z_PpJ&!gsA+aE8H4%|s~M>tZh2~>oH^<$OqPHgJXv^mgL(KWGMY9P
z!Ic?r)B|wXWxdxsJFF+Q5!sSxcksB)utGAbS&x4@3hiyPQMfT-&*kkvnh}lZBQaX>
zia%|$OUsCE)*CDzVpjiO6n56XpVIxe3j4GA`;Q9y?_D7O!{RCHGtSE2z*IoHG5(#x
z{%^JP|7agJ>%Ut(W%^(2!&XE-Qy+8JmO7wDR$Qh!-IW2lEy{9b=t
z&a|m<@gDPTqKtz=+;6!e+z)MDZN!cRxdrL1i>9>_jswpb%6m%MM)x~uT_81CQGjwP
z^bJ?X9>;E2SXM@Xd7474c^s7BES82S72H9Smqwb)F85sU6bk1#*rRAVaMvZ+FN?5f
z=q6%5T_t(uV*}F>ZeYVYsinCQ$RteR)8K()Ii@?$*2_ZROR<;0Z+VjO<`f#mrCc
z{aI%W#2R;sb;FOAui$K^tV1%*zwD+#*5vvxoW_00Ps`Lmu-p>K^3SHWkAipwNwB-o
zJD===(e8cxE#V^rEfc%atl9j+o$Kj5;i#E2dZ|}Sut2JDF^q|*Udm?ihKRrG=K&SZ
z)9Vesw6UJ3xlc7ebP9r)Tnl1to20wM?e2K0UfzYOUA(Ewtn8P%Dgh@e3>W&SfIOW{
zR@E+>!Sn{#5ozA?nlXb!_x6&ijw+GITbc5KC-}{gI)&7@fo0=`^_^|Kvg?tHxOdio
zRWfAl3+|*W3?#;ktnq;uPd#1&mZa1Oj2b_Bc6D=IsJI$>=euCC!b(c86uebVNH^b)
zwXff!&D9vlH+_}o-qojaCf@&=v6ybtnZ$y36a-5LR(NA;M`Ucb3P8
zMD=M98l}&(R;yDS<*jyIH-1@yLT@oA?78RSvWxubS^0aNYgtJ%zopB!bLes85Hedq
z-f4@uKD4t@EtWMyqXApf+;FKv?R5q9#$0g7R`R@oRh%pGgbgT{0BB+XXqd37IE<~n
znxUSk;Gm@z3`_ocdjbjT7p#@kft>~GGml2{KwWSfeT7hE>oKw!M+k4i+flji2Udok
z@aUwBQ2a{ihsoYAzAu>T{%H+qtUnz-&*Jz3Vo(928H=IFNh3W9m{qw6Rq~B9)xwB_
z)YY@-K9id`TX%Yp$o?fcSc8<+O#xa1i>Pes>-VzT{qSG(MkJ(RW{{=)gnY?*7=5~m
zj?vzko94}ICt5Ub5xU}APL4*=z!{jCW7;9Z1MXoYMW?x^`Kv7b@Z#vEy8uaXNa#Vmt^+dskX`Rn>VYr
z3D=Un5PPMMs^|y5=&Z<}CRCX{YHJdSh%@8g+eS?)_Kv^9HLR|kK2grFKfPuJ>`ds0
zs-Kf?dPn8MY^r$zvTswdYTQU+2o-)HvSnmI
zRa+ZzbyCOMkg-Lq2-ZG-TI)K_cr6H4uWPBRP8ip%EGMlL83cyp``0}@2xGR5mX3C>
z<4Y(AEtoErm11EYoahCpZ2GZ|;lz*L?CKA->7X0r47aR7FTESjQ0y(XpcRccjHQU^
zX@M+}*$fa{bpVc_#87?siFNP}M!HoSgI0uRrXTSqwVU!nw&HFuX-6==g2xef^bouf
z3m1-=MoLB9GB(UB5smw8906hWiAuFu)j-oCoq)`bQ$_-N@Iu0B1*H3z27DbE28cv`
zSK#U-%f%BPF{9>$ZeHbmO3IK!+8I`I?J_4p(dHDN%IFa^$Wan6wzsK!A;`Z@q-nv4
zH`w;#yOEB#yN!$9$BnQu@a9LCAqzHj+(k4;r(SaF?oS!kj8S-+pXx0<_vY<6b}^EG
zY%32lQh8OVMwO6HII_=%GmVZ^^V$@C-e?gSB
zOqsFBjY-tNI9HF>Mz!4GUeeygbu5D{>l;zHKpsztRhzesJ&t+AyfcEJ9dmKUfd*^yJt^P`?TNEcl!_8D^MGrl?o%^RI^sOX1YE7T@(
z>m0Dj14_E&X>Y#y4nMkGA~-yXrIW1AEaOyB{P0sqZ(VS2pF?Ch~U>F#vaMZ=$%yrK=F=
zp{_00vkx(-oHBU^((K(ADkwBN!P5l+qiVsrwq?ceR5HCS*JMGrJKy+U-6dU(<)-)r
z299nV8}{nmV>J&W`qS>O^#=PZd}@!2ZbXqO!#QSKe<|Pm@}V$qH3&n%byLP)@d(Xd
zY?OJj-wMkJtdC3TT_;#+plwo1A01xl(8rlvwXMMb@z2-yWY9Wxaky;|$GbMuSP)Sm5X
zM>5t}P&*nf6S%6gCT7<64bW1E_)rs6L+{wZr+n@-n1(u5BrAMh+S_Q_WCQ*Q=p?-#SRIN@=+4Y`W9{XK
zyk2*Z_vx+*qH&64=`H8EOD(VM4}0rOj@(=?+Q`u#4BxJ&uAWP$iV@zR4K%+|RuK0S
zS=?_RM)s4Gi|wIkFQgA=2v-myQ9?E6g^1QGKr=s{gmOmb9-&U>bg~Lx<#nXSbz#4e
zKNIGYGWuR4n%}-@x9vhI#3Ey+BCq2=l|yS8JtB0Bkj*M2PbklaxQDgUr4
z^-Z1oxY?HNd+l1^9MBG8dX4g$kt|758{S$wmr3bOEf$tlE_yo0f6{h!90
zG*@yBPYhcO?Oe~4ID3iXzZx_}Zu8c;oG$X6AS1ngH>tFg%HQcwHKkcI(76-xjfh8<
zSn#`n_t%*pJkgydX+^dn<%?;MP|`1MVOh}H(Sw1_4V!f&B^XeIQ;CU~?zW$pJuUj6
z#JcSQr=fXk>5K1u;K>h|0x72BdNbo44n`)5-6E}CTS-2;FuMvjTnI-WQ1t~RnGIv9
zoNzSpeY!_kXs@LT@aD}G0G)y9+f6>1O-?ehj$fCXjX_jN>}zF^H`b9+kqsr5RNJb<
z?tvFtNh63A7bVYXN(^=^G0a(QU>FxBG;4O=fO-Wu2dE^Z`ZdG?Z{XU85wB89sr6vnwF6`+0LyRJQ(1k10m|aXpQ*M^OUX
ziYlSIx0*k!YB5B4#2DTy3U89`$~Je}u#vmwQP3AUY%`-+2Sf>Q_k^!lNlf(G-1!`c
znQ(z&I3XetAG?k*Suz#dG)6SEDjHrC55}PESsFwO4#jmVI;yShKzEUwbAyV$@=|{7
z>PlX+_L~AR7H!tsj$rcoYXYqBFK5jGpW!U7+R}>Uq4$9)^77E%sk91({7o%%UKu}s
z7eSGIUuEvRTtc#8vRJ+=2_(0|8w>2JeODsbp-Z(@_KFd}GLdMyEIh>ZB5Px`4Svj~
zVx{oL$ipw)Oqhy74*UM#%juXBP~L?l@AjD%!(Z+5@E3I&;vhVpkA11N+5>FUZ7aJK
z!#Qc|d@zpUn|G3~ZVe(3-??iFT`?Z;N;k1TYEspVq#$5-4rsP*Wf800j;d^+L>~Yt
zr0m_b$i=T~$#guR#M##OIr|}0P3#>@@{`PUY{AIbz_=B41U{w(h72UokK>rmz^#x_
z?mt#RZoc;`evF2x{SXa#bi{dkFod>c5{|e?%F=GQ9lhKk{9!^x!dMS8F}~fBqBe1)
z@r3OqkzMlb{zJJhHAY||TX{!%_Bov&{Hw05*M0-S@8SR$U=8WSEEH={t(H#4+j=>4
z!h9BD7-y`S8ZG{U&EMMnC+s(&&Gb?>0Bf3G}go-Ryl?f9Qf5A9l-7CmSSPk2tn|Q8uv5x_kmq)%POqBG&2Hq)frqfm=^~
zudsc$V$EyveeGHcmN&7558~?}7b#5PIhmCKUxDTNc*XTfs@*N6zz4*qz
z>S6!m!s)ZJ^ACC$>wmUz`Y(FebH0CFIQ>fx`_FW*zaJL(y(-H52i@sE*l#Z_BW(>%
zjQaJIb^{hh1Uk~UzRpZR5nZ~0jVEEUF4>_;jf3Fk5Z|f3C$eqTa|jii7xEQi5~@$M
zs2xF($U7Kl!J~GRH1|O>N|O+ouDx7h)jEBEuJ`zvcXhqfCmsrOT$HAnHJA3+clKMi08JVV4-Jk542~nF*Kt7{rU;CRmRa$9|S}ECx^I|
zk+qPifddiivnc@7+|qgucBUY})~B^{#%CN?Mi5}Dz+WJ)te_QlcDARu|3Z6xf;|3j
z5Lf}iE)JrK4j|0qXJBLnRwA}vJGOawLHz$lcvWCxBw~9i@EOeYsn`DQhitzr5C0Q*
za!RS%lKm<>ve&duX1yf0JpP1E!`i>5L3`pexGMKjs
zCFV-{FpUTi%t;Qg^1Oao*Y!1jmza;vn3XzFbVk)6D3LiFNYsy5OBhT!q`Kzv3L8Ap
zPeHR$+^00C;ltR#66aY6^JWIJ2vL-C>7h7X)S0Q6sRYy|eDDi0rO2qmQ8RctRoV!T
z6Xy0xe#Otxnt4twSp6Q)*c&T6{*FpsWe?Lns
znBuh)n0J|X(cx*m+O$7o#I6{`t1-rmqOq+H#3%Jcr((#l!eQe-O$u1U5UX&)Wg7JP
zol=_-7=i-_3vk5i^Jmb%q6
z%Fg8oTIR}c-}y)i8|7D3O#y-H;{zwD_`!+4UFcp)OK*ocGR}fO^22^tTpU2odgGm5
zoI9(T#dot;o9vwBK3xqR`J{t5P|cDjlklc|Jc!Knfl5uRV1+V~#TcYy
zR8*a5d}Ty$o`e{r%tEZC1#NXZvchkeJZ5Oqpod5gwE&ZtU|
z{^AJREY-$eS@?DkEGDeF`iSZW!-wHm*;u5_$Yfx6u{w3W|2cw&!nKUl`Pi!NyBr9b4Ao$tyM$QUy}jLUB{DWau^PBTU3~33+nVZS|x%dHQ}_vv&Vht-v5lr%of
zlfU57Pd6GsW($W1ee;nd@<*t?I$Z{-JN}0uruYO#10(8A$P6AIP%kOu#vH-=P(n`I
zKa_Okpqz_kX}7D4C=u%>I-Azz
z7D|P3)owGbl#Kjq?J3lf#2xWwgmTByS6qXR^%RAmwi$Gv(@!T<^WWurZ+d8YqKtA2
zBBOLrY@>)N(V;xKI@*<$xf)|cW_ijiyXi=F5F(FhaSf|!oz<_&2acg&GBBCf_0QKZ
z?-95wTdHXIoGtG=EzcLjAVm^8_cv);$9(6kIpN6~3gXok>Lw=x5Ba?AuE#m*AB()2
zOT7zEDw5N`r61@ay)jW(%NkUu@b>(1-?Xib%Rdcqb0O$>{&2m%e>9{+vs-h0ueT&I
zU}1MV9e=x@xlEbDyFvMA;)D|6y~pipbl$&ou`jf-*s^$#b`6KC>}B`$@Bk7rq4I8R
z^}*}%dN}(>dgJ{(0teTG*lD_T;9L5Wxt7J+FcZ2QSi<6hNYF{8RMH1b8n1b$rhP)$
zp?3}heL|ajyFmzgVDUkC|9Iyx9c1CH3Cy(K
zALzdAm-S=#d1>WrVwE)AS;yV?%7+!{v+GG$dAarWqM;9mbQ8w8zRT3wy4GQub9n^w
z-p;tBef>DlgTA`nO_oYAZsu6So+E(=Yo!V{tqtXN4=LWZgP3am=f1&stvgU~g?*X@
zNXW+SIrEi(B6e#$d1C;#+GQaj6ufvZTq(enmoV9B>AvvIet^fj0wkd>2unhYvR6EQL{GXI_2@6D3W
zT>Lfhe4-0VE}&-*e4f1lL_usR^OG_Qhyg*nq8&^P^ymewEet`${GO;$b_f4zHp&jN~L;{+W*{+<5m{&fAx^FQ66X@KZF(fD(*
zgHVRJIM|-cdMfYF#m>$F@@v3E!~vrJhmR-vPtQ;BoFF>SX+iIw(*8;JSH7oopm;7;
z*5~`5_$MCz@c+yMh~}RMY4dtV$W^xl>0Y@e-g=GME+3z@9xj}{=EE$z*FQuQT=}%Wa0StZTZCgQwROU
z4?Ej`&&N}@`X{db64i73Q|O7@bH8R{dUD0_*Fg3AfImDls44!6{JjVKdARwnvE#pn
zieE#=uaV@x9dw>XmtU$X5j_(J3lTjFBL|32F3<-P7Z*rp1$o~w2DvK%jWPW8Ac!-N
z3ldN~$dQSjjTq3>*aS4HvVumOr&6C@K4pgoh&$+6m>Td~8Cw8}KqHu<1JF{Lh=~pK
zC~azQZ)#=yJU~5*10JCIn+_Z66IsxJ_T-g>3gBR32LuB60RjL)fDk|!AOa8thyx@6
zQUGay3_unj2T%Yg0+ayn0m=XsfF3{}U;r=#00Bk-V}J?36kq|c1XuyA0X6_zfE~ad
z-~ez0I00M$t{~n1snSoHKd9UPLJa?#sDB%y{>>mR2+H_xToV1IdH>qo4CwZ#Uf~EMdLZ^^(~($A|M9>TqW%z+D6fKUY+~+L|Ki5Z5#l
zhwd9)rWvZo9RZLk2DpYXG)-m}Gs#q02;cO99Hi@qqyOMd+}qg5gPm^i%03E`knbg<
zvgjwGr;&-lCX9=GqsWv#7#^x9N~OP26I@9hQqk8}AZDtjtExGn?yrEKFC#UPL%p{#
zPs^LYYFdSUow$-8j-nHvb%Y`*Cht)YS@F0ICoCUM_sxY(kruH=6b_GYCCY&Z0EPLx_eT)h@~n+9zRZM`3ne~bu?b^lV<4}BbR&6Yy-
z#0BmY*31;4^0j%|d8?>$e>V;R?zV`YqV21&)Yqzo=9#B6!?Qs7sPBhcIQF6qJfDrGFgjf|-aR=zqoC!$7+nwNnCP92K#s(Q#C8kFmfUh%6RsEeF462|_gFd5Il8^Jag<$b{yA=tV&stV
z=LdSw9O{iZ8aQ}1dMvmTup;+m0d{cQG^X)zk2`9NSQ{c<_Osklz+=JHEK$+1t@(Iy8xtL!ZdA65QWXT8{!by$ag6tIv-$*cVhv
zOCwH|^)lbvK|8x~6Q)ifiw&8#I7581#S1)QI#wjX4(4Oh(PzYf4u#S!RWWtgj<
zjVo#V;`W#vvs?KIk08AC*d-;bx(z;p((D-els7?*N>_to+EqS-C9$KBz!yR4{z8S3
z%Uu(9ag`e-x~um_zG_kVWO
zn2ei|0qrm04QwB6Z&8
zO)>k%aV#d!K`?s7%Zr{5|EQvNdqx+q{H6Rsu>0+W6W3k0*+K}-V0Kt(z
zkLV5{MyOC*Cy!d@k;YPzbrfhyqZ+CKlkn1#94GM7{Mxqf{r+xk8;ctrNw2Qj(2V&v
z+M+bJ;Zgeu`z?vyxC?9T3}e`Yt$c(=k2l6j5&W9GbsKsksmE6k7JQwU$$F=WUm;`1
zoxkcbAO@*Qfl1z9Hi{wd)cL4GBg#+z$onjb|6Ue#1oqQP8r=b+mAtY(DH^`(VaD!y
zf7MI65iFzSLYI(7Iw=J2&IS`v&h-8=Q#CI+g_tnXx;86&nhiW8%y0&C6&+#xtur4%
zdOTi%*!bQR`~;Gm;vU4H9k%TwQYNH$n1~z0`rN}Aii7oh4wl={hlctBGHd&&<_PH<
z5_Uca+wYmh6mSR|oJ&)#J_2Wr1L|k8WM1_?hJ_pNBj#r54j@9y$Y4nH1ngktG*};uyuFb94{tBp`hR?;q8J
z++##NV>s_w+?h$cxk)Rp-}Ka`31r#}xC=8uKjDASJ3d%ISh4(o?sb4GCeGuk#yBXK
z@=jUSxU;Nm;ZU-#+NtZa3q_hE3T}G^ZEY^SF_g0tzfFO$u7$8?9bSaNNv8x_cn@KO
zrIvR*)%^J~QUzMqDcwUz#hl}mXY)NX!1*w4KO#Ntu1;`7gQTuz+v!Ied?F-)JuhNQ&_p-w`$Rc0u;&n^2hhS?8kW;WVVSnxVRzeI
zm9>1xaI!RTc)##tq+Boxjec5)VI|`eeu2+A*%5hpN)sf7uDK1iK6!^5XHTc*#@qv0
z>#&@R_UFiWpBgIM!WK;xoPqWUPbo9ogXsb{^peGz#fETFBOR@!lW`jd>aScETS=X?
zl70DOVc3ylFd}QP))$;}wWcZ?WrH1tIUun
zBW$S`-S&+PNw*0No%e!M&Zck0K-5Dbv0Cg6xc1vfN95SBuXh&MqRn0@V-4t!ZBpDX
zFC^cuf3}PRO*49vIo8&R7~VLGX1~Z6+xDmFc+9IRcFT;MeCb9R_i^Z~7dJ|pu0M!O
zf+pdm3Tq5T=dBHL4^GH%H^6lwZ5CA#%fqS%Y(>
z>1jk7t5yFoH|~1ZJi>9#Qoavi)Da0|vnX1u)A2RGMeS6-=M0#FqQ
z8hYS*w>H6q&>08`EZX5~^il2?o%bw)a;;{$(98m|L5OLyAJSC?4Y;8P9KhKzp-d4&
z9ocGH+SEe7LUdP1ydh1u$Q%G_33R5U2t+K#gf`?bZdCP^>f4SNZA2dAxNi~c{Fq;u
znUJ%EJGaqixJ8Bzc7gK~h~_z+_JYOipc1U~piFPR9WNjCLQYa+OmGLA*PrKVnPFBz
zCQa}~XmWNWIauJe{KS5lJ%}VAf-LaO;xQ!?-f~1??SnF==z9v~Ql7ZcAobDfuhPbm
zMv^keqIgoeM;1Yw#zG}$s!?)k&LvrU3Q^6K6?s340TOjc_0VFx9yJTN1<>sQhU-6%epyCzA(LGz5Q(RL4x3I
zMoLe00`<%!lxFTrm!Ecf{b@UO90Bc{G@dDyhIj%9$aBh>cLiNp
zp*e-nV@0pcTy@~0dY$?+W~%;6@lxsb&~^uE1%5#;Z;>g^4Fy_jIn()zN*ZGPu6
z^$RrGh1dS)F<-6j_mjEjK}~X{+G{TfShTJ_R*Egs
zPsO(HRYfi$CM2|(C^AyBviOb84-Hc9ct@R*G5m|oitYCXUU(KpPPRWf-SZMC$RhQ;
z^7m{mdUogd56jcP*EXLWJ&Id`H1KB_y8rAvFWbLaxn=`FI{*6n_xjD>6q>oP*n#k5
z##aHCuJE6Y^eDo8-sSXXDO7J|?U3x`l}5G)cG1lAwXu5M-b8ORvcPqMUtgroTYK>0
ztDwL{L%GK0vvN*>WwOrYsYO~|h6rwp(US`3Sp1}o>gAHj8PYlCp%KlDlL?b3x
z)_Ao(Cy$4(jEMNvL6)O`Cw)*to-AA`H~D3S`wm8XUm-+0R4z<*IPrjWHK
zXfR8Rga#Bt#X^(n$`=hFY82LWyLW~yN`vre@5<++gsBa!ZObn_TM;BqGLDD%FA23%
zqqe!?Utf?f+Py;Z7G6#CC-_~|V%5yzC(!Rai
zZ{yid_HB9Jf@wF$&fCF%A~5N{ugY@G%jXixczro{)6*X7nmS;9$sPLRyS~G@Sl%9M
zN2lJHK@6+J7;76-81qgw_Y8vr4%^2Bg%I{8)LvWPmP9RdA0q6AL>kRmfB0hxL8s1T
zrlS4Ci@c@IPb4Eqh)t>Nx7jHd3V}7>S8eLN>ViaNeBVY?aU$5@;Py8o8(PXLr$-cf
zt20%Z{%fcH*QNM>@_YK+1OH(w2gqjhpWu=G%T|u(eE<4gVS2KnJTHv@Z($VK{|=-0
z+Z=2K@;VH%SAYy0rbb3UkPFr)SFIW#@&HjiOG`bF-NV8`51?cM1bs^BS?U|=0aQ%^
za;5-9Q)AHU_cr#X7S>h(ITKK%y(vJ}5@@Uk(6h6%c6PM+bvLwjwgT8XS~~!r3?mjm
z0MG@r5yA3zG|=9`)EZ>12Bk5ycCY|~tS5gS%mBuAdQL!qo`ItS5MZEZ2P#M34hV{|
z)H8Sr*n#K)3_$*5EiCi^#y~quJu5?f3wwYb5MT}Zw+C%^Fahuj17tz}zXDFC_NMxv
zLTvO5{_ydOHBdf))$=|GdsE{lRzM`}05(97=UOWVfSx7j$sR=gS7N&-vk<`M$ziPp
z(CGKmuSk%U8fa%~Zw|1raI^;)SU-u-#MQ*s7eLtW^NF=W>C+OSStI`deuZrWS@k
zfT1aX(a|Eeed5e^X(
z7ssDU{N2+r`_ulKKO@PneKeq-Rc2O@CG%H^nTU;x8DyRO6#`iepO;*Ido1T*X8KQm
zKKJ8*nz0(U1^lttk;vPbQ8)pyXxq_|R@k~9h>diy-w5$XMn``}-WE}KLl$dAz5N+A
zXKnDut>@*%-ea#lZ|YK_G4I>yuLq9hyK0wiMl%UD4lY0oRAo6~NW5-@4Tw;P^$ADE
z_PckCgnX@ZAHt-O;rugnVZXeEgu;OF?e2p7z|0_)@c|Zf5Wqn8;Q@T+gB>iEHQC^I
z0;siE_bbt8D!HE5As^KQf4*~s{YJ3!sZ0;a7JP@OL{Bc3gHW!gAFcr86^kQ-eF7=u
z?8Dp|Nm
zL#pwa1Am7DduAqAGw_I1>W!u1;A}Vcf%8%RVL0`Ma23@!A?DJX8TC8E0TmTmL3o7^
z$?~HN+lL48eB@f62JT09B`7FcU$ud@{gjM%53K@>9Ks1$6Cwx!Rdp~~kq0AuopN7>
zTJcA?FZ02~-b;!qYKPtON#EPcA%Od3JABo8D7z88NJo4K9G&<3=n5u%HF}F6LYY#=nG;ovM$^Yyu2o8&^Nw5j9-u@5Ma**5d5xw!oG5YNBv6q
z@J<5~?MPJA|H_953l^7M0sDp){z3Mv{NWxS6WVw0tJYrTQLA7_#0MhtkLWAu
zL9lKN6O;Ih)2*c7?;;loVO}qWJoX5Up0YkrFgfTs*U*d&;rM)fK*xmgC3-}43z_62
z3(&r@20zVC@7&^v)p_VQm<7TnFedGtL377`uI(k5V+~kjo
z6Lz^ifd_Mig7o*v^40xG>gDzL$Rmw{Z0{IY)Cb{zeSJ*`2Y>gv)=4yL7-nZU9r}#+
zNT<#}a|w%T0>KO^Ayq(C5|W$D6P6lYt`DEOrPrU@T!k=Oew_E=C7%fe=6Ip*yP-Jq
z%DJN|vbaw+t3RN#H`z)xTtb(sheHAZng-F7-j*(^RyjAUV>kr)MegNIjOj4@Xl&jI
zoQvo^6L=bUbZ+Lc(=iWUojEi@6}6qwH)+vD0)s^@Mb#^U3NrPJVzjr4&7AltELRMt
zv!(?y5z@M>R}UBvVW2oauCUt-6B4U0hy(0xt5>|$bTw5CM4Mkl4kZ*4Iz_@uujhKz
z=+am_SE8DVKy(#Qn?O{m<;FSZ#EMF>wi_o8Z_bG*J0umoX=b7*Yh6~e_00RR=l(62
zg=E91B0@3CMd$XQ?2CYB;hV~ATxGS$6Iw5yOZa>x>)dhH%ghD-9#?TZ|5`HM`ZF^f
zU+kSAzUC|mJX8O#CG}9q5)95Ku?$pQgaghddoOWr3!SRF?sGaVqS&s4)rno^>d-^o
zC7Dx%%6MxxqsJF3ojMY~4%AT$I(5IBC1~R!OY!QV+T8+Fuf9WV)*Z@>iHvyP9M*Ia
zd?~*{g*A>z*Ui1kWj?Aw-eS>VC;OJWdFqy|+o0&K+9t%qjqIm{uZu+kYY{1KL7hKq
zZ-O?xK#ohTTxcWgki;@tUw4_iv#GXbwF^qLA#*|v=BNiLT67Iq9kWe@D)*1C1RZ=|
zmbkVpZWkMp&ZEF7Pk(;D?49TK*B`|@x=<_I4AaKfzTDt1dQiCzDhxr5*Lo>E*uCXb
z3}4Xdse+lnI`wW%QruR^%bgn67r)J!Q1mI5bLmt;`s`Zz*?8I)MeqKS@~OyB$wem}ko?
zEQl#^`l~`}x(3_PLf#IPa@zdppJh{>lx)B=?c3BY^sp_>vlh%+Ud36}q!sodf8(^X
zS|On_1X~e-l}PbIk+1D?r%Pn7cwRnU#55gc1YzyyCw`mp;>V*t;DEf6WT9}jop>&I
zf_A&)=bomjG@EqImIH}|njpAw>d(6Z7;)Mk_9axw`&T17PF7g(eM)z-=?K~=B;@av
zD{VGYb29}SFsv9?oDR|;^p2QX0ZwWlEJ%lgz%Fc-I7Tb1PAaAJ)1I~?#qIHrhpX>(
zy)oM3)^n=1#c2R3JviMcZbHy=i-y%Red$2Oo+cV5%Tv!#=MuC|2!ue-WO}EDH^I#u
zn-^8&0AV-sr_@y`qPT2>L2(YG$xJ?Dk9vO)J#JW)5K!uR9wQDZ
zI580Mg2}jE^hzA*$DUh6@;6yVrB)$7ReEUf+Qj&u9X$Im$HO6Q`McaAcuRr`NasTu
z;It^oRdl;xUvzL4?hUfCG*eBD?wrEo9fNi|oK$M|rkT>s#4^HZ-$94<3eWn&-BNzH?E(kdZP9&MxA
zZ|^w!wuS=Zj*!m~2ipl6(KG*(2b%32gJ|V`1^Gn++6=zoMgNaL%@ghK_OADeoIDI=
zR-djW*o`xaC~ORVqc`E?2N#S|>+bgbOn2&F*KV
z>#(~>Uc!QjAFOf&5Vl-xS)V^dHg{LJs+nMnjI$dCkg@XWDrFo5TNxO#DFUR_R|mXZ
zLc}@`Y*>$zhPy{LbQ37>?g+|N95Ct&7GE<71~}u2KS}b(TpDUtQKU_ob~Ad;6jX6o
zjqASr=$l@#dxYcOuujMPu3$(lX=nV1fPpMt=tK#Z13}7Gf+m*V*(?cZll(nGpUCW9
ze`n+T>5fSZc)65*5@LEuE~Im@KY`EOEGG2~3EMHg3?fN7J@g_|%R$22>=Dh(=B0Fy
z*0)fo@!ubTFP-ZciMa&clo
z_*3z4(PZu-bdc^ZXvLzi!B;=;r(c~E<5x>*4)FKgjRxK3%Os*WGhe=3tOz3oLui>>eRftVbD?-
zuRVZpfXY;-sZCmQExIA|j%UBl5i&Ik;yIuW)Lm5=%9=xo0dHiy9SnP|jMZ_C=9bpq
zDd2DDO^)Bit~uA>_JyAqD_xD5)Bx=gK+{XQ6WqrGsnt)^~AmLkx^77$O
z`EcwSdP4nN=#OZ@%Mq
z#>}FK7th<_C4GLC*%jI!Z$DHF4h&QDR;sj5{z{DBh2n|M=J5(T0Ve)Y|EQ0+){=3N
zgqu(?*x8L2fQz2gtwe2acPUL-bQix2{_U$*glzp?U3QKa`ow<&}ywJq7+0nf5anih@M
z!FR`s;UxY0vZ3W}@knb|fv*+fW$?_A}OqCRi@1Uh+MT)0^CbH4q!Jz8{9hc%tHbl=deUa
zxrv$4J_o8EFRqEz4p7;P&uX@7P=B=0G(al1jyKx9ovl<%)gawe-OSf%_VLM{`P}|6
z{s1O@zr**VmBI5hELsOELIJ->gL~9A5*l?=(*@_6NbAxCSzoNQG0@PS6H}jNCxq-N+`=_&lF5;I
zxYv9dT_)M8XHq1oXV5TO*e>2oI6oF#;4}XtNaRg1JT4Z6ziS2K)p7iC&xSq0^|H$e(6Uie?4!jmwfi-b+Q5SU37
zZ=!l$CT(c$w`S=QIcT@M=8^%p&AV(^9^sIfyyll!EYxLvIlpDNKs~QMM}8SIv~anw
zetws_0T%yeya0>tmB=Ynjb7MlNy<5ft}7i0P3UfSxV7UvR&Z9W{!(Xd%LUzRD}0--
zHXj}!SXU*F9zlN-TsehiQ6lV|^GxQWz}fl-CG6iAJF}
zi3BQb^3E)NGkg3ks%8h52#{p$i+1(G=9lDCE=rS(>$J8(1KUWc^4X?!6U);Av*o5t
z4HOoJAfcQOJ`^z)p<`;;rkV#hrth8AXa|t#4k6syp*$F1*W`($X5<3yf82TLVSUa!rp
z!X@q~xe8AKov^of=WS54@s>?H2F0S{8_)+fRqMUDE)vp|IB`vek|9n}s&rbL2VEwD
zP$&_+h@#V3EMI8Q)I??u<{k9~*~eG5d0gzizMg_zuO)|z3+0?M@6EFcvGZe2Yg9N`~K(7IY0HXO)S}Ud;65s_Yu3DI#(9gDXq);
z)4lqMtD(@Du_U6=$&WUslOpT>FN2Hg=ONlt@PzsYtL}N;%CKV;`v?sh8x$oih1*+;
zMSLWo$uKs?CXwPP8-10YbhV~Oy79_VVtMi8TUb-M_autx2?KTUMdy+bv>I+~G32!o
z*TRLmcj@>mVF9nk)5@RG?;P=0oX7CU8o`e=*fNm=3&s4Xu*yxN}3o)3Bx$?3$L$|MsSJ0#@$mwo?}XpHdZxno}Q>~f$iY8
zhqiu~&|(kGe}Cff!jtr^EBvzcML=5=d;gbkCfdQR+vrY7gC0))R|_6EkX2yr-3OlL
zg;;ss4I7_)!zDw(paxL$OkcIPJAkxMLVafe~dlHfziyyV~mLZtT|7
zP>2?cwIJxWFfuarHKgrE1aGWaD9^GbUOLSL16QJyU^GWAfQzzwN9E;2N&9;7H@-~p
zg54ZK%S%Ke)HgTH<4NWQcJujgVr3&u3v9)a~7|C#0L)o~k8`>wM4pd7^yabLR?^rzP?0g4w_@+{F)4n94+ta9`Lc{D+_^+`B7s714wOm%}N-f
zT*ht#oRcQoiwVQmT|wSc9k=p5%#j{cTLigvwGJjqiH>7qFHV^WK)MxHr3BxwOp@`$
zFE$xj^r9lwxR06=D;^nUKcO|Nl<>&>nne)AFdwNW2a?zk(E=Y=QzmCj&ZYaTZ$HTL
zJ@ax`l7VlHN?p0l{GMoPXfF5h9XoT-j81QS8p0|jE)aYNS#p;|Nin=T06Ev
zA=8k2P18ifqd0`D}D=w1?cCt+=y!GFU!}H3B8KYEB+G7dAY;Ulo(Z**p
z-^9u-eDTMNZy8K5nl$xE8SC`eH1y%|g4aw$ri~UuKSiFN-<3$AM|X-#I;pz)@!*%Xm-OjX$tEZeOEI&hUv1Rs--=)6c*`#+(5tZ
za~_$&`Fj1R@D<~^G;A^My)2_lc0^Jf#ivC%Epjuy?Ya(2AzLD;KQMX2y(#ebjbsF~
zT5}L)==*nSeh4K~P{0kW!1BUSb%B@_CHX_T@5CLF;d(&8QsM2w<@$7StJJH}er(8L
z6yNL$>ziweisq#8ph3udtnt{j}kZEgq5Pl@*McZppLhL#{qVy|2RVYF0f<3`G`uJm*bsK?Cz{O
zZk;z0d&OT(wmxb0^^@6^Va@{WO--#zUW?|MLZ3MF81F>ZbO#uoeQ{3N-ClNq7)XPe)@
z5$0)zvuzhjhcD>*3|*(Vvw`9oQD&IGjnY4AfGs=!fy;^u>Y3wLmrs^>9Th?3wJAZ;
zn?lOL7t8F(2Oz&`w)t@qP60V-V=K}tH9*|%?@1#QN{SID`T5&3$(42C^P
z!RQIKL;?{4)y$!yp_F|u%#y*moQguE2~IHwjW_e`b3KcvNz&dYZtuOb6m~m0zeC{j
zxzm}y3kDmPk%K}GunEM+1-*m8nbNQNhjjZxqWf3k>Q4&5Z?yOON&ucO@-vT(|B&Mb
z((iw=hnN9W{vYfi2F5?xL##iEMt?GwB;*ujg(azkq?MHZ88PWm8jqJ3{};XLcg?$>
zP>;7Xh_xDuppQM9_MDB021N?`C
zqrz|UfDhLi=tyWy7Dg8Q`^SAs65!1IQF^q0fBv=YukZN1#qZ@`t$(lkd-<`=ul@bc
zULViwpSAecTYk0H{sneo{YTx8U*q(DFW$n)2)F^$8ClR7-M;{@{|!(Idq*1^AR`Z8
z5FX$hf50dnGxxt?Dee&w7WYVv|C7-azhP1yzYo8#6n}xK{O0teLNC~Fj@m;@=672W)-CsZnMh3Qj0h9pvZU5n=9d{I5xueaW$CC76>Qlf+
zs*+0@HBZXw8-Jsp+U&|TUBwMl$JEb@s8En6L={4fESGEJh~AcbM?jP==HD|;AwQRC
z-*Dua!q>Q0w0BlNc-3kD_qGAQ`Qc=A3{7FwLT@Y}vA<~e@
zqv3=ldt>1~urs5NgtdY8gO(>GfX;)wK(zpG{SIvXAzQCyc{BN7B6Q!s*-T7XS1ZKaW5
zt8)=P-Ju31_@LlAQPI0M3X|Dw84nhnh&=j9DJXXcSZM;E!kN!p`+@*8P4t02dMClK
zseySm9xnLt@Jp!Z;WKucHeYm51t@0tDZI1~>jd)}I0gz{N`(8?(vaWA+Peqfi_SVt
zCO&{8zwbVH3kwE1s~WZDKS?Tt?2EjrZDz;i;Gf
zD+My4i!hi!DEuc->?p98Ag}%T8IYR|Y{5lAe3|4yV%h^elBm8T@+)l2>6y#XQxI>Y
zLeZE6gWo|wcX@dWlvDV?YWjIMHok|WS}f-MK8ZNQ)&vEf{yr3xNxv-D+t`iKv^gDr
z0KO?rj$HA{NCi1ZbpO+YsV*dsR6jJvR-
zj25luoT9~*+BnUpS6jtux0N}fb$ngGIqGo}eY}Wz2-%RvTR(U9N_Rshhr3Pv`g;20
z1r}=;_F&CtH?P9nC%jv7`r+b!ZFa~?=nEI?#o}7&_D>eq7(`@m-*eE)P_|coSG6@_
zT-~gP3qF8Oj3;2L0&
zMDl7JQLM&A1@y>a;ucwjjJ!o~XdC^mGHNszM_)#5uMVR)UtW4DYaO3hWW5}XW#P0&
zVe+u(GiO4|sBsiWf`Ce~FhJOAg?LZ*ZCvDQJLfjPzzEp&lq8`spM_p+7sFzGehM8!
z$dhTeOs&vVF%rxM6BqXaP(t`{Hx?|6VQ;6rU6ikh-A-=={jj;8lzh~<9n+7&ZRX*?
zdWyQDYdBiPyfLOq-q4`s=S}q7if`2GCFb)E*A8omgm;xE+^+(*8peKfrqI5^{@(C8
zart_<-?fz$Cb+BqrEkdlv9Sd_pqt>$F?$NSd_o@1nsbjZNvjyDI}CODJv
z=6=OtR~f!Q_OL7BQT9)g-a^G~x2z%R-`acD-SA1S*-L6Dj7i(dk!?9jNhO(ss7ymJ@WEoPHi
zwUQh=68olSe}tyy469u}iJWJ?FZxc<*kj^IJ?26tF
z+49b-=(N_x`dtBeO#g&PUG&UVY`cnD;KV^%am93hSt5e-Ca6*#YeEp0OVO<2D$&QzF{&>aVxxIuG?Z261pnCC%
zp;&CCIe6Kqmr@RfCU}q4H~)*H4z0B{ySJwPcKLdj!L02bv)+y3#sMprXKeuaUAsx(
z7fXI(bgVPdaXR*ph8
z`ueXt;kmS6hsl_jLyxDorDF
z3++N84=_8kP%5J8DbUfnZLN6ms9JEOd?IdD#%eIcPaIp8=;DY@IFlmAKw9L&%R|EV
zubV18V@FLIy1NQ^ZMFoFKArb`OVwwZ+O18vXJb^6G}HOq
z7cHGHXiTNghU1P@3#<|AuH}mSSdu-y!OWZ))Fmb>Blp9b4Od~Q`KDeeF`vws?dE-S
z41whwXK~rAP%||c5$FD5u9He1XGzik;g-w1>B~A{>S&WW%Au;L1-@6@!x+StCE{EQHmIGUg4&HMMJ7USBXHa`%ok
z8L#deN_OKjwC+hmlU6fE*zue4S*aK)xR_u)4WuKm~t6&z!L1f)K=mda1>wQ
zHHt=&D-*gK`ZmF;ZlS=2o-)pdDqw~#3hRfPqaedKRs?ez%|)i&@mV`tLX*{#+US;B
z+hkd~_CmQ3#7$)GsGzr$ujP^Zw`rM^r?6uaJSdV{!sk+Sg#3(FaUiCVvpLImsWYEn
zmKO=#1mP{;LEMr-M{M8+_vU9D22V**?-vccV{>onhEx%Y-nk?-1T&VvW{#gJ~kjL|6eh3^rKL6`8zu*36kmbHd
ze|6x$CACO_M=Hj>e+`fd4CpK!EgVd3EL`b+u{SJr9ZUemfs^S2@j%xSAgDZ4
z-IF>V_wexi$ngNk5Eh2N6(1YiJM$KXf9(0u{JG1YhP4CTV@coI#?{V&&f3P%O3x7p
z3;^tl#}^