From bd178531f80f8bc41c998d1c4588f9e18cc29389 Mon Sep 17 00:00:00 2001 From: gburri Date: Sun, 9 Nov 2014 12:03:27 +0100 Subject: [PATCH] TP fractalTP fractal.. --- .../BUILDER/makefile/private/cuda/cudaGCC.mk | 2 +- .../header/namespace_cpu/CalibreurF_CPU.h | 6 +- .../cudatools/header/both/CalibreurF_GPU.h | 6 +- WCudaMSE/Student_Cuda_Image/cudaLinux.mk | 8 +- .../moo/device/FractalDevice.cu | 37 +++-- .../moo/device/math/FractalMath.h | 141 +++++++++++++----- .../02_Mandelbrot_Julia/moo/host/Fractal.cu | 120 ++++++--------- .../02_Mandelbrot_Julia/moo/host/Fractal.h | 65 +++++--- .../provider/FractalProvider.cpp | 27 +++- .../provider/FractalProvider.h | 7 +- .../src/cpp/core/mainGL.cpp | 12 +- 11 files changed, 260 insertions(+), 171 deletions(-) diff --git a/WCudaMSE/BUILDER/makefile/private/cuda/cudaGCC.mk b/WCudaMSE/BUILDER/makefile/private/cuda/cudaGCC.mk index 05d41b3..c5d3a79 100755 --- a/WCudaMSE/BUILDER/makefile/private/cuda/cudaGCC.mk +++ b/WCudaMSE/BUILDER/makefile/private/cuda/cudaGCC.mk @@ -23,7 +23,7 @@ COMPILATEUR:=g++ include ${ROOT_MAKEFILE_PRIVATE_COMMON}/commonLinux.mk include ${ROOT_MAKEFILE_PUBLIC_CUDA}/cudaGCC.mk -#include cudaGCC.mk +-include cudaLinux.mk include ${ROOT_MAKEFILE_PRIVATE_CUDA}/makeCudaGCC.mk diff --git a/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/CalibreurF_CPU.h b/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/CalibreurF_CPU.h index 6f40d89..3de7d95 100755 --- a/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/CalibreurF_CPU.h +++ b/WCudaMSE/BilatTools_CPP/src/core/tools/header/namespace_cpu/CalibreurF_CPU.h @@ -26,17 +26,17 @@ namespace cpu public: - void calibrer(float& value) + void calibrer(float& value) const { value = value * pente + translation; } - float getPente() + float getPente() const { return this->pente; } - float getTranslation() + float getTranslation() const { return this->translation; } diff --git a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/both/CalibreurF_GPU.h b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/both/CalibreurF_GPU.h index 5d5183c..664758d 100755 --- a/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/both/CalibreurF_GPU.h +++ b/WCudaMSE/BilatTools_Cuda/src/core/cudatools/header/both/CalibreurF_GPU.h @@ -41,19 +41,19 @@ namespace gpu public: __BOTH__ - void calibrer(float& value) + void calibrer(float& value) const { value = value * pente + translation; } __BOTH__ - float getPente() + float getPente() const { return this->pente; } __BOTH__ - float getTranslation() + float getTranslation() const { return this->translation; } diff --git a/WCudaMSE/Student_Cuda_Image/cudaLinux.mk b/WCudaMSE/Student_Cuda_Image/cudaLinux.mk index 9d18821..be1d351 100755 --- a/WCudaMSE/Student_Cuda_Image/cudaLinux.mk +++ b/WCudaMSE/Student_Cuda_Image/cudaLinux.mk @@ -3,10 +3,10 @@ # # Notes # -# (N1) Les options de compilation sont souvent les memes d'un projet à l'autre. -# Ces options communes à tout le workspace se trouve daan la partie public du projet BUILDER. +# (N1) Les options de compilation sont souvent les memes d'un projet � l'autre. +# Ces options communes � tout le workspace se trouve daan la partie public du projet BUILDER. # -# (N2) Vous pouvez redefinir ou surcharger ces options dans ce present .mk pour customizer un projet spécifique. +# (N2) Vous pouvez redefinir ou surcharger ces options dans ce present .mk pour customizer un projet sp�cifique. # # (N3) Pour modifier la configuration de tous vos projets, modifier directement les .mk generiques dans le BUILDER # @@ -30,7 +30,7 @@ CXXLDFLAGS+= # nvcc # ######### -#NVCCFLAGS+= --ptxas-options=-v +NVCCFLAGS+= --ptxas-options=-v NVCCLDFLAGS+= ######### diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu index 8281489..f9a8926 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/FractalDevice.cu @@ -23,7 +23,9 @@ using std::endl; |* Public *| \*-------------------------------------*/ -__global__ void fractal(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float t); +__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n); +__global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i); +__device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath); /*--------------------------------------*\ |* Private *| @@ -41,10 +43,21 @@ __global__ void fractal(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineM |* Private *| \*-------------------------------------*/ -__global__ void fractal(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float t) + +__global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n) + { + FractalMandelbrotMath fractalMath(n); + fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath); + } + +__global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i) { - FractalMath fractalMath(n); + FractalJuliaMath fractalMath(n, c_r, c_i); + fractal(ptrDevPixels, w, h, domaineMath, n, fractalMath); + } +__device__ void fractal(uchar4* ptrDevPixels, int w, int h, const DomaineMath& domaineMath, int n, const FractalMath& fractalMath) + { const int TID = Indice2D::tid(); const int NB_THREAD = Indice2D::nbThread(); const int WH = w * h; @@ -57,19 +70,19 @@ __global__ void fractal(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineM int s = TID; while (s < WH) - { - IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) + { + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) - // (i,j) domaine ecran - // (x,y) domaine math - domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y) + // (i,j) domaine ecran + // (x,y) domaine math + domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y) - fractalMath.colorXY(&color,x, y, t); // update color + fractalMath.colorXY(&color, x, y); - ptrDevPixels[s] = color; + ptrDevPixels[s] = color; - s += NB_THREAD; - } + s += NB_THREAD; + } } /*----------------------------------------------------------------------*\ diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/math/FractalMath.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/math/FractalMath.h index 2bd56ae..aff0b79 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/math/FractalMath.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/device/math/FractalMath.h @@ -6,32 +6,20 @@ #include "CalibreurF.h" #include "ColorTools.h" -/*----------------------------------------------------------------------*\ - |* Declaration *| - \*---------------------------------------------------------------------*/ - -/*--------------------------------------*\ - |* Public *| - \*-------------------------------------*/ - class FractalMath { - - /*--------------------------------------*\ - |* Constructor *| - \*-------------------------------------*/ - public: - + /** + * n: le nombre maximum d'iterations afin de savoir si Zi diverge ou non. + */ __device__ FractalMath(int n) - : n(n), calibreur(IntervalF(-1, 1), IntervalF(0, 1)) + : n(n), calibreur(IntervalF(1, n), IntervalF(0, 1)) { } - /*--------------------------------------*\ - |* Methodes *| - \*-------------------------------------*/ + __device__ + virtual ~FractalMath() {} public: /** @@ -39,40 +27,115 @@ class FractalMath * y=pixelJ */ __device__ - void colorXY(uchar4* ptrColor, float x, float y, float t) + void colorXY(uchar4* ptrColor, float x, float y) const { - float z = f(x, y, t); - - calibreur.calibrer(z); - float hue01 = z; - - ColorTools::HSB_TO_RVB(hue01, ptrColor); // update color - - ptrColor->w = 255; // opaque + // Z courant. + float z_r, z_i; + this->initZ(z_r, z_i, x, y); + + int i = 0; + while (i < this->n) + { + i++; + + nextZ(z_r, z_i, x, y); + + if (isDivergent(z_r, z_i)) + break; + } + + if (i == this->n) + { + ptrColor->x = 0; + ptrColor->y = 0; + ptrColor->z = 0; + } + else + { + float s = static_cast(i); + this->calibreur.calibrer(s); + ColorTools::HSB_TO_RVB(s, ptrColor); + } } private: __device__ - float f(float x, float y,float t) + static bool isDivergent(float a, float b) { - return sin(x * n + t) * cos(y * n + t); + return pow(a, 2) + pow(b, 2) > 4; } + __device__ + virtual void initZ(float& z_r, float& z_i, float x, float y) const = 0; - /*--------------------------------------*\ - |* Attributs *| - \*-------------------------------------*/ + __device__ + virtual void nextZ(float& z_r, float& z_i, float x, float y) const = 0; private: - // Input int n; - - // Tools CalibreurF calibreur; }; -#endif +class FractalMandelbrotMath : public FractalMath + { + public: + /** + * n: le nombre maximum d'iterations afin de savoir si Zi diverge ou non. + */ + __device__ + FractalMandelbrotMath(int n) + : FractalMath(n) + { + } -/*----------------------------------------------------------------------*\ - |* End *| - \*---------------------------------------------------------------------*/ + private: + __device__ + void initZ(float& z_r, float& z_i, float, float) const + { + z_r = 0.0; + z_i = 0.0; + } + + __device__ + void nextZ(float& z_r, float& z_i, float x, float y) const + { + // Z^2 + (x, iy) : + float z_r_tmp = pow(z_r, 2) - pow(z_i, 2); + z_i = 2 * z_r * z_i + y; + z_r = z_r_tmp + x; + } + }; + +class FractalJuliaMath : public FractalMath + { + public: + /** + * n: le nombre maximum d'iterations afin de savoir si Zi diverge ou non. + */ + __device__ + FractalJuliaMath(int n, float c_r, float c_i) + : FractalMath(n), c_r(c_r), c_i(c_i) + { + } + + private: + __device__ + void initZ(float& z_r, float& z_i, float x, float y) const + { + z_r = x; + z_i = y; + } + + __device__ + void nextZ(float& z_r, float& z_i, float, float) const + { + // Z^2 + C : + float z_r_tmp = pow(z_r, 2) - pow(z_i, 2); + z_i = 2 * z_r * z_i + this->c_i; + z_r = z_r_tmp + this->c_r; + } + + float c_r, c_i; + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu index a92b86b..88a225b 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.cu @@ -7,47 +7,15 @@ using std::cout; using std::endl; -/*----------------------------------------------------------------------*\ - |* Declaration *| - \*---------------------------------------------------------------------*/ +extern __global__ void fractalMandelbrot(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n); +extern __global__ void fractalJulia(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float c_r, float c_i); -/*--------------------------------------*\ - |* Imported *| - \*-------------------------------------*/ - -extern __global__ void fractal(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, int n, float t); - -/*--------------------------------------*\ - |* Public *| - \*-------------------------------------*/ - -/*--------------------------------------*\ - |* Private *| - \*-------------------------------------*/ - -/*----------------------------------------------------------------------*\ - |* Implementation *| - \*---------------------------------------------------------------------*/ - -/*--------------------------------------*\ - |* Public *| - \*-------------------------------------*/ - -/*-------------------------*\ - |* Constructeur *| - \*-------------------------*/ - -Fractal::Fractal(int w, int h, float dt, int n) - : w(w), h(h), n(n), +Fractal::Fractal(int w, int h) + : w(w), h(h), dg(8, 8, 1), db(16, 16, 1), - t(0), - variateurAnimation(IntervalF(0, 2 * PI), dt), - ptrDomaineMathInit(new DomaineMath(-2, -1.3, 0.8, 1.3)), title("Fractal Cuda") { - assert(w == h); - //print(dg, db); Device::assertDim(dg, db); } @@ -57,73 +25,79 @@ Fractal::~Fractal() delete this->ptrDomaineMathInit; } -/*-------------------------*\ - |* Methode *| - \*-------------------------*/ - /** * Override */ -void Fractal::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) +int Fractal::getW() { - fractal<<>>(ptrDevPixels, this->w, this->h, domaineMath, this->n, this->t); + return this->w; } /** * Override */ -void Fractal::animationStep() +int Fractal::getH() { - this->t = this->variateurAnimation.varierAndGet(); + return this->h; } -/*--------------*\ - |* get *| - \*--------------*/ - +DomaineMath* Fractal::getDomaineMathInit() + { + return this->ptrDomaineMathInit; + } /** * Override */ -int Fractal::getW() +string Fractal::getTitle() { - return this->w; + return this->title; } -/** - * Override - */ -int Fractal::getH() +///// + +FractalMandelbrot::FractalMandelbrot(int w, int h, float dn) + : Fractal(w, h), variateurAnimationN(IntervalF(30, 100), dn) { - return this->h; + this->ptrDomaineMathInit = new DomaineMath(-2, -1.3, 0.8, 1.3); } -DomaineMath* Fractal::getDomaineMathInit() +void FractalMandelbrot::animationStep() { - return this->ptrDomaineMathInit; + this->n = this->variateurAnimationN.varierAndGet(); } -/** - * Override - */ -float Fractal::getT() +float FractalMandelbrot::getT() { - return this->t; + return this->n; } -/** - * Override - */ -string Fractal::getTitle() +void FractalMandelbrot::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) { - return this->title; + fractalMandelbrot<<>>(ptrDevPixels, this->w, this->h, domaineMath, static_cast(this->n)); + } + +///// + +FractalJulia::FractalJulia(int w, int h, float dn, float z_r_min, float z_r_max, float z_i_min, float z_i_max) + : Fractal(w, h), z_r(z_r_min), z_i(z_i_min), variateurAnimationR(IntervalF(-0.8, -0.7), 0.0005), variateurAnimationI(IntervalF(-0.3, 0.3/*0.15*/), 0.0004) + { + this->ptrDomaineMathInit = new DomaineMath(-1.7, -1.4, 1.7, 1.4); + } + +void FractalJulia::animationStep() + { + this->z_r = this->variateurAnimationR.varierAndGet(); + this->z_i = this->variateurAnimationI.varierAndGet(); } -/*--------------------------------------*\ - |* Private *| - \*-------------------------------------*/ +float FractalJulia::getT() + { + return this->z_r; + } -/*----------------------------------------------------------------------*\ - |* End *| - \*---------------------------------------------------------------------*/ +void FractalJulia::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) + { + fractalJulia<<>>(ptrDevPixels, this->w, this->h, domaineMath, 300, this->z_r, this->z_i); + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h index 4dbfe65..d9b8841 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/moo/host/Fractal.h @@ -9,36 +9,59 @@ class Fractal : public AnimableFonctionel_I { public: - Fractal(int w, int h, float dt, int n); - virtual ~Fractal(void); + Fractal(int w, int h); + virtual ~Fractal(); public: - void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/; - void animationStep() /*override*/; + int getW() /*override*/; + int getH() /*override*/; + DomaineMath* getDomaineMathInit() /*override*/; - int getW() /*override*/; - int getH() /*override*/; - DomaineMath* getDomaineMathInit() /*override*/; + string getTitle(void) /*override*/; - float getT() /*override*/; - string getTitle(void) /*override*/; + protected: + // Inputs + const int w; + const int h; + + // Tools + const dim3 dg; + const dim3 db; + + DomaineMath* ptrDomaineMathInit; + + // Outputs + const string title; + }; + +class FractalMandelbrot : public Fractal + { + public: + FractalMandelbrot(int w, int h, float dn); + void animationStep(); + float getT() /*override*/; private: - // Inputs - const int w; - const int h; - int n; + void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/; - // Tools - const dim3 dg; - const dim3 db; - float t; + VariateurF variateurAnimationN; - VariateurF variateurAnimation; - DomaineMath* ptrDomaineMathInit; + float n; + }; + +class FractalJulia : public Fractal + { + public: + FractalJulia(int w, int h, float dn, float z_r_min, float z_r_max, float z_i_min, float z_i_max); + void animationStep(); + float getT() /*override*/; + + private: + void runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) /*override*/; - // Outputs - const string title; + float z_r, z_i; + VariateurF variateurAnimationI; + VariateurF variateurAnimationR; }; #endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp index 9a824a2..125225d 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.cpp @@ -29,21 +29,34 @@ |* static *| \*----------------*/ -Fractal* FractalProvider::create() +Fractal* FractalProvider::createMandelbrot() { - int dw = 16 * 30; + int dw = 16 * 50; int dh = 16 * 30; - float dt = 2 * PI / 8000; - int n = 2; + return new FractalMandelbrot(dw, dh, 0.2); + } + + +Fractal* FractalProvider::createJulia() + { + int dw = 16 * 50; + int dh = 16 * 30; - return new Fractal(dw, dh, dt, n); + return new FractalJulia(dw, dh, 0.01, -0.745, -0.32, -0.09, 0.1); } -ImageFonctionel* FractalProvider::createGL() +ImageFonctionel* FractalProvider::createMandelbrotGL() + { + ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0); + return new ImageFonctionel(createMandelbrot(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + } + + +ImageFonctionel* FractalProvider::createJuliaGL() { ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0); - return new ImageFonctionel(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + return new ImageFonctionel(createJulia(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel } /*--------------------------------------*\ diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h index b749235..27bd73e 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/02_Mandelbrot_Julia/provider/FractalProvider.h @@ -15,8 +15,11 @@ class FractalProvider { public: - static Fractal* create(void); - static ImageFonctionel* createGL(void); + static Fractal* createMandelbrot(); + static Fractal* createJulia(); + + static ImageFonctionel* createMandelbrotGL(); + static ImageFonctionel* createJuliaGL(); }; #endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp index c853716..97952e7 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp @@ -46,24 +46,24 @@ int mainGL(void) { //Rippling0Image* ptrRippling0 = Rippling0Provider::createGL(); //Image* ptrRippling = RipplingProvider::createGL(); - - ImageFonctionel* ptrFractalMandelbrot = FractalProvider::createGL(); + //ImageFonctionel* ptrFractalMandelbrot = FractalProvider::createMandelbrotGL(); + ImageFonctionel* ptrFractalJulia = FractalProvider::createJuliaGL(); const bool isAnimation = true; const bool isSelection = true; //GLUTImageViewers rippling0Viewer(ptrRippling0, isAnimation, isSelection, 0, 0); //GLUTImageViewers ripplingViewer(ptrRippling, isAnimation, isSelection, 10, 10); - - GLUTImageViewers fractalMandelbrotViewer(ptrFractalMandelbrot, true, true, 20, 20); + //GLUTImageViewers fractalMandelbrotViewer(ptrFractalMandelbrot, true, true, 20, 20); + GLUTImageViewers fractalJuliaViewer(ptrFractalJulia, true, true, 20, 20); GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte // destruction //delete ptrRippling0; //delete ptrRippling; - - delete ptrFractalMandelbrot; + //delete ptrFractalMandelbrot; + delete ptrFractalJulia; return EXIT_SUCCESS; } -- 2.43.0