From 1a59c41c272e22861e759e660a64578b47147d85 Mon Sep 17 00:00:00 2001 From: Gilles Perrot Date: Fri, 3 Jun 2011 16:24:59 +0200 Subject: [PATCH] initial commit for lniv --- Makefile | 49 +++++ defines.h | 29 +++ levelines_common.h | 24 +++ levelines_kernels.cu | 315 +++++++++++++++++++++++++++++ lib_alloc.c | 136 +++++++++++++ lib_alloc.h | 6 + lib_images.c | 306 ++++++++++++++++++++++++++++ lib_images.h | 9 + lib_lniv.h | 4 + lib_lniv_common.c | 99 +++++++++ lib_lniv_common.h | 4 + lib_math.c | 227 +++++++++++++++++++++ lib_math.h | 27 +++ lniv.c | 186 +++++++++++++++++ main.cu | 198 ++++++++++++++++++ main_gmem.cu | 143 +++++++++++++ obj/release/levelines_kernels.cu.o | Bin 0 -> 20176 bytes obj/release/main.cpp.o | Bin 0 -> 3040 bytes obj/release/main.cu.o | Bin 0 -> 80544 bytes 19 files changed, 1762 insertions(+) create mode 100644 Makefile create mode 100644 defines.h create mode 100644 levelines_common.h create mode 100644 levelines_kernels.cu create mode 100644 lib_alloc.c create mode 100644 lib_alloc.h create mode 100644 lib_images.c create mode 100644 lib_images.h create mode 100644 lib_lniv.h create mode 100644 lib_lniv_common.c create mode 100644 lib_lniv_common.h create mode 100644 lib_math.c create mode 100644 lib_math.h create mode 100644 lniv.c create mode 100644 main.cu create mode 100644 main_gmem.cu create mode 100644 obj/release/levelines_kernels.cu.o create mode 100644 obj/release/main.cpp.o create mode 100644 obj/release/main.cu.o diff --git a/Makefile b/Makefile new file mode 100644 index 0000000..2801dda --- /dev/null +++ b/Makefile @@ -0,0 +1,49 @@ +################################################################################ +# +# Copyright 1993-2006 NVIDIA Corporation. All rights reserved. +# +# NOTICE TO USER: +# +# This source code is subject to NVIDIA ownership rights under U.S. and +# international Copyright laws. +# +# NVIDIA MAKES NO REPRESENTATION ABOUT THE SUITABILITY OF THIS SOURCE +# CODE FOR ANY PURPOSE. IT IS PROVIDED "AS IS" WITHOUT EXPRESS OR +# IMPLIED WARRANTY OF ANY KIND. NVIDIA DISCLAIMS ALL WARRANTIES WITH +# REGARD TO THIS SOURCE CODE, INCLUDING ALL IMPLIED WARRANTIES OF +# MERCHANTABILITY, NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE. +# IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL, +# OR CONSEQUENTIAL DAMAGES, OR ANY DAMAGES WHATSOEVER RESULTING FROM LOSS +# OF USE, DATA OR PROFITS, WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE +# OR OTHER TORTIOUS ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE +# OR PERFORMANCE OF THIS SOURCE CODE. +# +# U.S. Government End Users. This source code is a "commercial item" as +# that term is defined at 48 C.F.R. 2.101 (OCT 1995), consisting of +# "commercial computer software" and "commercial computer software +# documentation" as such terms are used in 48 C.F.R. 12.212 (SEPT 1995) +# and is provided to the U.S. Government only as a commercial end item. +# Consistent with 48 C.F.R.12.212 and 48 C.F.R. 227.7202-1 through +# 227.7202-4 (JUNE 1995), all U.S. Government End Users acquire the +# source code with only those rights set forth herein. +# +################################################################################ +# +# Build script for project +# +################################################################################ + +# Add source files here +EXECUTABLE := levelines +# CUDA source files (compiled with cudacc) +CUFILES := main_gmem.cu +# CUDA dependency files +CU_DEPS := levelines_common.h +# C/C++ source files (compiled with gcc / c++) +CCFILES := + + +################################################################################ +# Rules and targets + +include ../../common/common.mk diff --git a/defines.h b/defines.h new file mode 100644 index 0000000..cddf7c6 --- /dev/null +++ b/defines.h @@ -0,0 +1,29 @@ +#ifndef __DEFINES_H__ +#define __DEFINES_H__ + + + +/** + * \def SIZE_NAME_FILE longueur maxi associee aux noms de fichiers + * \def SIZE_LINE_TEXT longueur maxi associee a une ligne de texte + */ +#define SIZE_NAME_FILE 256 +#define SIZE_LINE_TEXT 256 + +#define COEF_DECROI 0.99999 +#define INV_COEF_DECROI 1.00001 + +#define BSMAX 512 +#define MAX(x,y) ( ( (x)>=(y) )?(x):(y) ) +#define ABS(x) ( ((x)>0)?(x):-(x)) +#define DEC 4 +#define DEC2 8 +#define CONFLICT_FREE_OFFSET(index) ( ((index) >>(DEC)) + ((index) >>(DEC2) ) ) +#define CFO(index) ( ( (index) >>(DEC) ) + ( (index) >>(DEC2) ) ) +#define CFI(index) ( (index) + (CFO(index)) ) + +//dimension de la matrice definissant les chemins des lignes de niveaux +#define PSIZE_I 24 +#define PSIZE_J 4 + +#endif diff --git a/levelines_common.h b/levelines_common.h new file mode 100644 index 0000000..ddb0ea2 --- /dev/null +++ b/levelines_common.h @@ -0,0 +1,24 @@ + + +#ifndef LEVELINES_COMMON_H +#define LEVELINES_COMMON_H + + + +#include + + + + +//////////////////////////////////////////////////////////////////////////////// +// Reference CPU functions +//////////////////////////////////////////////////////////////////////////////// +//extern "C" void fonc(...); + +//////////////////////////////////////////////////////////////////////////////// +// GPU functions (in file.cu) +//////////////////////////////////////////////////////////////////////////////// +//extern "C" void fonc(float *h_Kernel); + + +#endif diff --git a/levelines_kernels.cu b/levelines_kernels.cu new file mode 100644 index 0000000..491c600 --- /dev/null +++ b/levelines_kernels.cu @@ -0,0 +1,315 @@ + +// chemins des lignes de niveaux +// longueur = 4 pixels +// une ligne = un chemin + +__constant__ int pathDi[PSIZE_I][PSIZE_J-1] = + { + // Q1 + { -1, -1, -1}, // + { -1, -1, -1}, // + { -1, -1, -1}, // + { -1, -1, -1}, // + { -1, 0, 1}, // + { 0, -1, 0}, + // Q4 + { 0, 0, 0}, // + { 0, 1, 1}, // + { 1, 0, 1}, // + { 1, 1, 1}, // + { 1, 1, 1}, // + { 1, 1, 1}, + // Q3 + { 1, 1, 1}, // + { 1, 1, 1}, // + { 1, 1, 1}, // + { 1, 1, 1}, // + { 1, 0, -1}, // + { 0, 1, 0}, + // Q2 + { 0, 0, 0}, // + { 0, -1, 0}, // + { -1, 0, -1}, // + { -1, -1, -1}, // + { -1, -1, -1}, // + { -1, -1, -1} + } ; // + +__constant__ int pathDj[PSIZE_I][PSIZE_J-1] = + { + // Q1 + { 0, 0, 0}, // + { 0, 1, 0}, + { 1, 0, 1}, + { 1, 1, 1}, + { 1, 1, 1}, + { 1, 1, 1}, + // Q4 + { 1, 1, 1}, // + { 1, 1, 1}, + { 1, 1, 1}, + { 1, 1, 1}, + { 1, 0, -1}, + { 0, 1, 0}, + // Q3 + { 0, 0, 0}, // + { 0, -1, 0}, + { -1, 0, -1}, + { -1, -1, -1}, + { -1, -1, -1}, + { -1, -1, -1}, + // Q2 + { -1, -1, -1}, // + { -1, -1, -1}, + { -1, -1, -1}, + { -1, -1, -1}, + { -1, 0, 1}, + { 0, -1, 0} + } ; + + +// declare texture reference for 2D int texture +texture tex_img_in ; +texture tex_img_estim ; +texture tex_img_lniv ; + + +__global__ void kernel_init_estim_from_img_in(unsigned int * d_estim, unsigned int L, unsigned int H, unsigned int r){ + // coordonnes du point dans l'image + unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; + unsigned int j = blockIdx.y*blockDim.y + threadIdx.y; + unsigned int pos = i*L +j ; + unsigned int ic , jc, ng ; + + if( (i>r)&&(ir)&&(jr)&&(ir)&&(jlpath)&&(ilpath)&&(jlpath)&&(ilpath)&&(jlpath)&&(ilpath)&&(j 0) { + if ( mse_cur < mse_min ) { + mse_min = mse_cur ; + } + } else { + mse_min = mse_cur ; + } + } + img_out[ i*L + j ] = mse_min / lpath ; + } +} + + +__global__ void kernel_trace_levelines(unsigned int * img_in, unsigned int * dir, unsigned int * img_out, + unsigned int L, unsigned int H, unsigned int pas, unsigned int ng){ + // coordonnes du point dans l'image + unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; + unsigned int j = blockIdx.y*blockDim.y + threadIdx.y;; + + // nb de points par chemin + int lpath = PSIZE_J ; + unsigned int ic, jc, idpix ; + unsigned int idpath ; + + img_out[ i*L + j ] = img_in[ i*L + j ] ; + + if ( !(i%pas+j%pas)&&(i>lpath)&&(ilpath)&&(j +#include + +#include "lib_alloc.h" + +/** + * \brief allocation d'un tableau 2D (tab[i][j]) avec data en ligne (tab[0][n]) + * + * \param[in] i_dim dimension verticale du tableau + * \param[in] j_dim dimension horizontale du tableau + * + * \return pointeur sur le tableau + * + */ +void ***new_matrix_ptr(int i_dim, int j_dim) +{ + /* allocation en ligne */ + void ***matrice ; + void **vecteur ATT_ALIGN_SSE ; + int i ; + + vecteur = (void**)malloc(sizeof(void*)*i_dim*j_dim) ; + matrice = (void***)malloc(sizeof(void**)*i_dim) ; + for (i=0;i +#include +#include +#include + +#include "defines.h" +#include "lib_images.h" +#include "lib_math.h" + + +/** + * \fn int type_image_ppm(int *prof, int *i_dim, int *j_dim, int *level, char *file_name) + * \brief Fonction qui renvoie le type de l'image ppm et des caracteristiques + * + * \param[out] prof profondeur de l'image 1 pour pgm 3 pour ppm, 0 sinon + * \param[out] i_dim renvoie la dimension verticale de l'image (si NULL, renvoie que prof) + * \param[out] j_dim renvoie la dimension horizontale de l'image + * \param[out] level renvoie la dynamique de l'image + * \param[in] file_name fichier image + * + * \return 1 si ok O sinon + * + */ +int type_image_ppm(int *prof, int *i_dim, int *j_dim, int *level, char *file_name) +{ + char buffer[SIZE_LINE_TEXT] ; + FILE *file ; + + *prof = 0 ; + + file = fopen(file_name, "rb"); + if (file == NULL) + return 0 ; + + // lecture de la premiere ligne + fgets(buffer, SIZE_LINE_TEXT, file); + + /* pgm */ + if ((buffer[0] == 'P') & (buffer[1] == '5')) + *prof = 1 ; // GGG + /* ppm */ + if ((buffer[0] == 'P') & (buffer[1] == '6')) + *prof = 3 ; // RVBRVBRVB + + /* type non gere */ + if (*prof == 0) return 0 ; + + /* pour une utilisation du type */ + /* ret = type_image_ppm(&prof, NULL, NULL, NULL, file_name) */ + if (i_dim == NULL) + return 1 ; + + /* on saute les lignes de commentaires */ + fgets(buffer, SIZE_LINE_TEXT, file); + while ((buffer[0] == '#')|(buffer[0] == '\n')) + fgets(buffer, SIZE_LINE_TEXT, file); + + /* on lit les dimensions de l'image */ + sscanf(buffer, "%d %d", j_dim, i_dim) ; + fgets(buffer, SIZE_LINE_TEXT, file); + sscanf(buffer, "%d", level) ; + + + fclose(file); + return 1 ; +} + + + + +/** + * \fn void load_pgm2int(int **image, int i_dim, int j_dim, + * int nb_level, char *fichier_image) + * \brief lecture pgm 8 ou 16 bits + * + * \param[out] image + * \param[in] i_dim dimension verticale de l'image + * \param[in] j_dim dimension horizontale de l'image + * \param[in] nb_level dynamique de l'image + * \param[in] fichier_image fichier image + * + * + */ +void load_pgm2int(int **image, int i_dim, int j_dim, + int nb_level, char *fichier_image) +{ + int i, j ; + char buffer[SIZE_LINE_TEXT] ; + unsigned char *ligne; + unsigned short *ligne2; + FILE *file = fopen(fichier_image, "rb"); + + fgets(buffer, SIZE_LINE_TEXT, file); /* P5 */ + /* on saute les lignes de commentaires */ + fgets(buffer, SIZE_LINE_TEXT, file); + while ((buffer[0] == '#')|(buffer[0] == '\n')) + fgets(buffer, SIZE_LINE_TEXT, file); + /* derniere ligne lue : dimensions */ + fgets(buffer, SIZE_LINE_TEXT, file); /* dynamique */ + + /* data */ + + if (nb_level < 256) + { + // fichier en char, on converti au format int + ligne = (unsigned char*)malloc(sizeof(unsigned char)*j_dim) ; + + for (i=0;i +#include +#include "lib_lniv_common.h" + + +/** + * + * \brief calcul du PSNR entre deux images int + * \author NB - PhyTI + * + * \param[in] im1 image d'entree + * \param[in] im2 image d'entree + * \param[in] idim + * \param[in] jdim + * \param[in] range dynamique de l'image + * + */ +double psnr_image_int(int **im1, int **im2, int idim, int jdim, int range) +{ + int n ; + double eqm = 0.0 ; + double err, eq = 0.0 ; + for (n=0; ndim_wind2+1)&&(i<(i_dim-dim_wind2-1))) + if (j==dim_wind2+1) + j=j_dim-dim_wind2-1; + + result[i][j] = 0; + for (iw=i-dim_wind2;iw=0)&&(jw>=0)&&(iw +#include "defines.h" +#include "lib_math.h" + +/** + * \fn void tic(struct timeval* temps, char* texte) + * \brief Initialise le compteur de temps + * + * \param[out] temps + * \param[in] texte texte a afficher + * + */ +void tic(struct timeval* temps, char* texte) +{ + gettimeofday(temps, NULL); + + if (texte != NULL) + printf("%s\n", texte) ; +} + +/** + * \fn double toc(struct timeval start, char* texte) + * \brief Calcule le temps ecoule + * + * \param[in] start temps de debut du chrono + * \param[in] texte texte a afficher + * + * \return le temps ecoule entre tic et toc + */ +double toc(struct timeval start, char* texte) +{ + struct timeval end ; + double elapsed ; + + gettimeofday(&end, NULL); + + elapsed = (double)(end.tv_sec-start.tv_sec) + + 0.000001*(double)(end.tv_usec-start.tv_usec); + if (texte != NULL) + printf("%s : %f\n", texte, elapsed) ; + + return elapsed ; +} + + +/** + * \fn void min_max_int1d(int *val_min, int *val_max, int *vect, int dim) + * \brief determine le min et max d'un vecteur de int + * + * \param[out] val_min + * \param[out] val_max + * \param[in] vect + * \param[in] dim dimension du vecteur + * + */ + +void min_max_int1d(int *val_min, int *val_max, int *vect, int dim) +{ + int n, min, max ; + + min = vect[1]; + max = min; + + for (n=0;n max) max = vect[n]; + if (vect[n] < min) min = vect[n]; + } + + *val_min = min ; + *val_max = max ; +} + +void min_max_ushort1d(int *val_min, int *val_max, unsigned short *vect, int dim) +{ + int n ; + unsigned short min, max ; + + min = vect[1]; + max = min; + + for (n=0;n max) max = vect[n]; + if (vect[n] < min) min = vect[n]; + } + + *val_min = min ; + *val_max = max ; +} + + + + +/** + * \fn inline int test_inf(double arg1, double arg2) + * + * \brief test (arg1 < arg2) inferieur a avec pourcentage minimum + * + * \param[in] arg1 + * \param[in] arg2 + * + * return test + */ +inline int test_inf(double arg1, double arg2) +{ + if (arg2 > 0) + return arg1 < (arg2*COEF_DECROI) ; + else + return arg1 < (arg2*INV_COEF_DECROI) ; +} + + + + + +/** + * \fn inline int sign_diff_ou_egal_zero(int val1, int val2) + * + * \brief fonction qui test si les arguments sont de signes differents ou nuls + * \author NB - PhyTI + * + * \param[in] val1 + * \param[in] val2 + * + * \return le test 0/1 + * + */ +inline int sign_diff_ou_egal_zero(int val1, int val2) +{ + if (val1 > 0) + { + if (val2 > 0) return 0 ; + else return 1 ; + } + else + if (val1 < 0) + { + if (val2 < 0) return 0 ; + else return 1 ; + } + else + return 1 ;/* val1 == 0 */ +} + +/** + * \fn inline int sign_diff_strict(int val1, int val2) + * + * \brief fonction qui test si les arguments sont de signes differents strictement + * \author NB - PhyTI + * + * \param[in] val1 + * \param[in] val2 + * + * \return le test 0/1 + * + */ +inline int sign_diff_strict(int val1, int val2) +{ + if (val1 > 0) + { + if (val2 >= 0) return 0 ; + else return 1 ; + } + else + if (val1 < 0) + { + if (val2 <= 0) return 0 ; + else return 1 ; + } + else + return 0 ;/* val1 == 0 */ +} + + + +/** + * \fn inline int sinus_triangle(int Ai, int Aj, int Bi, int Bj, int Ci, int Cj) + * + * \brief calcul le "sinus" de l'angle du triangle ABC + * \author NB - PhyTI + * + * \param[in] Ai les coordonnees + * \param[in] Aj + * \param[in] Bi + * \param[in] Bj + * \param[in] Ci + * \param[in] Cj + * + * \return le sinus non normalise + * + * Cette fonction est utile pour determiner si un triangle ABC + * est donne dans l'ordre trigo. + * Signe > 0: sens trigo, + * signe < 0: sens antitrigo + * = 0: plat + */ +inline int sinus_triangle(int Ai, int Aj, int Bi, int Bj, int Ci, int Cj) +{ + return (((Bi-Ai)*(Cj-Aj)) - ((Ci-Ai)*(Bj-Aj))) ; +} + + +/** + * \fn void recopie_vecteur(int *in, int *out, int dim) + * + * \brief recopie le vecteur out vers in + * \author NB - PhyTI + * + * \param[in] in vecteur d'entree + * \param[out] out vecteur recopier + * \param[in] dim longueur du vecteur + */ +void recopie_vecteur(int *in, int *out, int dim) +{ + int n ; + for (n=0; n + +void tic(struct timeval* temps, char* texte) ; +double toc(struct timeval start, char* texte) ; + +void min_max_int1d(int *val_min, int *val_max, int *vect, int dim) ; +void min_max_ushort1d(int *val_min, int *val_max, unsigned short *vect, int dim) ; + +#define min(a,b) ((a)<(b)) ? (a) : (b) +#define max(a,b) ((a)>(b)) ? (a) : (b) + + +inline int test_inf(double arg1, double arg2); + +inline int sign_diff_ou_egal_zero(int val1, int val2); +inline int sign_diff_strict(int val1, int val2); + +inline int sinus_triangle(int Ai, int Aj, int Bi, int Bj, int Ci, int Cj); + +void recopie_vecteur(int *in, int *out, int dim) ; + + +#endif diff --git a/lniv.c b/lniv.c new file mode 100644 index 0000000..6757945 --- /dev/null +++ b/lniv.c @@ -0,0 +1,186 @@ +/** + * \file nliv.c + * \brief test de reduction par MV et contriante de ligne de niveaux + * \author NB - PhyTI + * \version x.x + * \date 6 mai 2011 + * + */ + +// protection +#ifdef __PROTECT +#include +#endif + + +#include +#ifdef __MULTI_THREAD +#include +#endif + +#include "lib_alloc.h" +#include "lib_images.h" +#include "lib_math.h" +#include "lib_pretraitement.h" + +#include "lib_lniv.h" + + +int main(int argc, char **argv) +{ +// protection light +#ifdef __PROTECT + if (ptrace(PTRACE_TRACEME, 0, 1, 0) < 0) return(0) ; +#endif + + /* declaration des variables */ + int ret ; + int Bin_file = 0 ; + int Prof ; /* profondeur en octets */ + int I_dim ; /* hauteur de l'image */ + int J_dim ; /* largeur de l'image */ + int Nb_level ; /* dynamique de l'image */ + char *File_name ; + char *PARAM = NULL ; + + /* images */ + int **Image_in1, **Image_out1 ; + int **Image_in2=NULL, **Image_in3=NULL ; + int **Image_out2=NULL, **Image_out3=NULL ; + + + /* variables de calculs */ + struct timeval chrono ; + + /* debug : affichage snake */ + int Verbose = 1 ; + int Display = 1 ; + + + /* lecture argument entree (basique!) */ + if (argc == 1) + { + fprintf(stderr, "USAGE : LNIV pgm_file \n") ; + fprintf(stderr, + "\n" + "\n" + "\n" + "\n") ; + return(0) ; + } + File_name = argv[1] ; + + /* verif type image (pgm 8/16) */ + ret = type_image_ppm(&Prof, &I_dim, &J_dim, &Nb_level, File_name) ; + Nb_level++ ; + + if (ret == 0) + { + /* tentative image bin */ + ret = load_dim_image_bin(&I_dim, &J_dim, File_name) ; + if (ret != 1) + { + printf("format non pris en charge ... exit\n") ; + return(0) ; + } + Nb_level = 65536 ; // quantif 16 bits pour les bin + Bin_file = 1 ; + Prof = 1 ; + } + + /* infos */ + if (Verbose) + { + printf("Image : %s\n", File_name) ; + printf("lecture OK : %d\n", ret) ; + printf("Image (%d x %d) pixels\n", I_dim, J_dim) ; + printf("Dynamique : %d\n", Nb_level) ; + printf("Canaux : %d\n", Prof) ; + } ; + + /* Allocation */ + Image_in1 = new_matrix_int(I_dim, J_dim) ; + Image_out1 = new_matrix_int(I_dim, J_dim) ; + if (Prof == 3) + { + Image_in2 = new_matrix_int(I_dim, J_dim) ; + Image_out2 = new_matrix_int(I_dim, J_dim) ; + Image_in3 = new_matrix_int(I_dim, J_dim) ; + Image_out3 = new_matrix_int(I_dim, J_dim) ; + } + + /* chargement image d'entree */ + tic(&chrono, NULL) ; + if (Bin_file) + load_bin2int(Image_in1, I_dim, J_dim, Nb_level, File_name) ; + else + if (Prof == 1) + load_pgm2int(Image_in1, I_dim, J_dim, Nb_level, File_name) ; + else /* (Prof == 3) */ + { + load_ppm2int(Image_in1, Image_in2, Image_in3, I_dim, J_dim, File_name ) ; + //RGB2YUV(Image_in1[0], Image_in2[0], Image_in3[0], I_dim*J_dim) ; + } + + toc(chrono, "temps chargement image") ; + if (Display) image16(Image_in1, I_dim, J_dim) ; + if ((Display)&&(Prof == 3)) image16(Image_in2, I_dim, J_dim) ; + if ((Display)&&(Prof == 3)) image16(Image_in3, I_dim, J_dim) ; + + /* sequence de parametre */ + if (argc >= 3) PARAM = argv[2] ; + + double poids = 15.0 ; + int nb_iter = 15 ; + if (Prof == 1) + { + if (Verbose) + { + printf("poids contrainte : %.3f\n", poids) ; + printf("nombre d'iteration GN : %d\n", nb_iter) ; + } + tic(&chrono, NULL) ; + mv_gaussien_lniv(Image_in1, Image_out1, I_dim, J_dim, poids, nb_iter) ; + toc(chrono, "temps lniv_image") ; + + wimage16(Image_out1, I_dim, J_dim, "output.pgm") ; + + } + else + { + if (Verbose) + { + printf("image couleur (YUV)\n") ; + printf("poids contrainte : %.3f\n", poids) ; + printf("nombre d'iteration GN : %d\n", nb_iter) ; + } + tic(&chrono, NULL) ; + + mv_gaussien_lniv(Image_in1, Image_out1, I_dim, J_dim, poids, nb_iter) ; + mv_gaussien_lniv(Image_in2, Image_out2, I_dim, J_dim, poids, nb_iter) ; + mv_gaussien_lniv(Image_in3, Image_out3, I_dim, J_dim, poids, nb_iter) ; + + toc(chrono, "temps lniv_image") ; + + image16(Image_out1, I_dim, J_dim) ; + image16(Image_out2, I_dim, J_dim) ; + image16(Image_out3, I_dim, J_dim) ; + write_intRGB2ppm8(Image_out1, Image_out2, Image_out3, I_dim, J_dim, "output.ppm") ; + + } + + + /* Delete */ + del_matrix_int(Image_in1, I_dim) ; + del_matrix_int(Image_out1, I_dim) ; + if (Prof == 3) + { + del_matrix_int(Image_in2, I_dim) ; + del_matrix_int(Image_out2, I_dim) ; + del_matrix_int(Image_in3, I_dim) ; + del_matrix_int(Image_out3, I_dim) ; + } + + + return 1 ; +} diff --git a/main.cu b/main.cu new file mode 100644 index 0000000..d9db4f8 --- /dev/null +++ b/main.cu @@ -0,0 +1,198 @@ +// libs C +#include +#include +#include +#include + +#include "lib_lniv.h" + +// libs NV +#include +#include + +// lib spec +#include "defines.h" +#include "levelines_common.h" + +#include "levelines_kernels.cu" + + +__global__ void kernel_debil(unsigned int * ptr1, unsigned int * ptr2, unsigned int L, int val){ + + unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; + unsigned int j = blockIdx.y*blockDim.y + threadIdx.y; + unsigned int pos = i*L +j ; + + ptr2[pos] = val - ptr1[pos] ; + +} + +int main(int argc, char **argv){ + + + //float coef_regul = atof( argv[1] ) ; + + unsigned int timer ; + cutilCheckError( cutCreateTimer(&timer) ); + cutilCheckError( cutResetTimer(timer) ); + /***************************** + * CHARGEMENT IMAGE + *****************************/ + char* image_path = argv[argc-1]; + char* image_out = "./image_out.pgm" ; + unsigned int * h_data = NULL ; + unsigned int * h_data_out = NULL ; + unsigned int H, L, size; + + cutilCheckError( cutStartTimer(timer) ); + cutilCheckError( cutLoadPGMi(image_path, &h_data, &L, &H)); + cutilCheckError( cutStopTimer(timer) ); + + size = H * L * sizeof( unsigned int ); + printf("Loaded %d x %d = %d pixels from '%s' en %f ms,\n", L, H, size, image_path, cutGetTimerValue(timer)); + + + //essai alloc mapped + /* + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + unsigned int * h_ptr1, * d_ptr1 ; + unsigned int * h_ptr2, * d_ptr2 ; + int h = ; + int l = h ; + int mem = h*l*sizeof(unsigned int) ; + cutilSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost)); + cutilCheckError( cutStopTimer(timer) ); + printf("Temps set flag Mapped : %f ms\n", cutGetTimerValue(timer)) ; + + cutilCheckError( cutStartTimer(timer) ); + cutilSafeCall(cudaHostAlloc((void **)&h_ptr1, mem, cudaHostAllocMapped)); + cutilSafeCall(cudaHostAlloc((void **)&h_ptr2, mem, cudaHostAllocMapped)); + cutilCheckError( cutStopTimer(timer) ); + printf("Temps cumul alloc Mapped : %f ms\n", cutGetTimerValue(timer)) ; + + for (int i = 0; i>>(d_ptr1, d_ptr2, l, 255) ; + + cutilCheckError( cutStopTimer(timer) ); + printf("Temps total Mapped : %f ms\n", cutGetTimerValue(timer)) ; + + char * image_1 = "./image_1.pgm" ; + char * image_2 = "./image_2.pgm" ; + + cutilCheckError( cutSavePGMi(image_1, h_ptr1, l, h) ) ; + cutilCheckError( cutSavePGMi(image_2, h_ptr2, l, h) ) ; + */ + /***************************** + * FIN CHARGEMENT IMAGE + *****************************/ + + + + // use device with highest Gflops/s + cudaSetDevice( cutGetMaxGflopsDeviceId() ); + + + /* + cutilSafeCall( cudaMallocArray(&a_Src, &floatTex, imageW, imageH) ); + cutilSafeCall( cudaMalloc((void **)&d_Output, imageW * imageH * sizeof(float)) ); + cutilSafeCall( cudaThreadSynchronize() ); + cutilCheckError( cutResetTimer(hTimer) ); + cutilCheckError( cutStartTimer(hTimer) ); + + cutilSafeCall( cudaThreadSynchronize() ); + cutilCheckError( cutStopTimer(hTimer) ); + gpuTime = cutGetTimerValue(hTimer) / (float)iterations; + */ + + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + // allocation mem GPU + unsigned int * d_directions =NULL ; + unsigned int * d_lniv, * d_estim = NULL ; + + cutilSafeCall( cudaMalloc( (void**) &d_directions, size)) ; + cutilSafeCall( cudaMalloc( (void**) &d_lniv, size ) ); + cutilSafeCall( cudaMalloc( (void**) &d_estim, size ) ); + + + // allocate array and copy image data + cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindUnsigned); + cudaArray * array_img_in, *array_img_estim, *array_img_lniv; + cutilSafeCall( cudaMallocArray( &array_img_in, &channelDesc, L, H )); + cutilSafeCall( cudaMemcpyToArray( array_img_in, 0, 0, h_data, size, cudaMemcpyHostToDevice)) ; + cutilSafeCall( cudaBindTextureToArray( tex_img_in, array_img_in, channelDesc)); + cutilCheckError( cutStopTimer(timer) ); + + cutilSafeCall( cudaMallocArray( &array_img_estim, &channelDesc, L, H )); + cutilSafeCall( cudaBindTextureToArray( tex_img_estim, array_img_estim, channelDesc)); + + cutilSafeCall( cudaMallocArray( &array_img_lniv, &channelDesc, L, H )); + cutilSafeCall( cudaBindTextureToArray( tex_img_lniv, array_img_lniv, channelDesc)); + + printf("Temps alloc + transferts en Textures : %f ms\n", cutGetTimerValue(timer)) ; + /***************************** + * APPELS KERNELS et chronos + *****************************/ + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + + unsigned int iter , nb_iter = 15 ; + unsigned int poids = 15 ; + dim3 dimBlock(8,8,1) ; + dim3 dimGrid( H / dimBlock.x, L / dimBlock.y, 1 ) ; + unsigned int smem_size = dimBlock.x * dimBlock.y * sizeof(unsigned int) ; + // init image estimee avec image_in + kernel_init_estim_from_img_in<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, 7); + + printf("Grille : %d x %d de Blocs : %d x %d - Shared mem : %d octets\n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y, smem_size) ; + + for ( iter =0 ; iter < nb_iter ; iter++ ) + { + cutilSafeCall( cudaMemcpyToArray( array_img_estim, 0, 0, d_estim, size, cudaMemcpyDeviceToDevice)) ; + kernel_levelines_texture<<< dimGrid, dimBlock, 0 >>>( d_lniv, L, H ); + cutilSafeCall( cudaMemcpyToArray( array_img_lniv, 0, 0, d_lniv, size, cudaMemcpyDeviceToDevice)) ; + kernel_estim_next_step_texture<<< dimGrid, dimBlock, 0 >>>(d_estim, L, H, poids) ; + } + + cudaThreadSynchronize(); + + cutilCheckError( cutStopTimer(timer) ); + printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)/(float)nb_iter) ; + printf("Total pour %d kernels : %f ms\n", nb_iter, cutGetTimerValue(timer)) ; + + /************************** + * VERIFS + **************************/ + //trace des lniv sur grille de 'pas x pas' + //kernel_trace_levelines<<< dimGrid, dimBlock, 0 >>>(d_data, d_directions, d_data2, L, H, 16, 255) ; + //cudaThreadSynchronize(); + + // enregistrement image lniv GPU + h_data_out = new unsigned int[H*L] ; + if ( h_data_out != NULL) + cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) ); + else + printf("Echec allocation mem CPU\n"); + + cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ; + + // calcul lniv CPU + + + // TODO verifier pourquoi les deux lignes suivantes produisent une erreur + //cutilExit(argc, argv); + //cudaThreadExit(); + return EXIT_SUCCESS ; +} diff --git a/main_gmem.cu b/main_gmem.cu new file mode 100644 index 0000000..a33fbff --- /dev/null +++ b/main_gmem.cu @@ -0,0 +1,143 @@ +// libs C +#include +#include +#include +#include + +#include "lib_lniv.h" + +// libs NV +#include +#include + +// lib spec +#include "defines.h" +#include "levelines_common.h" + +#include "levelines_kernels.cu" + + +int main(int argc, char **argv){ + + + //float coef_regul = atof( argv[1] ) ; + + unsigned int timer ; + float time_cumul = 0.0 ; + cutilCheckError( cutCreateTimer(&timer) ); + cutilCheckError( cutResetTimer(timer) ); + + /***************************** + * CHARGEMENT IMAGE + *****************************/ + char* image_path = argv[argc-1]; + char* image_out = "./image_out.pgm" ; + unsigned int * h_data = NULL ; + unsigned int * h_data_out = NULL ; + unsigned int H, L, size; + + cutilCheckError( cutStartTimer(timer) ); + cutilCheckError( cutLoadPGMi(image_path, &h_data, &L, &H)); + cutilCheckError( cutStopTimer(timer) ); + + size = H * L * sizeof( unsigned int ); + printf("Loaded %d x %d = %d pixels from '%s' en %f ms,\n", L, H, size, image_path, cutGetTimerValue(timer)); + time_cumul += cutGetTimerValue(timer) ; + /***************************** + * FIN CHARGEMENT IMAGE + *****************************/ + + + + // use device with highest Gflops/s + cudaSetDevice( cutGetMaxGflopsDeviceId() ); + + /* + cutilSafeCall( cudaMallocArray(&a_Src, &floatTex, imageW, imageH) ); + cutilSafeCall( cudaMalloc((void **)&d_Output, imageW * imageH * sizeof(float)) ); + cutilSafeCall( cudaThreadSynchronize() ); + cutilCheckError( cutResetTimer(hTimer) ); + cutilCheckError( cutStartTimer(hTimer) ); + + cutilSafeCall( cudaThreadSynchronize() ); + cutilCheckError( cutStopTimer(hTimer) ); + gpuTime = cutGetTimerValue(hTimer) / (float)iterations; + */ + + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + // allocation mem GPU + unsigned int * d_directions =NULL ; + unsigned int * d_lniv, * d_estim, * d_data ; + cutilSafeCall( cudaMalloc( (void**) &d_directions, size)) ; + cutilSafeCall( cudaMalloc( (void**) &d_lniv, size ) ); + cutilSafeCall( cudaMalloc( (void**) &d_estim, size ) ); + cutilSafeCall( cudaMalloc( (void**) &d_data, size ) ); + cutilCheckError( cutStopTimer(timer) ); + printf("Temps alloc global mem : %f ms\n", cutGetTimerValue(timer)) ; + time_cumul += cutGetTimerValue(timer) ; + + // transfert data -> GPU global mem + cutilCheckError( cutStartTimer(timer) ); + cutilSafeCall( cudaMemcpy( d_data , h_data, size, cudaMemcpyHostToDevice) ); + cutilCheckError( cutStopTimer(timer) ); + printf("Temps transferts en global mem : %f ms\n", cutGetTimerValue(timer)) ; + time_cumul += cutGetTimerValue(timer) ; + /***************************** + * APPELS KERNELS et chronos + *****************************/ + cutilCheckError( cutResetTimer(timer) ); + cutilCheckError( cutStartTimer(timer) ); + + unsigned int iter , nb_iter = 15 ; + unsigned int poids = 15 ; + dim3 dimBlock(8,8,1) ; + dim3 dimGrid( H / dimBlock.x, L / dimBlock.y, 1 ) ; + unsigned int smem_size = dimBlock.x * dimBlock.y * sizeof(unsigned int) ; + // init image estimee avec image_in + kernel_init_estim_from_img_in_global_mem<<< dimGrid, dimBlock, 0 >>>(d_data, d_estim, L, H, 7); + cutilCheckError( cutStopTimer(timer) ); + printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)) ; + time_cumul += cutGetTimerValue(timer) ; + + + // iterations + cutilCheckError( cutStartTimer(timer) ); + printf("Grille : %d x %d de Blocs : %d x %d - Shared mem : %d octets\n", dimGrid.x, dimGrid.y, dimBlock.x, dimBlock.y, smem_size) ; + for ( iter =0 ; iter < nb_iter ; iter++ ) + { + kernel_levelines_global_mem<<< dimGrid, dimBlock, 0 >>>( d_estim, d_lniv, L, H ); + kernel_estim_next_step_global_mem<<< dimGrid, dimBlock, 0 >>>(d_estim, d_lniv, d_data, L, H, poids) ; + } + cutilCheckError( cutStopTimer(timer) ); + printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)) ; + time_cumul += cutGetTimerValue(timer) ; + + printf("Execution moy par kernel : %f ms\n", cutGetTimerValue(timer)/(float)nb_iter) ; + printf("Total pour %d kernels : %f ms\n", nb_iter, cutGetTimerValue(timer)) ; + printf("Total execution : %f ms\n", time_cumul) ; + + /************************** + * VERIFS + **************************/ + //trace des lniv sur grille de 'pas x pas' + //kernel_trace_levelines<<< dimGrid, dimBlock, 0 >>>(d_data, d_directions, d_data2, L, H, 16, 255) ; + //cudaThreadSynchronize(); + + // enregistrement image lniv GPU + h_data_out = new unsigned int[H*L] ; + if ( h_data_out != NULL) + cutilSafeCall( cudaMemcpy(h_data_out , d_estim, size, cudaMemcpyDeviceToHost) ); + else + printf("Echec allocation mem CPU\n"); + + cutilCheckError( cutSavePGMi(image_out, h_data_out, L, H) ) ; + + // calcul lniv CPU + + + // TODO verifier pourquoi les deux lignes suivantes produisent une erreur + //cutilExit(argc, argv); + //cudaThreadExit(); + return EXIT_SUCCESS ; +} diff --git a/obj/release/levelines_kernels.cu.o b/obj/release/levelines_kernels.cu.o new file mode 100644 index 0000000000000000000000000000000000000000..69cdd7a017b1f4c226686c8f1203a4c9349589b4 GIT binary patch literal 20176 zcmd^`3wTu3wSW%L=>GK2UCHMGQi(BK8(hD_*@7AFUuhBH9}3gS*x~YtO9g%t#x1 zzkc7nJKs5H?|-ko_S$=|z0aO=PL@w-nAo>ZA5Dfn+E^`{l2p?U56i|45jC_kwL#j* zFCCM=T-UY7ao#FHAw7|&{Ug8hgM?p*)N7hZPt2x4{&Fr4={lYk@|S1jMY@iqSH4^J z1lD{6Pa@`DcS;8MH(ZSIU*c(kYu~sXCm{IOJ+lje|Ec=Y!7@@JBZYRPUnyRNR{zSq zWzfKOkzp%5igCwzhUSO*{p*&MmiRZqX`NMCylzHm;Vx9-I*1kMC@Ive?y!nT6_ndB zqg2N-#k=}}fPdwag(Q@IC6iYu^NMBO6Lt%q5-rSVKDIS1`=1Pho(6>dz4G0{pT2aY>tzkub)KF}%5cC367+%%hd7HR~sXF(i}dEZ0}AiW!o=O8W0K;D_i;kx^{ zV;03{$vlulK96{&4DtURd2A+hf3}(AXF1yP9JVLh25FzPSG0+>McOOc#P~m2-frvB z&M!`!aGVEYbS2p~a9(N2afkE)Xe+uXFpv(rd@;UVF`WxU3ohRJ+=n6zIf6<4slg#enqrdtR zGuyEPPd5K?TyV{Ukt()l>*@~BQe5g-eYHIy#b7zUzm%s|{-Lj!tajdYzpzKmHiSl6 zcmbFTqizo|u2$Sya)STA;|og#bURMx`3In5o)qgbM9_iJ;c;or+@~Dtm&2gKv=H6q z{}qlcuE~D5J;5}C^|&Oxf5|?G$>~SK34|6v7*tpdarHZJY_nXiGa~{KoS=Wxu+eIq znR--_Sh*;k$veraE2r!~R=&rOY?57K=gix!-j7c zc_Ft0uzp%zJ2!J*#{q)M^||04aQ}zghSEX-z)kjHuJD^~rJSHXeg2Yf0VHq0Lf~gG z!uD}&tK9#8gk)%$T>TA{uF|sAx2Wj#w+gP)LV-y)hk$S~?iuu5?wF zS9;4UYx<*9dGCX$q)r$bi8bpLR!ahpwy+hbgy0TURfz~D%4-5qv(=17VrD9^$V|q} zXv!H%Bs|WDUM5P0LgO#_mNVr^x@Jxv{&g+;y5(^QrX(yFu+*r(hc11F81PXEt8IjS z$ra;9(a2ATHARB4QPGHHMa>cV)OaLjnMwVM$P#mu7aRr5(95TIiptQKmdKJInxu=K z6nzA_&6P!E^=4zbxjSxg7nMyn6Y->_2b)?WsdzGCrv487>I{Qh7P*RsjR6;J3C5br zaZJY!1E&@_?xJB|?FimnjsX4~Hr>wVL^`9$1Bw*4Zw`ejs+=CD?dF(jtMvxFRpsf} zqFDTz7@Cr9D{qdaNh}fzMbk}YMapVwGQ-Y=8Rb@=3Z~238(m_8M`8hMaiW{Dn%*d*W~#bsd!;KGfyJ+f zT_by=E0s=$gV6t~VXn&FXtT^VE1k@t4cuGxy_35~s~NK5J=^KR8KF4x$(7LnM1$#A zXkj1`Oa@y_SQ)sNxN^*plPg9iPJx}lW-}0q22jQ4pq0nIUU{@o7lQkY6^vPYFxi>->BC8m1=wC-Kf?RXo{yBqh@Yvd$sa|T6g7F zTHCAT-&PB^{A47SiiEn2ey=8xTpn`{f3K!bcW)0gHAVxlhjFoK>4sU9V`~RDa@b!w zO=#R=Jut^>@A;-)b=6gRS^uIUr&!W-XWCn(2LiZk0cZ(y6AVOJnn!S1k2}9?!(T1w zhh5N8FfAGn6@l@dqRRfI4fR(xn$2+c>}CLL{iKu^q>9Y8h&5UWo3W-IqTkbDx7MwE z|Ni#wU(=S`&;LE}d`5)6+OBTT|8FUPSu7LR9?K6fP4&CR@N^7oGyg}>mYH7wVMd&8Df`L>(`vAg()j(1sfqI zuZvqmI@aoJj>a2=7;K(QI3! zphve!u6BpyJyZ*jTN*nx+g)-Xu|RG-FQ7jPZ zw_`#%Dw1e^e8DN;q|*1Kyu3=El~-QwR46BXM#`&Kod7GR3}U0F(yXoW zx*NlurW&s|sO4IJHC@Ap2x-fNXABzN7%SD9vgtl^{%|g)WStwQP1%yo{rkttzKree zskPH5m~o?F=e_TnTw;E+<~LzqyQAXht!s|?7LC53?D@fcj6r9=-MU~>kuhN7r#C%+ zN3l_RJaFxI)^+&?ET|qopy7AEKO|F!zH`m}zVFUX@3{Trt9`Q%RNi&|g~NPpPruyw z%S-wi{Rhvu;ld3AjWMQqdd)q9j9ndXzkKUqhtd4P-qp8l|J3*SwCe+|XMX2f^QJZF zoGthJ23>ys_|~qgeRJ;o=qIB<&z+Anj=$;QN5+ml^1+H*;o7||N9=a!@A~e4?%Zo{ zUa{Zz`LA}*`KJfYGG_j3{||q*d*UCTdg!;l<2#>uX-<5lZ|I8YrT5=E+}QNRl-rj)T5p_wcYN+muUu-} zz2?&`_jVe_>x1t&-n!-6#t$y5-L~a2%`n;zMeaNJlyBK`tNMofmfH4cZ}{F?xOL0n zH)=W_e0Sp~TNW=fMtA+kv}-@wVm#O9Kev_EY%?ZJn19wMLmoEXYdmw!tJ{8RY+w6t zm;dqNl=1q^*Z!P3IKvoK)Lb+9v+IqQE*}5G{gL~PnkPP(@`2+Q#?omY|M&3DZN?{$ z-nePVgFiLaZn*odmma^G#_=Iv1ZN8r^zNK@|r$c=gy}kCJr(^rrfr%fUapa(9?DMtz z-YXyTsqYu(c5V3K7n;#=+RKl(y#9{wT(fP+yg%*q9ol?&L-|u~`@M}nfu|*7q9jd())f?=R=5dPA5A+0j#850OC0hsdH4+C zc<+?~<)f5P`6}Y7{@e5LNAu{xcf1%>{cqV=i{Q0>fLKk?p3jNXvMN%@&k!xwv^7YW z{0QQF5l8uY;uqTT;+Bkj9&ypflLlnCMf_##>-qr(7DM^#^T^*tyc*^GS^Ib9kjJ`S zApRM|wQ&PVpb0DG{sHUxI7d0;#U%7Lma8A2;{|ysao!GXQeE9By=>CtnZtGQWJ`nn z0qqMi`-IHSA!$Z~Vo#tsMYJhs1vO`5Dh27dov^RC$8b4aB!-VAPy|cIS|O!%>j4<9rq{*o3&VsXoKg30CZn>DluoA9Pv?GRa zR!y6vaB}_3hH*ONBU9~POXX#Ir5ZJTIb`JOza!Lta-6XLu&gXE`Fvox>fcQDw?jS; zI!{1+HXOO?*BWfWtt5~27#P;ChQY&ZqV|g_ zdnbsg3`<@}?Wy~*~v0&Gz4776L6^Qhi!e% z`u#XDAr$hP9i-Lzg;0(Y%JtJeJdq;{>dURAAg3j0pJD_l&Gj>Dl*m(8V@;bOIjt+) zt^#;cp$6P5^aVTHDMRcXAifWdK4K=>8SVB(`#oGyzLV0j8ti~R%zoMbC{~ae4Jrhe zv!H8GUdqX zU3U>*Ag@F?-A3QrO*QFt5iSqi_F_+W*v68xK)`?qYK{T#Gd>yUu1H%dJUzlHch zh5v-OuJAjEClr1+@iK+qN4!np4-v0a_;%v$3V)1vt-^N`U#sw^h#Lz34e?D1f1bEs z;V%>ak-}diK2_oP;Uqp7+^q0p#Ahk|L*iQ%{t58~3jc!mLkcetx13^ahr)}9A5#3s zLHw}7&m=DQX;R$~;ztzup~O2CegScN-b*c-R!+Q2k@pZ+_m?Ay+v{IekB_(^`XkFH z5HD2tWyI}uA}c?Qc!?r^CGn{W|1NQRz0B$f5uc^VM~LeRj}y1om8_mL@q{A3l=wo0 z+wUEUHGBQY>RChb?TY+H;%gOtGx1Fd|1t3&DI9Am=&|}}+8rQ2j4hspWYW)xmrz>r zt;BWW^fOs|fOw_Cw-YyrOFh3NF6X1`Q22h}(niV>mpAk9cYqIq_HH77 zmOJCFJn|0sag2%5o`Hd$1aMunr28pN@&`zs ze&%Y+h|75`o{d4#m3j26$-_5F{nRm1|E*F#aq%n-^#3%Ep1bq#dr40v{&0dI^=~6S zmAH8J2KhVk=;_GA_vYcxF5p&-%d5H}dFtI}d-C^emtW zB=vs;JeR*|dB3l8<^8PAh2L3f0r<`gKV6#vuM`8N**8O_N{)C@1aMA6?Z`kNWkmuI z@csOBct4(kPmsn%qk(F7AfF+Ds-A`fzyMbu5((5+H@T`j)itfC-&Gw5go9Rn#w>(x zuU*^BSd!|4hn2zP;+a+?DoVk37{VNQj9Hh45A^)-31ZZQqI29evW;MNy0QEFdHnET zf(I%THM;}&6~hGl@WM<9QiWPU9cN%)6Mk4ptWpSq4l^1isZIvJ%VNaDH9$`i=N0l^Az5R`g``)=)Clgi4YRqh ztVZN|g=sam9${OJohxjsv2%rOH6piG#A`)vt*EP3=&u!;YlY5QuhtTYX=yw~qatcw zT7qpr;RH#H+L!LL%6=*jj;;mar9XTI(EYW&B1~V$!PLc>nVl zV>}jw52Y6eS}V2Hc+-G@%P;@-WWAvFv%EHl!@pyRfX_!UpgeAgIxzv^U+jxuu;gas zZ)1Eh<99O72Pm0=liyghOIS1|eIkSFbVnDG^iKhF4BjK9KoJL8`-em3L%umd3A zh3lEf_z=jG_Q>A>n7xW|i^=0XSIS?(LGmh`7q@Dd>zl#C0FMvbxV#doDKZo(*jF&QgA>&ny zU&Qzr#z!#T#CSR57UNFFmoZ+!_$tOL8Q;V>zH^i9y`6D4k*J^1`m~k$@fpLS$-@*7e#_wT#JmU{CUdQ-jjMp>%6yp;Zf1dG)jK9YC zB*u?1?q~cH#xG&K;1ozez;U@04%t5r#xG-h2;&WmU%)uNLzQ|wj8A6V$M_V+FJpWv z5qFJhege+}c@|LYj%{y&9r?*DTc=l&P#6OCC->X48RvdGjB&2t%lItNCG8)}_-w{6Wqc0f(;1)3cq8NU7{7|~ z?=rrG@%fBj$9RD8A27av@t-grWc)tH8yVlmc!=>Q8E<0z8OBY<|BLZ3<8LtD%=mG} z7c%}A#v_cMQ3MGHIPQELM>D<#^5i(qV|*>+F~$wXuVMUp#@8{vk?|)Ozk%`R7{7|~ zPZ?jtc)$LTfPn47-+|1sly zT=p@}`}0M{`R}Ip8NZq7|1;w)jDN=XEsPg@0}>F_^`VOK805=-9>X}de-Yz9VDeq4 z90KZJ28Yyx-%nsb{yjJ(AIip^-}g}#0~SMRkM%Rpd zu?_Gy1qS4}#WUdG0Xf>*PH-??k)u5WD6P{~^7AOo>qU9|je&vJTVNw`#@`4Sq&+2e zGGmwII`KBTB1b(sr3cX!hhIJ$K_u-#j_s1cporwM-dPIQDV<}-w}nkQ(ANw#HIb6#N~T?$)#K+b{GW7<@+P~eSzegNp686FW=Y5=PgoR vzORw`BtJk-x>eEBNxVbh@_ou7h0FITT`I?lAe7Lc%6jE{jWUJH_Z0sF_2S|= literal 0 HcmV?d00001 diff --git a/obj/release/main.cpp.o b/obj/release/main.cpp.o new file mode 100644 index 0000000000000000000000000000000000000000..da796c7262d46a63d1c4c53e90d2d469925f1324 GIT binary patch literal 3040 zcmbVOO>7%Q6dq@j)MN|VZL?QyeUd#&Ad za8Q-1awxKlDF=iiK;p)=TnmUsN~_$tAS47Fkcc3M_JTMd!h5qb&U(A%&?oJ@nQy-D z&F|aY%d>MY_k}`4ScpuMy_P@;iAsC*8CIPk2S|i*_w99&c`d)IHFjjJwfOJfT6<0z z(b^d$sVs@hDCSJPKnt=wT9E5lRE;_<*>ce`%ktQ$ zJvPCaBU})AH+_bL>cjikxk`ik!ymC@pMZXabG|@-GBSKJ68%7G_a!58JTjh&Z1gt< zuJ&^d#>?>fr|9F6;Z$U@0jqK?l}Jp<(a zI<)Q(wV*pXQ7_mwQ9G!bJ9p%0zUHK;b6T&biY2pZC+V`0r>6?bz%x;VVPk5nyZIuI zLYHzyOD|KR=FM`MR$!ot=};Bs_nhtM`L~IB(K2gQG5h=WxOWEv)H56lb3X)r5FW9g z!pnf+fp~KZnk0<{7~5;N%F1C+9ji~kBVu{GG=2}l?6l$TLMNV;m`{|Gp8<~)PujDd z9p3ME{;}{LVa)s0zW`dc{?0{mN4(=Y!tN%oqw}vn2AF?N$V)$jCFH&9*!ppg{p)`N zFdY8~=LvU>@7n_c;GV-wgU}K@LHc!0oaZ{>#P}l4C9d}^pWNYust|H7^a6edDnaXC z4_H55385F~e-$t>7ZDTvn|wkCy&%;6+-m_CN-Y0^5Crq{8;Jt`BBSDR9O(kD!x8d>|+c z;CQ3&HL0`8?mmh5Z?oGkkVcmD{Geyw#C!4e9{hX{{H-2%y$Ak&5B$R(_$NK^&wAjb^C$4X-@JUSpu^YD{Rk$^ z8hp^YptNOHspS~dW&xqo!9Rs`>a;(J>!#dC~t?>N*!QU^iX#~OWCSmGLVgqc}b&QX>F8Hr_ zaLh}AzvjX5-6QY?5B`V;cRctZ55B^2oR2q4Ry{a>MKI(&j^nx>g-48w?>98WAA?8W t3W8vWW0nY9MGy?}2s{EWA_(Tk#rqZK`8Z>^#W=4c?#=$Q9vou{`wx@1d>H@$ literal 0 HcmV?d00001 diff --git a/obj/release/main.cu.o b/obj/release/main.cu.o new file mode 100644 index 0000000000000000000000000000000000000000..20ef00afbeab84646c758e4387eca156cdbc11f7 GIT binary patch literal 80544 zcmeHw3w&Hvwf>o8(k7*8($Ye!pr@EW)X;B1ZnkF;tU|uAX+9Hn> z%A*vxRZ;mXO1T&BMezzEDn-Qd&=%xsMO1uOMFByr_^MYc|8K3m_RN|+XEM`Jz3ROu zznOE+x7S{K?X}lld+&AjnOu7E^4FJzLY_TB-br5e&Z3@o&(qoK7+=P`BfXhk=U-}0 z`^%R757oSKGo!eFy}s}2{7V-ox|n3C=kfjPv-hAsRny~siN0t0Q?vAZzeL?%{(Qq1 zsN`qpi?6!2RAhjzt*@i&9rSf_UGUjz%{`~B&Ef>|*c3nqbTheuvTUOPD@23h!9^W#xj%$0x&h(_N z%O9To&^98gyq8j$L{dk8OnCTS>O}X4@n!KxyS6-i?@SlH+w{|2+)t5c$e;UA%_}_- z44xuc^Gfm;=-(utPX+iSAH-c%S8D}aUN?2=7PQaOEz9e|R@Nzym1<1+k!^Oqygwkk z*2qhhvGRV2@TvEbSIDWyCkU^+_b3;CAK~G9eFv|(;zBqXN^82eE-&poqy=i=*`-L4 zhLTf&K2&=_9aEu=O3{X{|Lko2r<(eQA&9Lf*T3@Ksg7);Uka4#xjWDpn?dUZwa?$N zJb#z=Tv%V)hE;zHXg}_60qyswvB$iAcvo<{esEiWeYyP&1~xO>-D;`3O|_O*DLt)F zKif<4%*^)-y9uw&_Y8Lu9?<`843>XWko$Q}bZKIc}G1h{_ypReJb^B`>&;W$6=Ld{uneU?oF9pUE+ zp2;KlUx%Km3H4{IDSkMhE?=YiWb0t%v+{DCR9&pRTqlM9Pt%*aKD+M#4LxY1w}@=h z&#kL{p~l((>PjXG3>4g&Y!xtNG->0`nroWT{54n1XCr9%GO&8nwCBOp8yr>HCe|vk z*jv%O|4jIKdulGf;m^?RHP_=C%42hWb`wG9zhpz)vqhbx;b-stOLiJir2kX9&>RUZ zm6@x7)zrG0D=sh{k{he$%6IC!%YTZwBHGBLUG`biYDN%gX|@80i&i&HOx{XMD@O3Y zw|(I%pmyrc=k0-F9%O%v3o>9bcyg_od3VhyoXUThAr< zt(^V$4~5dN$q01$Ps(aGC$Z45+^Lg4dnT_uVrGFrw)EV$c>PpNKTyx?T=mR1RmrHQ zb+j3XnA3S!tbjXp4O*bYF!sarNgv>H*5j09^kkL&u$2;L=m66i8c)#YrZ2CnWg=Tk zxiQn9uB}i6XTf;YfbRL0FLyn#9}W4|Jy-0+Hx&6+>Z5cP)maPX0ccOn7Sagvef`71fYs8+4T|swt(Z{kcMLd_jFF)yasdIMi? zy_+dc+42l5h!->WP~&tx@RPDF-^OHq(Uu<=a<>*=s>;+NTjqZN{6eLz#1%VhuJ|VI z@-(o+0dtDBONzW)7ai3b zJdXD6+Paww-W9IhjRLzae=yv&>Ay*U(Wmkxy%}t~pjncXB&G;h|DC-Wg zi2J$EZmwGgKrYXcOIxwchfT^=C+qkUr#gX>TW#&AeGM9yzStI-GjnMK7pt>^L@$!w zq0l`L3{|2N9VbVgVn&hQ!-*J{h9-oHo{P$$ZBNAMlGC<)@pDwSnmc~}+>vITyJ*Yz zE`KO=!cENV;`v*?xa9|yiR7xeV>2`9eCE=}v!q*|x#;UlrL??exgK)nfpq|vZuuR_ zza4^;OQ16yU0Z%hx7*~Ft3W3IedwF$ZzTFHzs%sD6r3_>{0hczs;dQ#b|$C2$TY6c zrTmaNY&9MSx-kytzPhUoeV6zGD*DyL3Y6rv$WDkP@_6L;h9BHXMwn+X)>F$7Mc0<^ zWZRY4d=`m$npi+!o)BaCLcD;`A+e^UJk}2&F`vtZd@h@X3nk@oSxH=!?+dhEO~BTA zxRZ;)edTE2t+yJFY&@~)g-cOzB^m%t9Ft4wg>;+keTyD=zWjmb%0bozoD6l#O?6GU zT)O4fx)z+NA?O|y9YSH7Ve=aRZ27^W9h7g;JyaZdVPyR=w+O7^yG;A+r9hM0sBPJ- zz>*YW+7>nNblYBH3Ja{jBjiS?0^mgQa+QpHXfl#7Biqeji>`C|pK5A0lQdK~+b#;H zwhbn1N{&Bz=-3HG?6kk?L;r@dB-nw&hans9`DZV<{myQ#;QOgOMJQ1>u%j+obH{yVj_|V2OB=5&!QKj30ag=h-%amu#o;k_V@>^$dfi_Ts2 zwngVX&wFs*@gm6Bxf`0V{ZXDbnA>5fPK4X>Z=gJ7 z%T9SJ#H&(%d=%Nqv|j+JQM1kP5tHG$|Jw0+mgRN$_QuJ#^i|-gyC7bt%5zLuxkYS7O)*x;=&FgXK8>p_yLi3DYTvOOD2dz42?o1LuoVisQm_*hc79;As;^LLoKpWxSLe~y3v_iBT|Glrr{Rk0 z^X#|jwvH0MN>}fot1sdT=a+J<36$TJVZ|JiU(?gIz9$M= zU)WlE2Op;v7`|@dfH{^Y4@biT@0!$(8@`ieULS zzuXk8=}As~PZVnk)cuy^1lm)#M5nmn}55i zKUOeRuKd>(X#bBNZ3@=(QvvCFqNplR|2rlBMTQN3Q+~JpXQTX}`pU zp#C=hKFCK0wf{8uw7mQU+W)AzoXsPd?LjL*H~(!?z7%6gy#~qOn?cM96686@5|kfQ z|7#_GgQ5;9|20y++x|BHc2$4WC9nRUMTJNvSN?f*reIxv+6Us+|1>Gz&EFyU&sFkM zekw~JTfs1>{x?W|Ul9kD{}|GVr(1uUzljOW5!-bWsL1eHI6#Gv}y{2i+PDCaO&{<8oC$=~~0Q!2@3 z4{}Lv{tp8S%HJpXx7*A(Q+~Jp4PY2l|Ecp##ttXHTYj7HblczN-zND<4WOxT^FLmo z|KB9#YkD6%^gXd&iBC}dw@d!Bh)50}XSe=$7U=)`CBK*fbD;cg`F|)Nzs({dN|#^23PjHv7he;);F&08(h=f*O(eerMDb1edHEUS_Z2U+l&r-BX|zuOdA#=wJdFmkln=% zadLRDKc4De8&3^T1LpEQ7Zr!vSf=W(tR9FMT?VXSzA#GS)t3j1o{l~k~^nU z$id{x6jspvYRwlVYVy1So*do7*1d2`t3EW2&4FtCCW@b3stQ0^Qg2ey7ieiE`gooO zUP(NUXD|B(QtLA}%-tL6POn>fzB#3`18%@odEP9-5S8c6;w>(dK!;Z_n!gse*~&fN zUV3(3ML@ZCOa|8%ZMi2w z0VlN_Zz2U9@4(l>O}np}+l#)lo$|uMnXY_qdF>djb@qAFD}NcA8#>MKIZ#yw z`cUuDOTE=&%Y3wXnCT9DZL;N)$>%*En|q+fV|q*H?+Zt~ec`zWrURcG>Mh?{#^sjZ zS;qOM?wcH`+c$acv11{hx^UHgDz`j*U~*N{?pae zcbj|&Yx8}@5r!`0#a`HqZawQ7lYXGOqUp(HpbaWu;65~o^#gt33_F5e znyB3mlwZ`e8+L$}*yDY%ES%9J>;N=-Ujz_7%=iWB?c@pQE?idWHQ`K0`Rtiq_{gaq z>IZV9?mIklUq(69eTc?1b3LybdCC7z->Pe;qskQ~oDx#*8s!Gh{DxN25`$-5iy%lQ5+!yWqK) zx-WC=%GeR?M(hr(kzecxUKwrcRf~_WtV5jk7;m~A1Gw*?eWovmJ;0uhJ`H-?<$Ewe zZuctj0pTlA=w`1rRJNFXXO;NS*%%XMK2NdV>e$qoRk68cQ=xaX?-7-3f7P|74_n_7 z@{o=`STPgnp!3Gc{L0HF`xX0)easS{0w3!mwV-c1SUdSg_AT^8oZ2wgUJj9;5&M`f z{spsC_7hdhssH-kfr>GVwK2>I;3F{@R=;rKT=LxqD#~RnH+V(aWf*&6UP$~+HTizn z)zPQHPZ6J*Tt1*yLlv)t>QF6_|_>UN3WqnVZMzXg*VW&8E89Wm94i?=aYZ$eUMI8Jm6e zWp(qtDwJF1bm z{CPKSYM4WI54>qAl=Qv|f7TJ8zsd_orcL7fv!+5`_^J7oW=x$U<%XnQ7!xt3W3E8; z+I0;4nW@*DsmKqxO?@#2OMRzv{pVt=G<5T49>(&|KLT>2Uh{1|W}5psRQ|5A@XRYK z{CQp|hShz~VK&(8SM zm*RiDkni~@%bk1AU0#NEL%*1w+b=49Dg9!4Xx#o)A^j^t{c9Qc%*o{wdS3X<@pAZt zq`$!bd(hKP+b>#W&V1}_^jnPk^QE77H;y%gw~f_R-r&_v{eV|nJGQA|=4G4eU>hlo$u5^+JvH0( z^9r$N=7WA({W9w3n3IM@4#_`#63Pc&Jsat}kYDDkHXr2Vbm+;Q@fGVEjkLQMy0)1^Ro02vfB3SxDv}HR!LA#ue>iHp zly{ja5Ay@Af1&bfub$*jMOhVQJbcQ|17?LyeeHZ;)+E#q{ol?Ds_fc%n(5aSVWeN> zc|4x%^7?qqhyGr^x`M}t1EHxHH@sfVv1-R2MEzzPxypx-j(Ucnn3rDLM~()Ysn%O9UAWPGZAg7opKa8-D+ zX?IyiOrJE#)I;Wb@Wb25k7oQI^+$|vp)lyk9>@>DkIoT4RJ$LcE7(FcjmuZf4IeG* z$?~zVPxYC(Z_?bU&^yy%-C^+qA6J?J6i z)ou>^voUrYJXwQs;b7lQB^~$hMed6dc8vjDx zTpm}-yo&Fb_L*~oReET%Ko4+I*~U z`th;HQ2#9Rj-6@ra;z`+v~ItuEPU)t*cD<%^)pTW`Wbw`8GImrS-4*K)Dxd&Q0ffB zr=IxmJW}t5bPp+n29^vH9pNN3Gwl% z!>4?9LuTBZM2K3?R}huD4^vKcU+%;M;5w0>UGI;#{%?TAlFhpF&uZ!46%s$P{kux~ zb%pfvko52AGQYC@yu$Wt+s|#BN9={?+sm@+k?`}Ox!wzi9V14$33lVg(hlXkn__PV zdq71V`T!r4S1xl-lx{ziV6XogXX0<=ApSvzP1k3-E_S%QEPs9eYt+f|C;3PQd+*9) z)AgCI{~CvugZR^ZcycYzcXfG$Q5?p_X#CX}!Vlcv?B06Z<{bNI5W!*3H1^7|x&r&c z?A~$<+VkC#&z^O@NyumSWLtXoOg@MXov)&?a6WoliVl0$`6eNs-2-i?^1tT^r1MqG zFPzVp^6XjXn}mFJ&vAZ92B`}w4?+~mw_5aR&pG8GpWWkIsJw=t@~|(nP`*kj&z^J2 zLq5BwworL~PbILhSxOsQ&2cMq>k zr2Ucpc%*5b-`_nHPqwwh(}VG3b9748f(3rhVE<66FVXA2Bb8p~FX-uCV73!2=pGsx zS1BrBr?CLc?wa@pZLLrEMZww18e;S>HZ=5Y)Gc#O*mSbdph}oZsxbg`@2&E zjXgs{&5bF)fr)#17N7E}#(_mcE0;!(eW}8{)cmCQ(}&V%rjg?h1uuUUH}mmB=?(Mz zQ&uiIUfTQQf!GE zTwodX45t_LCf28V67k{D0klFwF>Re7({+O*Y3#vH()MY^u5E(sxId=`;^~b;IsDou z$WQ9)vg?>AyWtd^UY>NF6J$3sI-H~p=&pE?rU^2nUCh{Lp3e|w?iyrCHnTp_lOD{k zaD+OCW8fw6llJKA9v$de7a!^#?(R>(PiWJL1azK~c$0mPK-%uLiFi+6_ec&?>NKMx z!Wxm+V8tGS>K&b^&C-YziFc<5`%^uE^^Z=}dWmy%T+S1ApDo0%iRP0%zsvB7T~ZENxEhVInB34|x0$Blkz`@;Q(DOFXC z>^%*yc$_>Wfc|)C9OAv^<(7GxQFGKkct*=#I*YEC&SDr7!goT8!-=)fNTkiLA6|E2 z+es&6Ak?jX_{5fW1)&Y1LV{xQKVpkCJCc`jqgbHCTu>+ zW-<@2iy~vs`gG%HbJQkNUYoVMx0jN9l-5kSC~Y{BpQlCgY;@$HOe}|+KpBkal+xF0 zY|)n0I?q4n;0Xu_M7H^NVH<84S zg4(E->p0iC7NjdJv{MqNQ-UjZ$AskWBDq~wFGU5TYpfzWNnA)(KTOqU-K@>%h8vS> z*@;Bx2CHGEo0e!>fLydRpFrCzKh$;1ZDy_ow~aJ&b52#-I>^pl75}|dCEHvPHdm4> zu)(O)40A0b)CbXGEhSX*C0J6UIIE5}IIF4=C1}=mNa5sqMvR+?a8pD`->Tb0SaeRx z8X|igxRxPE2(My%3w0}0tW{U6BhpE*m4;nsUTYq*YHG1rGV|qBhu7y(AIuo2E`qygP)d}|fHX%%IDX9(&s-zad)fJ$7Kp-#*sEn5C2DlregZ}u@_Vgq z%`L}I@fR#Ov2Sp2=p=tRzWp_Wy&HXsVE7&2j@5XYqF23%q4c`reemR0#k$iN%u-lz zarkBAcwd4K+1BZxflX^9YAm7yQpV=i54Gu*ZSClUs98Z)bR$SJnDj|3N1%0y?q0vI zdrg7@fZ1x~7R6>3A?4GcHayUX@mPFe()6unq$%}Bxp5*~#s^ zcTz3tlP&nP;nQ(3#**Sq1-SxDhFU_|YKLR1&1@u5l2jWBqmo8?(`jtoWaF7$I*=J* zxmz0(<92~#YuDya(pVdzuApsXB3(hrC4M+qF7_MZU&vk{SBI91`~zUj1{y~PQg*h~ z8i{74C6{bxqd9IzdL~Gb{8#tboe^$8%-{M3?VP8X`vL6P8H1X+r6`q6KR1sVp9}?Q zq&J0Wwy9V%y9n~CFilcRb{>HR0A|L$NSeTuNET2uF&bY_FJ+!$IR`3=y%g6uLIZ}f zo@TaO{%W?~6iZI4l3rGEGj)^f+}UWW*XR&gNma~Jw8SbVT|r}{;*qrdi4IouIe=SU z+w?Y8Gx->0eHgNuJ6bZvN7HsPK4;qA%w|2#yxkDwHldOFrKniR$!iLpY?^M+jyW_0 z%^J*0E=~BNjN>u3Krby#ogJ#Dbv9Fn$R8WbEQ8uBJ2db-GfA_De0nnzb;c5&q9u9m zso9{#oe>~^`gs;ZW92mJmW9J4Oit?bJQ5Zk$RD?y5xL1+fMM7Rhau1aj z?TSXTVJwO`QnX7&wrB=O!7da3lC@$dtrc?uTjR?@jvPLTLTM!*;rDh z2bg?3tZ+cg#I!U^>fFI$GU|&3!(=?hsnLYHV>V1iTq#UOVycl-b#L2+vw5f=287nQ zV{;#+DNBGeMz&$<*~;sZbUcU8Cue6yqqZU^v67DWMTw5z{@YFYX~s_ze)8Rs{P{ll z>YVw0W{MPPit>~%ySOz&6*;%8@0M{EN^*Sn__i45FjcO@wgC-l15&DLYZKMr7Uz-8 zb=Yi_5pKAQxr)PP&{I5xd;@@m8@f8xeg)k_8napkww zWQH7p0o&;t(jEQr*hC~~X&St=0(p3nsF$P@Wa5WIhmdW%OkZ->hzG$dHITz9Z@F0_t0D`|61jAh7E5hw z&5xxt@jH@THYpO(cQ=?9yiD|*u zH3!HF2byU~5fEtRB>^utXq}<7OiQ>-ZL&KdS;Ho#MzJbXN9~czRL9jbE&ZIeK#ARg zR|87XvcC)bGQwrIBJ#xD2uCY42&ZnVzq-P%TfZzN?w$-%YSfiMk zJ8-p{GjU;49g4ZTOB+)Nsewc=TXSV;!-~j{3rl$uwtDl=2i(A}iEna|y%d>6<2!*} zDi9BSD9TCFs-SyziK|iABaDk4Qa8uj+FJ2sI=fNF*<==>6}hwP3C5xKmPP5UQ0obs z&0ZZSEc-Au8zZKu03^&%xA1dtREVdG)&bEno9YuW^~p;nV@Gn-B&!xRD_Cw(UWMzV ztlcKc!jSO!WE|H3s4I=_$+BFN6&IB z4ZG8rE<>$3nVxQ}gM`X*3bexJj1{&}uX3%B_L&(g%8v zlfYt*yJ?|78vB#HtlV0Zj;#5uWj2Wz)zY>M-Oy&2h?E`JT84LUYq1#+i@eriBe1PS zsd=r{Vq2@VSl?>lzD4w?UeVTS;l5?4Tbw9ii)knRZlTs1SvPFGc`Iw4$|B|Gy24g& z8%sg?vYXdiS?z{%tBqI7_VrdMkETbM3N|J#ppxvy_Es}Av#hPuCIgAJnKV{jD`|}~ z_R0G+#!OmS4W>}unyfRjqFSlagTsw$B!el#@N_13tLvN3STF+mgz_W&kP% zidGxiji$5~+M1jvF>NeN8@FNBl-sz^QkwV?c|c*sq@Fk7)iPGz8P8=IvfcHXha$af zxB0s9~yN%o1-6mGzvP=chOIu!`;JjzxyjLtI z$T9x^?HNC|F@D~_kOsPMP9qgzF@E@=Nahlw4(*X_grV>buI833Rd5`ljRxeLU0vh; zOQe62@Iy<<;fG?aMIBHLL1st@Tne*)Fz&$k7Zr8L$itrMA3o}ksSc(Zxld%_eRKU? z5&a&zU_(YGngj>DKtv;deiAnt4ZPei*2zkA;Cu%MaVBb@7th@2X1kagS}_ffC3Q>D zC3PF-D8-l5RHnSVLq-0*za4}@3afkCFoc?*1=0S z)lBlJo~CT>%8lL;q4|EH@B$j%EC)E6qShp4$gN#a;qlItC~CkY`2tZmh+CxM0uq9o+U z7wu?{wv_bhVyT|VHB1R7N>=iUe66xy&TJQk#S%_=^5jEl(mU zrC6RJZ7yng6o^6Pcw)tY1?_!|&odM>@adqC886{V1A;;(q#_iQPX&dHw=+CBb_T1n zK_OTN2^vsMmT78H-YP?X>DwL zZ5)SDK_PZZlol8iLZctl8i}3(N+`%mf{}(MH-~2VWtrOQq`5XZ_kETSyx(uD~+qX{k{;vXbU>nEpBit+QfG2 z2DipvlxAaEmB69I zAQs;F%EqS(N3l8-m5pMV4L0%&If`F~bfQ?rh5RQ8WwpjTT3Tt4a63EtkmE**>*P@F zqR9xwgV{fPJS$TjSI7CgXiDhl=4C`RMVa^xnv7uiGhd!-OdzP0f+=yUWrx$8Ei!ie zHB%FEu&q4d(1W|PBLv-ZBH=6h1WzXIG+ekLE!em-?f%=n=_;?RV5TmCEHz!nKE?_C zxD9J6FLq@VtmYZbZtTiDxXOeWMPoGd{$socM2j9;P7h=CkFWe+dI)+8sAwRlRBiHH ziELhMTQ=J8%p|)+<;6grI%z`US|z=A}{;uxE12{wZK5R>sjyadg*8pSksLp+VUD4wLvq3A}IlcL(t%vb%E^Q%XT2y)I+0>ZYHy) zHVXzfBE!DQ4d_X*#nJY`23u64(mtu84aJesB=6n{*0u@`LfZkqO9La%#AoQLwD92g z-7Bl@yzn5ezM36(#WCDG`{k`u2wKV@FGfi1Ra`A>gl-TQk<#K>c`H?tB@EOk0*25d znJ$8P_-fCyU$W~YgvS>-td>LUuwFGGRZOk@UxC@o@hlVu+gHvsQqQs!=U zD#>mnnV{$FBBL0XafN{ad+iYt|SKDM| zMODV8lFV$0lU~nT;)*?KVOvY`zPLbbBc_SCCB4UuU?sj<^13cXw38IkPC^WP{Vkj! zZ{boJ#cQ{irsZjrN8~(>mSTjwB`?X5VO>lq&XC(Y*-Gv@Sfk_z5gf!98F5=jo<>m{ z=#fm9EIiqEh@GK{+i}1%*W0zTC7VY_9Jf>UEc-Hfp2fet?au7lkz~9bKQziq7H1fJ zt@})=oxLk%M`sRxBxwTc?aq9oodb;RS%;SI?Mx88o&A1xI>(##lCBMLL_68$44rMI zcsmXQCyho`F{Kz9$5)dS%M7-bcCwasvX*wTmiByWp?-<~w?$uMGO#D?Svq8EC)eAx zlkXaCr}0Bshw*kM&GmL{Zh6{fTg%fnSwT^6$1R?0cLpmmqvBb4L}Nn=?3*35M(rS% z*Fi3?qdBi{qBi(JT=I4h17Cj&Z`VPy3+0Aj6YzGf{lZ8Hb?g^bQ-WG}$ZdD#?9`k^ zH42TP)b@hju7jq29c0A#acX?!PyNWklVst^cJg)|huGV7@Oq>pf9lu4mTbIT2TlEy zJ!@~*L3L9PDm&RrlEt&CbDoxTvKO-ylrQ_BveP&!Lt%|q9t3uh@e;ZGI=-2m%I~YQ zf9%YE5ZD^0pD8_beZ9rJkda;Hb+Uu)q`qzh0V8HF|oVyHZ zh%6iEO^Se@M8SuO69%t#aWX(s5T?oKm_ zyhoX!LJ~2}f0oG{$j7h~kv8RyP};-HPBA%7M(hhndB- zqn>CMC(e1SX@AfzJxu?tF7s@Xe$I7e9M2@rFY;8gXnuaE>HI7J)R(&*zc`oP)Nsx( z3rv88SZEHLycYwCOH$`xvey*%y}oFXr*AmnaCI=r>e(AYeB}1Bsd-nYj}UG(kH*%p51kS@IkRnaNdVWfp&#Kd=-RT8%5(_43Tc z{MtSLW_?~+n1dX&f~c(2fcr|8@$x*UDex<+7+h#&EHC2+Y-zDSpU*!mp0RW8CQ3`D zWdj~b7uOP80*O7G=(n>p^Gur6#G|23;xUq|l{ZkD4BU#D--MSrEImKV%we5tin`kjzA-mycLf=QTVgrdnl>}1%$LtLnlil&a>qh7Gd@X0o`sB0oU!Tqff zdHf1i8M(4QfuXcTeF1W*-{;84rRyj5@6?Y~YxGM~5b#kyrIF#2u@2(}-Re%Hw$9Xv zy{lTc@Y6{8?^OYHXRBwK`PUU7vwC|qU*>ja!YBQHB0^)pcRM0Gt6ngpi8M3_-Q@l% zhJDS0emEcL8CA9iZlT%64a)>|D*WGIz3R-sq=_WM6z0 zEs|r%?BvkCP~PV^S$nxlvo+d$FyPEsSZq*LWkeJb&sbNm*nvTNW3-1#{)Bvl^yh+Svo7cK9RB8Rd zCtUVnf28HL?jtVqv|qiEqbwum*D^SEkkc^Pe+Ps(WUft_cV1*C0oiL_!?2Iwe+HD_ zFx>5k1vQLTmGOYGkmL7R0#p@kl9#RnRYgy0MOAM6TYl|=t%+YAaMestZ>1CE@v=X4hyQVm{{Wu5nhO%q zHYm~a&nxi=BpXM`(*AGx);-*vnVnOQFW!bcamx*Iu|MDrPmWXdclHE4pdX3_o6N0g zVN+a8m>P}!aBD^4h(U(K{U|eakWtZ0OvK3;*10o1UL;`}DANu#o0Yw|9Pi_)P@Y)i zsf$+%-8<@v#m=;}e>=gmH+FHh~x5{Us~T%#h*57A9<2)3K(zDcRQDhV4`>EuEh4 z^$hk8ji$l3$-_UK$2b1hw)M>g zqPVV4@W0E-@XPEvPYtBf=Kr_Hlf#4k(t`AM$oZbPI{F$)ZAkQu_{80RP5sDg{KSB1 z=y{oZ^dIL}b|(`{(5?-~`e%$H!q%Vg7p*vDymHB4M!0ZeOj!Q<$-eHjdG;j6)-wv% zW=toagf(OM7`nCKHZ+(Tz<;>)3y@nQ9_}6(NhXHVBUH6j=KY$HtT}t$vf)%8v<7)a zF}(?YF)Wyg1e))!T-QziAf61GDF=Jfi8Q){=bgME(KDJ(4G#GIgB!878us6w7#>LU z=`yGStGGK24UP^|Nromz5UyqZSLp1YzJk!OBwhJguQ8wM&>?rcXY|5+Zu=H87Wyz!&#qav$r%pQQnO|*sFV4AjdB)Vke}3WZ zUwg$R@7eU_g@3;HfwPbJ#GKffAD{B;58v{-*b`SSzWKt(m&I0m;^+JCSi2~;&~w(eBr)FzjgND3+r!rb?n`5?6~WeH+r$y*psP`J-%zq zum9LPjXd6u7e2M_wp;f8xP9jx&tLb4TQ*)4J7NEq-gL=tZi#&@^p(5n+V6@jJNdjh ze|Y&PW4~B4>#FbF^`Y3Oul~AG`0a*dOlxmu)Y<<3q8lw|?Z-Z+-p(X~#PkHdS2zRO5@U zeeIM_-+IlpYks!yn9tmI(iv?pF8swW>pC90^ydqIv94+VH-32E!Y`h1&g;MTx=JDo}zTWZ3FXo^4$AzDL#s01DKj6i7PX6}i z`+xNFg|A3#c=@@%yKmu>+xK>U?yDOX*57y4tKab7a~FQA@^klhA%DdOe~_U3(bE>+ za?Y>c^SW!k*z-v`Z{Pjg6~B4xN0*xR+q-b~{H3e@G`A{tCLiDUf}9t=Xbz#vwDwr$ z${Sil*GCBux^TrAUT8z@q;dkvr%tZi7NV<}_|(w_j!JXlZ^BQMdlz3N_+fIU6Vdm{ zy^C)a+~spq0sQU)d>$)+KPC9FRJWzmn&_6~`7>SMm}*Y^P54Z%x#wL+$kJ&(od};N zcm)A+5WZCK7YI}Q&n-Zo7W^HC-lr4sxwZiPtp(`sD?tB*;I9177oZQ9kp!OiFZ8AA zdz9dAeOm;-!_a5+xx4^AzQi z{(_R92nfe+IEVqOY5Ni!SXo@?~##y4!19GctnvL33ly zlsD2Cv9P)AkX$|K!C}fl%m&uKzB|1*HLx5?s+>Z_|gs^*-5>mDt*M{rqkD`=dqC~=;gS!a}EW2 z4b`S|jWrl_7EXj%pEg^FUn{sYsi2-W=9|k{1~C`DbN_1G2bF(Yt0{kj&{O$ThFg9_ zFt_}jJ`!-Oc9{E&X205SNhMkPw<$Ki|2d18v!8qG(ZqxDM;(0ccMb5QgA*+s_I#>b zyZP_=9+P>$as;#WxB0szzgz#?erWQ$?QQd$RP8r<6a2G=2E`2G{HG2`Ybulz*GhQ~BgB-1090M&-}IM<21h=5Bid{aHR+K^ByM z?~hG>-<&evR-g9#|Bw=t-}{Nl;?~XH+w<*`-x1TRl>B=f#ag+oK>v;l%gbVP-f#6M z`AL5yv(0bY@6*6Ye)o9tX9jWg>eAj1G8dzqKbQQ5OZMBPosJ}@icDw4;NgDV}C-|uM&Jr z64uF?x&`sU7ataO)wu}0??O)cQ6jUiL!+G z;Ff;7@Ui#y{3)T2IrLOaxmQi6yrbNkmWL5g?$s2)t#rA`<(?=(!C!(rUO>5zF`i%U z<3exS)v_|}TAeq)`hn2zwuEwiO7PXN7k&I(@a+z6+PgZ_Uu>S=3;mEo|7XE(bMP|O zd$qU0!7By7)4`_+KIY&v1;59^Uncm~4sPtD+S}>iuNHcZxg8L4i6@u?|_`F&028aG!!Jl#H*9hL^(61AGpF=+&c!xuu z7JR=$f1%(phki_O&%v)0yvxC_6}-~H-!1qG2ftD9S_l8I;Ac7bt%A>X@Q*P*Ew3Nl zDfBy?`hLdpaq#;DU+3Tt3hq1jZo!8f{40VtIQSESZ*cJM2;Su2dj%hJ@Sh3Z;o#2; zezk+YAb8Bd_Y1zw!T%z7mxG6SN>T0I=-^WXU*X`jf^T>5BLqLo!H*LBHU~dO@YN39 zAo!gQ-YEDz4&Ea8laBoyFZf;uj|tv@@lzkK7yKEAe!1ZL9Q+Kyn;br83BKQn`GUtBd`R$02VXCEmxEs7@JAUhG_P=1L!S3F&>y3m zYNfR1Hw3Shdy9WZa9?n{yx1#vlY>7kcua81=Q+V`KehNi!B;zcelPftgC78HWt1d7 z&B78Q2tNw=4Aj?*haP5}=4dWJ{|2F7ErgbSP;io$4vW86@G+b<{u#l2S*lq2y@Ky_ z@Jig1oMWOGOMjf;J}vxkSo}@EgX-H;0N+>ue~<9r$@4FXxucv13(!AN0RN%zSF;# z4*Yb%w;>-Lmi{{6WdBW~PmXaT{XMw%iO@N26#CUdPw{9vZWeq@aE^Tg|99N`MBp4B zD}diAd{)a+!t%dI@G-$TRt}y!3-H-p0DnyQ>=a|M{J$aiUcouG4*uUQz~?6g@TUvl zzbJtJwgA2#IJL`O3b5d?a=s`y8kvu$&B?34LvCNL^@;YM1mQKniT^&~WBDHeJg6Vh zP6M*5)$||%2gm+#jufDe6~Ip`fS+3cUsnJhEr4HH0KcvPzP$kci30ep0{FKJ;6E*Z z?=OJYVn7Mf|0@dM#}>dl3*c`ofUhZlZzzCYQUKpl0KdKfeoF!TwgUK_1@KQ7z`b<; z(1v6>-h`v2xu=sa=w^N-E*q(Ov27;O*fTWb#pAt+^{Ji&Z9s|F@dl?l2h47OLdx08 z+gsxCk#s7KgLdGYkyx7=!FHELeSPuf&Um4+@!QN)b{u6!;;B^J>;hX)3?r@ac(Oaa zG$RLr(Kf^W%zf(!=WWl`tg5BI45%}S|zs@CGI|BZDfULg<2dwO=gdgY#Eebq`(i=|3 ze?(6uMi@cY=pkd&80&G(=vu*O{~(Mn9#8Zo1&ZWDOjC6Q+S#*?t=&*E&vrw}oZFeQg;QG?Z#CD9w=&Yg z99j(}b7*Cxg*mh`(!x@;8fjUcRz|R)60)@##MF`TX2v6oNBKL#NYs2YWix+AS)OM8 zj{2gU^n)y4*@-*{zl;vsW?N{uKDH?+nK(dDLc$HQ?@fr2jlHb*2<@krP468z&*K6uF-)W+BXR5OgAgM)3vm1uAmx!lPh+s zmiDjvAl0GT*Q!Jx#|vOM)C1j{yN4ZQmD#*0Xm-TPY9Au!`Qo64Fe8b_J6jX!(V<1d zYe#tx1|iG4(Jt#y{g*bt5QB*hEyz2OL4}HZX~PL|cnJdu$vwd#gi4~k`nzrLmO8OW)aLgd!>hKc-1;6 zvv=i2bjjhtfz*XG^&IH#S4=#-HegDJR43C;_?4I`46OA=(!JOi?O_#`yXsUUhoj6>`S%$vaxE>2!@;Q#kNyeEuvOPHmin=qc?^iR-)%*K`m@ zA)4M0-Fe;z&57|XxVC)$RpFZcKLjWK5?W!z#=8l(vQffb`oclwZpFv)x9fOUp67`e z2l2tuKXcG?b~*^ZmaoiLz0||yZ}*D2{3Akd`8No!#R|7ECrUcM0RIaVuJyBC@gezF z2>+`Te!XCHxUDO~eES@9wH_Xw{yEBsxG|0Rl#mj4QcYyQ_LKE!{&@cN*_->vvR zr1)t5k1AaA|El6c_J5@C+N*G)9~1rBy{?q}9(=6*|EU0fY!cvu^i2F;VQxLzFF^<4 z#DBZwuNU0qKTqLW{)pm3{Oj?dW3j@C-X5nZJ~sfg^3y&II;dRD|7^vF_**~rc7+rD z7(R5+J`FmEkLG`g!ZrWP6(8bH`x)uDLE%J?&>bJ!6(7x?_I1$V%Ku}9zaQzgzFbM6 z*Z6EM%<$28V*$KfaB9bE@Ue$^?w^(C9+6|Q&>sduO}|X>zY%Fx&g&Jf^q~EDOwp4F_PNAT!ATFA{uD)@qX&g+dG0KL z?=FDT^D;U}PCYI>spv^gyRYhdf|Hz@{zr;l%lQgI0N-8!|ES<3zaAg&RP?0(y`ulS1t)z*B-#z4W*|r0A($cHduGaH^N4->B$yz3x)DuGgLd_%j7?J6Uw~P(_mf z93;OU_iF{Wc4hXfndgZ_Px5Q}If`D(zewS?AdmIKrz*Tb;cr*?c7w_*i}3r|?r1 z{)EC$Q}|N~|A@kWuW(KO@~I%gLFJx~k1cnB!dED~UEyz1_(FxBq3|~d z3jeOc;|hOT;j0z?8-;f({6&SYQTSoixWPf~+k=m7-y;>?tMF!pYdbkv;d*>|yTboX z@wr~%w<-MJ6t3;5J4zfriB|4$0n<66x$FvUT7)AoO) z!nOUsQsLVEU#oC!|LqFb_P<2o+Wt>hxVHau6t3-mt-`hak0@N*|D_7o_PiG|1pJYyS+={n*Wy+zE<)1io(|^{5uLy zDg4I@KVRW5DE#dT|8Ir&DSS#TZg7x3`|+`Qo1yRlg&(8vL4_Zu@F9gCukZ^LzF6VI z3O_^PBMLuT;c11RukcZY4=a4V!Y@O{&7KPIfqgehoDEtnEZ&&y`75*`W->L8~ zC|vi8uPI#1|D3{aS9}hefg2p82d%fG6|VJmtipA>yj9^^p9zH@i!!VpE>ZZW6n>?` zFH-nch2O34_b6QJ;bDbetmq$8xR(D<3cpLyA9pxzaFBk&`0SVGI4urt>vfWYQ(ii{ z9GrO3afX9e8rVBW;hzC!`Cq7T;z50ej?E5Ex})P?99+7C_W=hVV`}`Cf`e}s{1Xm- zwctA(e74{ZIk+$QZiVZ5{Z!$sXQBU+Em7q-1@oVmuq!y!XA$HDFU zNp%i>qwsHZ@a=*(D_qy>B!yGG?E6e_aB%zH(kch{Mb3*Aeh>Iqdw93PcPRXG3cpw3 zPbvI!3jd?R?^F2Evv7lh^!a&wY`JZM)AG0qAN&5+az%f?qCZRFTF!ojYdOah{sqP7 z28C-mKdNvo=O+cXa_$y=-mmDjoR2D8%ky)E?^NsjJ&2Fh+a`r;dETdREzgGqCwa(suEd9q+ZDZ*=Rt>_=2CR*QuGhu-pcc|Lw_+- zd*1T}=xfCRQoSBld|syTFDg7<0KZz{UsCipC|viWk1Aa2`TGjrrTF}=0A4)@L^w!( zP2aBY&mhg}=PZTqR`_KKe?;LQ5S-eJ{GmN=R`ibowERD%aBUCYSGX?scMktCDd4{q zz2^U_ml*=8ujX@t;2Al^Z^RV6=CexS+CI-!e8^A6s306EMX&91RN*^7XZ3K6!tYgh z)ys)nem|NgILS}t+V|mF6}>L^JcaA_`iR0G1RpEUcNMPbk2wlNIH;PM{$z!J5NVdr zpu#_-@QW4xWrc51_*dpTa-t@ZUE8nIyQ?r1@J@x{9g;;pDBPpS^)n+ z0sPMe@Hf4Zv}W|CEB=(Qi4JDFj|Z^EVJ^MJxXr0`ADnQ;)G594F$=aTW^ zApB;0EdEJZxDl@7$++c%x(~6WF`Euy&&waGiv67)pAvCcda_+Q8q6s}FSnWfK}1g~ zjOjre_k@#7WBAaa%O%|YPK-=9GUeXQ)Sl1KSR#sAHFWvbh7dgPW7@!OfD=g&E%cs;J(}sIk>IQ)eat$`|S?C zLhkQz@YQm^$H8rR&p7y)+_NHaZRNN7QSExv;&z`;6E&2gx9@FtIrwVP^J*8Dezn2D z?fb{u9DJY9-{#<+^uwJFJ|_5+4sPF9-sj-H^v_CZR9j#B9<%S@JB7Z(!R`CaD;#`` zj2DNUe_8(az2`B9-k1Jwe@DR5+xMaGbm&(L{cZ;z6MV0O+xM9FJGgz1xmG%rE!V!6 l+~DB$z2q(zmpJ2U2e