From 2daed2f9eb59d0b9f9637800ab16954abaae0f00 Mon Sep 17 00:00:00 2001 From: Klagarge Date: Fri, 14 Nov 2025 12:34:05 +0100 Subject: [PATCH] feat(rippling): add lab Image Rippling 2D - Full compute - Demi compute (optimisation) --- Student_Cuda_Image/.clangd | 7 + Student_Cuda_Image/.zed/tasks.json | 17 +++ .../INC_SYMLINK/PROJECT/Mandelbrot.h | 2 +- .../INC_SYMLINK/PROJECT/MandelbrotMath.cu.h | 2 +- .../INC_SYMLINK/PROJECT/Mandelbrot_BestGrid.h | 2 +- .../INC_SYMLINK/PROJECT/Raytracing.h | 2 +- .../INC_SYMLINK/PROJECT/RaytracingMath.cu.h | 2 +- .../INC_SYMLINK/PROJECT/Raytracing_BestGrid.h | 2 +- .../INC_SYMLINK/PROJECT/Rippling.h | 2 +- .../INC_SYMLINK/PROJECT/RipplingMath.cu.h | 2 +- .../INC_SYMLINK/PROJECT/Rippling_BestGrid.h | 2 +- .../INC_SYMLINK/PROJECT/real_mandelbrot.h | 2 +- Student_Cuda_Image/out/Rippling_justesse.html | 140 ++++++++++++++++++ .../out/Rippling_performance.html | 140 ++++++++++++++++++ .../bruteforce/Rippling-Cuda-uchar4_db.csv | 10 ++ .../bruteforce/Rippling-Cuda-uchar4_dg.csv | 10 ++ .../bruteforce/Rippling-Cuda-uchar4_fps.csv | 10 ++ .../01_Rippling/device/math/RipplingMath.cu.h | 56 +++---- .../01_Rippling/device/ripplingDevice.cu | 61 ++++++-- .../01_student/01_Rippling/host/Rippling.cu | 23 ++- .../01_student/01_Rippling/host/Rippling.h | 5 +- .../01_Rippling/host/Rippling_BestGrid.h | 21 ++- Student_Cuda_Image/src/main/main.cpp | 26 ++-- Student_Cuda_Image/src/main/mainBenchmark.cpp | 36 ++--- .../src/main/mainBruteforce.cpp | 40 ++--- Student_Cuda_Image/src/main/mainImage.cpp | 6 +- Student_Cuda_Image/src/main/mainTest.cpp | 21 +-- 27 files changed, 493 insertions(+), 156 deletions(-) create mode 100644 Student_Cuda_Image/.clangd create mode 100644 Student_Cuda_Image/.zed/tasks.json create mode 100644 Student_Cuda_Image/out/Rippling_justesse.html create mode 100644 Student_Cuda_Image/out/Rippling_performance.html create mode 100644 Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_db.csv create mode 100644 Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_dg.csv create mode 100644 Student_Cuda_Image/out/bruteforce/Rippling-Cuda-uchar4_fps.csv 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

+ +
+Designed by CppTest +
+
+ +

Summary

+ + + + + + + + + + + + + +
TestsErrorsSuccessTime (s)
10100%1.000000
+
+ +

Test suites

+ + + + + + + + + + + + + + + +
NameTestsErrorsSuccessTime (s)
TestImageCuda10100%1.000000
+
+ +

Suite: TestImageCuda

+ + + + + + + + + + + + + +
NameErrorsSuccessTime (s)
allTests0true1.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

+ +
+Designed by CppTest +
+
+ +

Summary

+ + + + + + + + + + + + + +
TestsErrorsSuccessTime (s)
10100%8.000000
+
+ +

Test suites

+ + + + + + + + + + + + + + + +
NameTestsErrorsSuccessTime (s)
TestPerformance_RunnableGPU_A10100%8.000000
+
+ +

Suite: TestPerformance_RunnableGPU_A

+ + + + + + + + + + + + + +
NameErrorsSuccessTime (s)
performanceOnly0true8.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 *| \*---------------------------------------------------------------------*/ -