diff --git a/Student_Cuda_Image/out/Mandelbrot_fp16_justesse.html b/Student_Cuda_Image/out/Mandelbrot_fp16_justesse.html
new file mode 100644
index 0000000..3ab67b8
--- /dev/null
+++ b/Student_Cuda_Image/out/Mandelbrot_fp16_justesse.html
@@ -0,0 +1,140 @@
+
+
+
+
+
+
+ Mandelbrot_fp16_justesse
+
+
+
+
+
+
+Mandelbrot_fp16_justesse
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 1 |
+ 0 |
+ 100% |
+ 2.000000 |
+
+
+
+
+Test suites
+
+
+ | Name |
+ Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | TestImageCuda |
+ 1 |
+ 0 |
+ 100% |
+ 2.000000 |
+
+
+
+
+Suite: TestImageCuda
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | allTests |
+ 0 |
+ true |
+ 2.000000 |
+
+
+Back to top
+
+
+
+
+
+
+ Valid XHTML 1.0 Strict
+
+
+
+
diff --git a/Student_Cuda_Image/out/Mandelbrot_fp16_performance.html b/Student_Cuda_Image/out/Mandelbrot_fp16_performance.html
new file mode 100644
index 0000000..f8de287
--- /dev/null
+++ b/Student_Cuda_Image/out/Mandelbrot_fp16_performance.html
@@ -0,0 +1,140 @@
+
+
+
+
+
+
+ Mandelbrot_fp16_performance
+
+
+
+
+
+
+Mandelbrot_fp16_performance
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 1 |
+ 0 |
+ 100% |
+ 12.000000 |
+
+
+
+
+Test suites
+
+
+
+Suite: TestPerformance_RunnableGPU_A
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | performanceOnly |
+ 0 |
+ true |
+ 12.000000 |
+
+
+Back to top
+
+
+
+
+
+
+ Valid XHTML 1.0 Strict
+
+
+
+
diff --git a/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_db.csv b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_db.csv
new file mode 100644
index 0000000..67498b3
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_db.csv
@@ -0,0 +1,10 @@
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
+64,128,192,256,320,384,448,512,576,640,704,768,832,896,960,1024
diff --git a/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_dg.csv b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_dg.csv
new file mode 100644
index 0000000..ccf8644
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_dg.csv
@@ -0,0 +1,10 @@
+68,68,68,68,68,68,68,68,68,68,68,68,68,68,68,68
+136,136,136,136,136,136,136,136,136,136,136,136,136,136,136,136
+204,204,204,204,204,204,204,204,204,204,204,204,204,204,204,204
+272,272,272,272,272,272,272,272,272,272,272,272,272,272,272,272
+340,340,340,340,340,340,340,340,340,340,340,340,340,340,340,340
+408,408,408,408,408,408,408,408,408,408,408,408,408,408,408,408
+476,476,476,476,476,476,476,476,476,476,476,476,476,476,476,476
+544,544,544,544,544,544,544,544,544,544,544,544,544,544,544,544
+612,612,612,612,612,612,612,612,612,612,612,612,612,612,612,612
+680,680,680,680,680,680,680,680,680,680,680,680,680,680,680,680
diff --git a/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_fps.csv b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_fps.csv
new file mode 100644
index 0000000..eae8ca8
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Mandelbrot-Cuda-uchar4-fp16-nMin20-nMax120_fps.csv
@@ -0,0 +1,10 @@
+11261,21977,27704,28654,28277,28547,28629,28404,28310,28099,28086,28083,27716,27519,27791,28011
+19038,28186,28672,29112,28426,28273,28146,28144,27978,27624,27759,27741,27325,27066,27341,27544
+27172,28378,27808,28260,28006,28038,27917,27990,27496,27172,27343,27313,26925,26740,26918,27046
+27677,28301,27408,28080,27452,27837,27162,27780,27046,26611,26934,26825,26609,26218,26593,26529
+27585,27504,27616,27705,27863,27826,27102,27073,26633,26380,26509,26507,26361,25949,26339,26248
+26661,27570,27783,28103,27632,27613,27589,27606,26069,25936,26088,26036,26011,25575,25894,25830
+27632,27463,27391,27489,27242,27606,27584,26880,25818,25521,25840,25693,25640,25316,25482,25309
+27657,27392,27302,27627,27038,27244,26594,27163,25196,24997,25035,25023,24712,24685,24668,24540
+25842,24673,26280,26057,26492,25605,26130,25341,24431,23305,24874,23801,24877,23471,24577,23293
+23284,25709,25062,26164,25511,27316,25176,24768,23921,22530,23529,22545,23692,22471,24333,22815
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/mandelbrotDevice.cu b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/mandelbrotDevice.cu
index 09fdec1..a8ea389 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/mandelbrotDevice.cu
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/mandelbrotDevice.cu
@@ -20,19 +20,36 @@ static __device__ void color(uchar4* ptrColor,int k, int n);
|* Implementation *|
\*---------------------------------------------------------------------*/
-__global__ void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath domaineMath , int n)
- {
+__global__
+void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath domaineMath , int n) {
// TODO Mandelbrot :
//
// entrelacement
// s -> (i,j) -> (x,y)
// appeler colorXY
- double x;
- double y;
+ double x = 0;
+ double y = 0;
+
+ const int TID = Thread2D::tid();
+ const int NB_THREAD = Thread2D::nbThread();
+ const int WH = w * h;
+
+ MandelbrotMath mm(n);
+
+ int s = TID;
+
+ while ( s < WH) {
+ int i = s / w;
+ int j = s % w;
+ domaineMath.toXY(i, j, &x, &y);
+ mm.colorXY(&tabPixelsGM[s], (real)x, (real)y);
+ s += NB_THREAD;
+ }
+
//domaineMath.toXY(i, j, &x, &y); // x et y doivent etre en double! Caster ensuite en real lors du passage à colorXY
- }
+}
/*----------------------------------------------------------------------*\
|* private *|
@@ -42,8 +59,8 @@ __global__ void mandelbrot(uchar4* tabPixelsGM , uint w , uint h , DomaineMath d
* optimisation lookup table color (facultatif)
*/
__inline__
-__device__ void fill(uchar4* tabSM,int n)
- {
+__device__
+void fill(uchar4* tabSM,int n) {
// Indications:
// (I1) tabSM a n cases
// (I2) La case k contient les couleurs en RVBA lorsque la suite s est arreter a k
@@ -54,18 +71,17 @@ __device__ void fill(uchar4* tabSM,int n)
// Warning
// (W1) Commencer d'abord sans cette piste d'optimisation
- }
+}
/**
* optimisation lookup table color (facultatif)
*/
__inline__
-__device__ void color(uchar4* ptrColor,int k, int n)
- {
+__device__
+void color(uchar4* ptrColor,int k, int n) {
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
index be49faf..aab33e2 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
@@ -33,25 +33,24 @@ using mandelbrotReal::real;
#define ZERO 0.0
#endif
-class MandelbrotMath
- {
+class MandelbrotMath {
/*--------------------------------------*\
|* Constructeur *|
\*-------------------------------------*/
public:
- __device__ MandelbrotMath(int n) : //
- n(n)
- {
+ __device__
+ MandelbrotMath(int n) : //
+ n(n) {
// rien
- }
+ }
__device__
- virtual ~MandelbrotMath()
- {
+ virtual
+ ~MandelbrotMath() {
// rien
- }
+ }
/*--------------------------------------*\
|* Methode *|
@@ -60,8 +59,7 @@ class MandelbrotMath
public:
__device__
- void colorXY(uchar4* ptrColorIJ , real x , real y)
- {
+ void colorXY(uchar4* ptrColorIJ , real x , real y) {
// TODO Mandelbrot
// Calculer la suite en (x,y) et recuperer l'indice d'arret de la suite
@@ -75,12 +73,30 @@ class MandelbrotMath
// Etape 2: Puis une fois que l'image grise est valider, attaquer montecarlo
// debug temp
- // {
- // ptrColorIJ->x = 128;
- // ptrColorIJ->y = 128;
- // ptrColorIJ->z = 128;
- // ptrColorIJ->w = 255; // opacity facultatif
- }
+ // {
+ // ptrColorIJ->x = 128;
+ // ptrColorIJ->y = 128;
+ // ptrColorIJ->z = 128;
+ // ptrColorIJ->w = 255; // opacity facultatif
+ // }
+
+ int itr = suite(x, y);
+
+ if (itr >= n) {
+ ptrColorIJ->x = 0;
+ ptrColorIJ->y = 0;
+ ptrColorIJ->z = 0;
+ ptrColorIJ->w = 255; // opacity facultatif
+ } else {
+ Colors::HSB_TO_RVB(
+ ( (float)itr ) / ( (float)(this->n-1) ),
+ 1.f,
+ 1.f,
+ ptrColorIJ
+ );
+ }
+
+ }
private:
@@ -90,20 +106,35 @@ class MandelbrotMath
*/
__inline__
__device__
- int suite(real x , real y)
- { // TODO Mandelbrot
+ int suite(real x , real y) { // TODO Mandelbrot
- // Utiliser dans vos formules les variable :
- //
- // DEUX
- // QUATRE
- // ZERO
- //
- // definit au debut de ce fichier. Est-utile pour passer facilement d'une version fp64 (double) fp32(float) fp16(half)
+ // Utiliser dans vos formules les variable :
+ //
+ // DEUX
+ // QUATRE
+ // ZERO
+ //
+ // definit au debut de ce fichier. Est-utile pour passer facilement d'une version fp64 (double) fp32(float) fp16(half)
- // Calculer la suite en (x,y) jusqu'à n, à moins que critere arret soit atteint avant
- // return le nombre d'element de la suite calculer, ie un entier
- }
+ // Calculer la suite en (x,y) jusqu'à n, à moins que critere arret soit atteint avant
+ // return le nombre d'element de la suite calculer, ie un entier
+
+ real a = ZERO;
+ real b = ZERO;
+ real tmp = ZERO;
+
+ int itr = -1;
+
+ while ( (a*a + b*b <= QUATRE) && (itr < n) ) {
+ tmp = a;
+ a = x + (a*a - b*b);
+ b = DEUX * tmp * b + y;
+ itr++;
+ }
+
+ return itr;
+
+ }
/////////////////////
// Warning:
@@ -150,7 +181,7 @@ class MandelbrotMath
// Inputs
int n;
- };
+};
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.cu b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.cu
index 1bf0ba6..32ce28b 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.cu
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.cu
@@ -40,10 +40,9 @@ Mandelbrot::Mandelbrot(const Grid& grid , uint w , uint h , const DomaineMath& d
// Tools
this->t = nMin;
- }
+}
-Mandelbrot::~Mandelbrot()
- {
+Mandelbrot::~Mandelbrot() {
// rien
}
@@ -55,28 +54,24 @@ Mandelbrot::~Mandelbrot()
* Override
* Call periodicly by the API
*/
-void Mandelbrot::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath)
- {
- assert(false); // to be removed once implemented
+void Mandelbrot::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath) {
// TODO Mandelbrot
// lauch kernel (you find at line 18)
- }
+ mandelbrot<<>>(tabPixelsGM,w,h,domaineMath,t);
+}
/**
* Override
* Call periodicly by the API
*/
-void Mandelbrot::animationStep()
- {
+void Mandelbrot::animationStep() {
this->t = variateurT.varierAndGet();
- }
+}
-string titre(int nMin , int nMax)
- {
+string titre(int nMin , int nMax) {
return "Mandelbrot-Cuda-uchar4-" + mandelbrotReal::realToString() + "-nMin" + std::to_string(nMin) + "-nMax" + std::to_string(nMax);
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
index 0e96601..bf08b1b 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
@@ -10,8 +10,7 @@
|* Declaration *|
\*---------------------------------------------------------------------*/
-class Mandelbrot: public Animable_I
- {
+class Mandelbrot: public Animable_I {
/*--------------------------------------*\
|* Constructor *|
\*-------------------------------------*/
@@ -53,7 +52,7 @@ class Mandelbrot: public Animable_I
// Tools
Variateur variateurT;
- };
+};
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
index 4fc20c6..f1e89a6 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
@@ -12,49 +12,56 @@
|* Impelmentation *|
\*---------------------------------------------------------------------*/
-namespace mandelbrot
- {
+namespace mandelbrot {
- class BestGrid
- {
+ class BestGrid {
public:
- static Grid get()
- {
- const int MP = Hardware::getMPCount();
- const int CORE_MP = Hardware::getCoreCountMP();
+ static Grid get() {
+ // const int MP = Hardware::getMPCount();
+ // const int CORE_MP = Hardware::getCoreCountMP();
- // fp64 (float 64 bits)
- #ifdef MANDELBROT_DOUBLE
+ // fp64 (float 64 bits)
+ #ifdef MANDELBROT_DOUBLE
- //TODO Mandelbrot
- // to remove once coded
- {
- Couts::redln("aie aie aie, your best grid won t build itself");
- assert(false);
- }
- #endif
+ const int MP = Hardware::getMPCount();
+ const int CORE_MP = Hardware::getCoreCountMP();
- // fp32 (float 32 bits)
- #ifdef MANDELBROT_FLOAT
- //TODO Mandelbrot
- // to remove once coded
- {
- Couts::redln("aie aie aie, your best grid won t build itself");
- assert(false);
- }
- #endif
+ dim3 dg(MP, 1, 1);
+ dim3 db(CORE_MP, 1, 1);
+ Grid grid(dg, db);
- // fp16 (float 32 bits)
- #ifdef MANDELBROT_HALF
- //TODO Mandelbrot
- // to remove once coded
- {
- Couts::redln("aie aie aie, your best grid won t build itself");
- assert(false);
- }
- #endif
+ return grid;
+ #endif
+
+ // fp32 (float 32 bits)
+ #ifdef MANDELBROT_FLOAT
+ const int MP = Hardware::getMPCount();
+ const int CORE_MP = Hardware::getCoreCountMP();
+
+ dim3 dg(MP, 1, 1);
+ dim3 db(CORE_MP, 1, 1);
+ Grid grid(dg, db);
+
+ return grid;
+ #endif
+
+ // fp16 (float 32 bits)
+ #ifdef MANDELBROT_HALF
+ const int MP = Hardware::getMPCount();
+ const int CORE_MP = Hardware::getCoreCountMP();
+
+ dim3 dg(MP, 2, 1);
+ dim3 db(CORE_MP, 4, 1);
+ Grid grid(dg, db);
+
+ return grid;
+ #endif
+ // default
+ Couts::redln("Oups, something went wrong on BestGrid settings");
+ assert(false); // unknown type
+ return Grid();
}
};
diff --git a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
index fc9a2bd..4f4b00b 100755
--- a/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
+++ b/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
@@ -14,107 +14,101 @@
\*---------------------*/
// Choose one of the two (either/either):
-//#define MANDELBROT_DOUBLE
-//#define MANDELBROT_FLOAT
-#define MANDELBROT_HALF
+// #define MANDELBROT_DOUBLE // fp64
+// #define MANDELBROT_FLOAT // fp32
+#define MANDELBROT_HALF // fp16
/*----------------------*\
|* private *|
\*---------------------*/
-namespace mandelbrotReal
- {
+namespace mandelbrotReal {
-// fp64 (float 64 bits)
-#ifdef MANDELBROT_DOUBLE
-//#define real double
-using real=double;
-#endif
+ // fp64 (float 64 bits)
+ #ifdef MANDELBROT_DOUBLE
+ //#define real double
+ using real=double;
+ #endif
-// fp32 (float 32 bits)
-#ifdef MANDELBROT_FLOAT
-//#define real float
-using real=float;
-#endif
+ // fp32 (float 32 bits)
+ #ifdef MANDELBROT_FLOAT
+ //#define real float
+ using real=float;
+ #endif
-// fp16 (float 16 bits)
-#ifdef MANDELBROT_HALF
-//#define real half
-using real=half;
-#endif
+ // fp16 (float 16 bits)
+ #ifdef MANDELBROT_HALF
+ //#define real half
+ using real=half;
+ #endif
- static std::string realToString()
- {
+ static std::string realToString() {
-#ifdef MANDELBROT_DOUBLE
- return "fp64";
-#endif
+ #ifdef MANDELBROT_DOUBLE
+ return "fp64";
+ #endif
-#ifdef MANDELBROT_FLOAT
- return "fp32";
-#endif
+ #ifdef MANDELBROT_FLOAT
+ return "fp32";
+ #endif
-#ifdef MANDELBROT_HALF
- return "fp16";
-#endif
- }
-
- static bool isFp16()
- {
-
-#ifdef MANDELBROT_DOUBLE
- return false;
-#endif
-
-#ifdef MANDELBROT_FLOAT
- return false;
-#endif
-
-#ifdef MANDELBROT_HALF
- return true;
-#endif
-
- assert(false);
- }
-
- static bool isFp32()
- {
-
-#ifdef MANDELBROT_DOUBLE
- return false;
-#endif
-
-#ifdef MANDELBROT_FLOAT
- return true;
-#endif
-
-#ifdef MANDELBROT_HALF
- return false;
-#endif
-
- assert(false);
- }
-
- static bool isFp64()
- {
-
-#ifdef MANDELBROT_DOUBLE
- return true;
-#endif
-
-#ifdef MANDELBROT_FLOAT
- return false;
-#endif
-
-#ifdef MANDELBROT_HALF
- return false;
-#endif
-
- assert(false);
- }
+ #ifdef MANDELBROT_HALF
+ return "fp16";
+ #endif
}
+ static bool isFp16() {
+
+ #ifdef MANDELBROT_DOUBLE
+ return false;
+ #endif
+
+ #ifdef MANDELBROT_FLOAT
+ return false;
+ #endif
+
+ #ifdef MANDELBROT_HALF
+ return true;
+ #endif
+
+ assert(false);
+ }
+
+ static bool isFp32() {
+
+ #ifdef MANDELBROT_DOUBLE
+ return false;
+ #endif
+
+ #ifdef MANDELBROT_FLOAT
+ return true;
+ #endif
+
+ #ifdef MANDELBROT_HALF
+ return false;
+ #endif
+
+ assert(false);
+ }
+
+ static bool isFp64() {
+
+ #ifdef MANDELBROT_DOUBLE
+ return true;
+ #endif
+
+ #ifdef MANDELBROT_FLOAT
+ return false;
+ #endif
+
+ #ifdef MANDELBROT_HALF
+ return false;
+ #endif
+
+ assert(false);
+ }
+}
+
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-