diff --git a/Student_Cuda_Image/.clangd b/Student_Cuda_Image/.clangd
new file mode 100644
index 0000000..819060f
--- /dev/null
+++ b/Student_Cuda_Image/.clangd
@@ -0,0 +1,7 @@
+CompileFlags:
+ Add:
+ - "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/INC_SYMLINK/EXT"
+ - "-I/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/INC_SYMLINK/PROJECT"
+---
+Diagnostics:
+ Suppress: "*"
diff --git a/Student_Cuda_Image/.zed/tasks.json b/Student_Cuda_Image/.zed/tasks.json
new file mode 100644
index 0000000..e2472a8
--- /dev/null
+++ b/Student_Cuda_Image/.zed/tasks.json
@@ -0,0 +1,17 @@
+// Project tasks configuration. See https://zed.dev/docs/tasks for documentation.
+//
+// Example:
+[
+ {
+ "label": "Run",
+ "command": "cbicc cuda clean jall runGL",
+ "use_new_terminal": false,
+ "allow_concurrent_runs": false,
+ "reveal": "always",
+ "reveal_target": "dock",
+ "hide": "never",
+ "shell": "system",
+ "show_summary": true,
+ "show_command": true
+ }
+]
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot.h
index c0735ff..b5a6a8d 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/MandelbrotMath.cu.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/MandelbrotMath.cu.h
index 9adb54d..321b02f 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/MandelbrotMath.cu.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/MandelbrotMath.cu.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/device/math/MandelbrotMath.cu.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot_BestGrid.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot_BestGrid.h
index d1d6859..ddd76a0 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot_BestGrid.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Mandelbrot_BestGrid.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/host/Mandelbrot_BestGrid.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing.h
index 8cda233..0a1de98 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/host/Raytracing.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/host/Raytracing.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/RaytracingMath.cu.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/RaytracingMath.cu.h
index fffe5a1..3c94b95 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/RaytracingMath.cu.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/RaytracingMath.cu.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/device/math/RaytracingMath.cu.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/device/math/RaytracingMath.cu.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing_BestGrid.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing_BestGrid.h
index 804f003..6a382d0 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing_BestGrid.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Raytracing_BestGrid.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/host/Raytracing_BestGrid.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/03_RayTracing/host/Raytracing_BestGrid.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling.h
index 037c8bd..6178c03 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/RipplingMath.cu.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/RipplingMath.cu.h
index 121bd62..7853398 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/RipplingMath.cu.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/RipplingMath.cu.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling_BestGrid.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling_BestGrid.h
index 6f4e4f1..50b5d96 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling_BestGrid.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/Rippling_BestGrid.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/INC_SYMLINK/PROJECT/real_mandelbrot.h b/Student_Cuda_Image/INC_SYMLINK/PROJECT/real_mandelbrot.h
index 18d75e8..64f135c 120000
--- a/Student_Cuda_Image/INC_SYMLINK/PROJECT/real_mandelbrot.h
+++ b/Student_Cuda_Image/INC_SYMLINK/PROJECT/real_mandelbrot.h
@@ -1 +1 @@
-/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
\ No newline at end of file
+/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Image/src/core/01_student/02_Mandelbrot/real_mandelbrot.h
\ No newline at end of file
diff --git a/Student_Cuda_Image/out/Rippling_justesse.html b/Student_Cuda_Image/out/Rippling_justesse.html
new file mode 100644
index 0000000..fbad5ef
--- /dev/null
+++ b/Student_Cuda_Image/out/Rippling_justesse.html
@@ -0,0 +1,140 @@
+
+
+
+
+
+
+ Rippling_justesse
+
+
+
+
+
+
+Rippling_justesse
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 1 |
+ 0 |
+ 100% |
+ 1.000000 |
+
+
+
+
+Test suites
+
+
+ | Name |
+ Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | TestImageCuda |
+ 1 |
+ 0 |
+ 100% |
+ 1.000000 |
+
+
+
+
+Suite: TestImageCuda
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | allTests |
+ 0 |
+ true |
+ 1.000000 |
+
+
+Back to top
+
+
+
+
+
+
+ Valid XHTML 1.0 Strict
+
+
+
+
diff --git a/Student_Cuda_Image/out/Rippling_performance.html b/Student_Cuda_Image/out/Rippling_performance.html
new file mode 100644
index 0000000..1332575
--- /dev/null
+++ b/Student_Cuda_Image/out/Rippling_performance.html
@@ -0,0 +1,140 @@
+
+
+
+
+
+
+ Rippling_performance
+
+
+
+
+
+
+Rippling_performance
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 1 |
+ 0 |
+ 100% |
+ 8.000000 |
+
+
+
+
+Test suites
+
+
+
+Suite: TestPerformance_RunnableGPU_A
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | performanceOnly |
+ 0 |
+ true |
+ 8.000000 |
+
+
+Back to top
+
+
+
+
+
+
+ Valid XHTML 1.0 Strict
+
+
+
+
diff --git a/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_db.csv b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_db.csv
new file mode 100644
index 0000000..67498b3
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_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/Rippling-Cuda-uchar4_dg.csv b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_dg.csv
new file mode 100644
index 0000000..ccf8644
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_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/Rippling-Cuda-uchar4_fps.csv b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_fps.csv
new file mode 100644
index 0000000..cfa2604
--- /dev/null
+++ b/Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_fps.csv
@@ -0,0 +1,10 @@
+144555,220822,213172,232589,222368,222022,221525,218740,215112,220198,212562,211353,203128,207854,210719,208961
+178200,227914,225423,225667,225283,223689,214874,217351,189550,187027,176077,174967,175908,173784,170924,182629
+212976,225907,222606,224129,213339,213864,204105,204116,170180,174156,180857,200021,186592,189103,180464,178768
+220934,223204,216938,218089,214767,210852,205644,204681,182606,186172,175972,178235,171769,171199,162172,162394
+212311,227349,214553,217104,211690,204241,202456,209851,166727,167536,161323,164796,154234,155004,147491,147138
+198643,222201,217982,217089,212032,203211,198104,196038,154107,155668,149685,148698,142482,142935,134505,133011
+216171,220305,209239,214638,200728,193354,184005,187582,144618,145480,136990,136957,131445,130896,123050,125646
+209897,217355,213187,210719,199503,185380,179740,178260,132461,134694,129448,129059,121525,123736,116539,123906
+206870,215563,213613,212218,194167,175278,170315,167466,125147,125812,119401,120362,114225,119599,114065,119403
+205266,207631,213265,202704,185616,170020,163365,159203,116459,118558,112154,116339,113771,115710,113130,115851
diff --git a/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h b/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
index 47b022d..4355d44 100755
--- a/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
+++ b/Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
@@ -13,26 +13,24 @@
|* Public *|
\*-------------------------------------*/
-class RipplingMath
- {
+class RipplingMath {
/*--------------------------------------*\
|* Constructeur *|
\*-------------------------------------*/
public:
- __device__ RipplingMath(uint w , uint h , float t) :
+ __device__
+ RipplingMath(uint w , uint h , float t) :
dim2(w / 2), //
- t(t)
- {
- // rien
+ t(t) {
+ // rien
}
__device__
- virtual ~RipplingMath()
- {
+ virtual ~RipplingMath() {
// rien
- }
+ }
/*--------------------------------------*\
|* Methode *|
@@ -41,8 +39,7 @@ class RipplingMath
public:
__device__
- void colorIJ(uchar4* ptrColorIJ , int i , int j)
- {
+ void colorIJ(uchar4* ptrColorIJ , int i , int j) {
uchar levelGris = levelGray(i, j); // update levelGris
ptrColorIJ->x = levelGris;
@@ -58,13 +55,13 @@ class RipplingMath
// Etape 2: Puis une fois que l'image grise est valider, attaquer rippling
// 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
+ // }
+ }
private:
@@ -74,24 +71,31 @@ class RipplingMath
__inline__
__device__
- uchar levelGray(int i , int j )
- {
+ uchar levelGray(int i , int j ) {
float result;
- dij(i, j, &result); // warning : dij return void. Ne peut pas etre "imbriquer dans une fonction"
+ dij(i, j, &result); // warning : dij return void. Ne peut pas etre "imbriquer dans une fonction"
+
result = result / 10.f;
+
+ result = 128.f + 127.f * cosf(result - (t / 7.f)) / (result + 1.f);
+ return (uchar)result;
// TODO Rippling GPU : cf formules math rippling.pdf (attribut dim2 = dim/2
- }
+ }
__inline__
__device__
- void dij(int i , int j , float* ptrResult)
- {
+ void dij(int i , int j , float* ptrResult) {
//TODO Rippling GPU cf fonction math pdf
+ float fi = i - dim2;
+ float fj = j - dim2;
+ float d = sqrtf(fi * fi + fj * fj);
+ *ptrResult = d;
+
// Ne pas utiliser la fonction pow pour elever au carrer !
// Utiliser l'opérateur *
- }
+ }
/*--------------------------------------*\
|* Attribut *|
@@ -103,7 +107,7 @@ class RipplingMath
int dim2; // dim2=dim/2
float t;
- };
+};
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/core/01_student/01_Rippling/device/ripplingDevice.cu b/Student_Cuda_Image/src/core/01_student/01_Rippling/device/ripplingDevice.cu
index e5b6194..0ab37d1 100755
--- a/Student_Cuda_Image/src/core/01_student/01_Rippling/device/ripplingDevice.cu
+++ b/Student_Cuda_Image/src/core/01_student/01_Rippling/device/ripplingDevice.cu
@@ -18,12 +18,11 @@ static __device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , flo
|* Implementation *|
\*---------------------------------------------------------------------*/
-__global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t)
- {
- ripplingBaseline(tabPixelsGM, w, h, t);
- // ripplingDemi(tabPixelsGM, w, h, t);
+__global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t) {
+ // ripplingBaseline(tabPixelsGM, w, h, t);
+ ripplingDemi(tabPixelsGM, w, h, t);
// ripplingQuart(tabPixelsGM, w, h, t);
- }
+}
/*----------------------------------------------------------------------*\
|* private *|
@@ -33,23 +32,36 @@ __global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t)
* v1
*/
__inline__
-__device__ void ripplingBaseline(uchar4* tabPixelsGM , uint w , uint h , float t)
- {
+__device__
+void ripplingBaseline(uchar4* tabPixelsGM , uint w , uint h , float t) {
// TODO instacier RipplingMath
+ RipplingMath rgM = RipplingMath(w, h, t);
const int TID = Thread2D::tid();
const int NB_THREAD = Thread2D::nbThread();
const int WH = w * h;
- // TODO Rippling GPU pattern entrelacement
+ // TODO Rippling GPU pattern entrelacement
+
+ int s = TID;
+
+ while(s < WH) {
+ int i = s / w;
+ int j = s % w;
+ uchar4 color;
+ rgM.colorIJ(&color, i, j);
+ tabPixelsGM[s] = color;
+ s += NB_THREAD;
}
+}
+
/**
* v2 : optimisation
*/
__inline__
-__device__ void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t)
- {
+__device__
+void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t) {
// Indication:
// (I1) Utiliser la symetrie horizontale de l'image
// (I2) Calculer que la demi partie superieur
@@ -57,14 +69,36 @@ __device__ void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t)
// Partez de la fin de l'image, peut-etre
// TODO Rippling GPU
+
+ RipplingMath rgM = RipplingMath(w, h, t);
+
+ const int TID = Thread2D::tid();
+ const int NB_THREAD = Thread2D::nbThread();
+ const int WH = w * h;
+ const int WH2 = (WH / 2) + 1;
+ int s = TID;
+
+ while(s < WH2) {
+ int i = s / w;
+ int j = s % w;
+ uchar4 color;
+ rgM.colorIJ(&color, i, j);
+ tabPixelsGM[s] = color;
+
+ // Miroir
+ int si = WH - s - 1;
+ tabPixelsGM[si] = color;
+
+ s += NB_THREAD;
}
+}
/**
* v3 : optimsation : defi (difficile) (pas necessaire pour test performance)
*/
__inline__
-__device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t)
- {
+__device__
+void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t) {
// Indication:
// (I1) Utiliser la symetrie horizontale et verticale de l'image
// (I2) Calculer que le quart en huat a gauche
@@ -76,9 +110,8 @@ __device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t)
// (C1) Utiliser toujours le pattern d'entrelacement
// TODO Rippling GPU
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.cu b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.cu
index d76273d..a1398b9 100755
--- a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.cu
+++ b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.cu
@@ -24,19 +24,17 @@ extern __global__ void rippling(uchar4* tabPixelsGM,uint w, uint h,float t);
\*-------------------------*/
Rippling::Rippling(const Grid& grid , uint w , uint h , float dt , bool isVerbose) :
- Animable_I(grid, w, h, "Rippling-Cuda-uchar4", isVerbose) // super classe
- {
+ Animable_I(grid, w, h, "Rippling-Cuda-uchar4", isVerbose) { // super classe
assert(w == h); // specific rippling
// Animation
this->dt = dt;
this->t = 0; // protected dans Animable
- }
+}
-Rippling::~Rippling()
- {
+Rippling::~Rippling() {
// rien
- }
+}
/*-------------------------*\
|* Methode *|
@@ -48,25 +46,22 @@ Rippling::~Rippling()
*
* Note : domaineMath pas use car image pas zoomable
*/
-void Rippling::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath)
- {
+void Rippling::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath) {
// TODO Rippling
// lancer le kernel avec <<>>
// le kernel est importer ci-dessus (ligne 19)
- assert(false); // to delete once implement
- }
+ rippling<<>>(tabPixelsGM,w,h,t);
+}
/**
* Override
* Call periodicly by the API
*/
-void Rippling::animationStep()
- {
+void Rippling::animationStep() {
t += dt;
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
index a5c41d6..b54e394 100755
--- a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
+++ b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
@@ -10,8 +10,7 @@
|* Declaration *|
\*---------------------------------------------------------------------*/
-class Rippling: public Animable_I
- {
+class Rippling: public Animable_I {
/*--------------------------------------*\
|* Constructor *|
\*-------------------------------------*/
@@ -50,7 +49,7 @@ class Rippling: public Animable_I
// Inputs
float dt;
- };
+};
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
index 6936b3b..41d9d6f 100755
--- a/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
+++ b/Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
@@ -11,30 +11,27 @@
|* Impelmentation *|
\*---------------------------------------------------------------------*/
-namespace rippling
- {
+namespace rippling {
- class BestGrid
- {
+ class BestGrid {
public:
- static Grid get()
- {
+ static Grid get() {
const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP();
// TODO Rippling
+ dim3 dg(MP, 1, 1);
+ dim3 db(CORE_MP, 4, 1);
+ Grid grid(dg, db);
+
+ return grid;
- // to remove once coded
- {
- Couts::redln("aie aie aie, your best grid won t build itself");
- assert(false);
- }
}
};
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/main/main.cpp b/Student_Cuda_Image/src/main/main.cpp
index 7d5a3d7..ae5aef5 100755
--- a/Student_Cuda_Image/src/main/main.cpp
+++ b/Student_Cuda_Image/src/main/main.cpp
@@ -24,34 +24,32 @@ extern int mainTest();
|* Implementation *|
\*---------------------------------------------------------------------*/
-int main(int argc , char** argv)
- {
+int main(int argc , char** argv) {
CudaContextImage cudaContext;
// public
{
- cudaContext.deviceId = 2; // in [0,2] width Server Cuda3
- cudaContext.launchMode = LaunchModeImage::IMAGE; // IMAGE BENCHMARKING BRUTFORCE TESTING
+ cudaContext.deviceId = 2; // in [0,2] width Server Cuda3
+ cudaContext.launchMode = LaunchModeImage::TESTING; // IMAGE BENCHMARKING BRUTFORCE TESTING
- cudaContext.deviceDriver = DeviceDriver::LOAD_CURRENT; // LOAD_CURRENT LOAD_ALL
- cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT
+ cudaContext.deviceDriver = DeviceDriver::LOAD_CURRENT; // LOAD_CURRENT LOAD_ALL
+ cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT
}
// private
{
- cudaContext.args.argc = argc;
- cudaContext.args.argv = argv;
+ cudaContext.args.argc = argc;
+ cudaContext.args.argv = argv;
- cudaContext.mainImage = mainImage;
- cudaContext.mainBenchmark = mainBenchmark;
- cudaContext.mainBrutforce = mainBrutforce;
- cudaContext.mainTest = mainTest;
+ cudaContext.mainImage = mainImage;
+ cudaContext.mainBenchmark = mainBenchmark;
+ cudaContext.mainBrutforce = mainBrutforce;
+ cudaContext.mainTest = mainTest;
}
return cudaContext.process();
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/main/mainBenchmark.cpp b/Student_Cuda_Image/src/main/mainBenchmark.cpp
index 4b836b7..5950412 100755
--- a/Student_Cuda_Image/src/main/mainBenchmark.cpp
+++ b/Student_Cuda_Image/src/main/mainBenchmark.cpp
@@ -32,8 +32,7 @@ static void raytracingCM2SM();
|* Implementation *|
\*---------------------------------------------------------------------*/
-int mainBenchmark()
- {
+int mainBenchmark() {
// Please, un a la fois!
rippling();
// mandelbrot(); // Conseil : use nFixe (by example nMin=nMax=80)
@@ -44,67 +43,60 @@ int mainBenchmark()
// raytracingCM2SM();
return EXIT_SUCCESS;
- }
+}
/*--------------------------------------*\
|* Private *|
\*-------------------------------------*/
-void rippling()
- {
+void rippling() {
const double DURATION_MAX_S = 8;
RipplingProvider provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
-void mandelbrot()
- {
+void mandelbrot() {
const double DURATION_MAX_S = 8;
MandelbrotProvider provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
-void raytracingGM()
- {
+void raytracingGM() {
const double DURATION_MAX_S = 8;
RaytracingProviderGM provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
-void raytracingCM()
- {
+void raytracingCM() {
const double DURATION_MAX_S = 8;
RaytracingProviderCM provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
-void raytracingSM()
- {
+void raytracingSM() {
const double DURATION_MAX_S = 8;
RaytracingProviderSM provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
-void raytracingCM2SM()
- {
+void raytracingCM2SM() {
const double DURATION_MAX_S = 8;
RaytracingProviderCM2SM provider;
BenchmarkImage::run(&provider, DURATION_MAX_S);
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/main/mainBruteforce.cpp b/Student_Cuda_Image/src/main/mainBruteforce.cpp
index 4d8f94e..72cf416 100755
--- a/Student_Cuda_Image/src/main/mainBruteforce.cpp
+++ b/Student_Cuda_Image/src/main/mainBruteforce.cpp
@@ -40,8 +40,7 @@ static void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const
|* Implementation *|
\*---------------------------------------------------------------------*/
-int mainBrutforce()
- {
+int mainBrutforce() {
Matlab matlab;
// Please, un a la fois!
@@ -55,73 +54,66 @@ int mainBrutforce()
matlab.play();
return EXIT_SUCCESS;
- }
+}
/*--------------------------------------*\
|* Private *|
\*-------------------------------------*/
-void rippling(Matlab* ptrMatlab)
- {
+void rippling(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.01; // 0.9 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RipplingProvider provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
-void mandelbrot(Matlab* ptrMatlab)
- {
+void mandelbrot(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.4; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
MandelbrotProvider provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
-void raytracingGM(Matlab* ptrMatlab)
- {
+void raytracingGM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderGM provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
-void raytracingSM(Matlab* ptrMatlab)
- {
+void raytracingSM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderSM provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
-void raytracingCM(Matlab* ptrMatlab)
- {
+void raytracingCM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderCM provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
-void raytracingCM2SM(Matlab* ptrMatlab)
- {
+void raytracingCM2SM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderCM2SM provider;
bruteForce((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
- }
+}
/*--------------------------------------*\
|* Tools *|
\*-------------------------------------*/
template
-void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotType& plotType , double durationMaxS)
- {
+void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotType& plotType , double durationMaxS) {
// Hardware
const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP();
@@ -140,7 +132,7 @@ void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotTy
GridMaillage gridMaillage(iteratorDGx, iteratorDBx);
BruteForce::run(ptrProviderUse, &gridMaillage, ptrMatlab, plotType, durationMaxS);
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
diff --git a/Student_Cuda_Image/src/main/mainImage.cpp b/Student_Cuda_Image/src/main/mainImage.cpp
index 995b293..082811e 100755
--- a/Student_Cuda_Image/src/main/mainImage.cpp
+++ b/Student_Cuda_Image/src/main/mainImage.cpp
@@ -24,8 +24,7 @@ using std::string;
|* Implementation *|
\*---------------------------------------------------------------------*/
-int mainImage(const Args& args)
- {
+int mainImage(const Args& args) {
gpu::GLUTImageViewers::init(args.argc, args.argv); //only once
// ImageOption : (boolean,boolean) : (isSelection ,isAnimation,isOverlay,isShowHelp)
@@ -43,9 +42,8 @@ int mainImage(const Args& args)
gpu::GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte
return EXIT_SUCCESS;
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-
diff --git a/Student_Cuda_Image/src/main/mainTest.cpp b/Student_Cuda_Image/src/main/mainTest.cpp
index 502ba3b..897378d 100755
--- a/Student_Cuda_Image/src/main/mainTest.cpp
+++ b/Student_Cuda_Image/src/main/mainTest.cpp
@@ -31,8 +31,7 @@ static void raytracing();
|* Public *|
\*-------------------------------------*/
-int mainTest()
- {
+int mainTest() {
// activer ci-dessous seulement le TP voulu (pas tous)
rippling();
@@ -40,36 +39,33 @@ int mainTest()
//raytracing(); // voir code ci-dessous pour activer la version voulue
return EXIT_SUCCESS;
- }
+}
/*--------------------------------------*\
|* private *|
\*-------------------------------------*/
-void rippling()
- {
+void rippling() {
VTRippling test1;
test1.run();
- }
+}
/**
* fp16 only
*/
-void mandelbrot()
- {
+void mandelbrot() {
assert(mandelbrotReal::isFp16());
VTMandelbrot test1;
test1.run();
- }
+}
/**
* activer ci-dessous la version souhaiter
*/
-void raytracing()
- {
+void raytracing() {
VTRaytracingGM test1;
VTRaytracingSM test2;
VTRaytracingCM test3;
@@ -79,9 +75,8 @@ void raytracing()
// test2.run();
// test3.run();
// test4.run();
- }
+}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/
-