From 135f425fd793227024a5aa75d1c701d758481805 Mon Sep 17 00:00:00 2001 From: Klagarge Date: Mon, 17 Nov 2025 18:00:54 +0100 Subject: [PATCH] feat(reduction): add reduction PII add --- .../out/Reduce_Add_IntII_justesse.html | 194 ++++++++++++++++++ .../out/Reduce_Add_IntII_performance.html | 140 +++++++++++++ .../PII/device/reductionAddIntII_device.cu | 30 ++- .../add/int/PII/host/ReductionAddIntII.cu | 27 ++- .../add/int/PII/host/ReductionAddIntII.h | 5 +- .../src/main/main.cpp | 5 +- .../src/main/mainTest.cpp | 20 +- .../src/main/mainUse.cpp | 39 ++-- 8 files changed, 400 insertions(+), 60 deletions(-) create mode 100644 Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html create mode 100644 Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html diff --git a/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html new file mode 100644 index 0000000..30e7ffd --- /dev/null +++ b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html @@ -0,0 +1,194 @@ + + + + + + + Reduce_Add_IntII_justesse + + + + + + +

Reduce_Add_IntII_justesse

+ +
+Designed by CppTest +
+
+ +

Summary

+ + + + + + + + + + + + + +
TestsErrorsSuccessTime (s)
00100%1.000000
+
+ +

Test suites

+ + + + + + + + + + + + + + + +
NameTestsErrorsSuccessTime (s)
TestReductionAddII100100%1.000000
+
+ +

Suite: TestReductionAddII

+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
NameErrorsSuccessTime (s)
testDB20true1.000000
testDB40true0.000000
testDB80true0.000000
testDB160true0.000000
testDB320true0.000000
testDB640true0.000000
testDB1280true0.000000
testGrid0true0.000000
testMonoBlock0true0.000000
testSpecialeMax0true0.000000
+

Back to top +

+
+ + +

+ + Valid XHTML 1.0 Strict + +

+ + diff --git a/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html new file mode 100644 index 0000000..d1b6225 --- /dev/null +++ b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html @@ -0,0 +1,140 @@ + + + + + + + Reduce_Add_IntII_performance + + + + + + +

Reduce_Add_IntII_performance

+ +
+Designed by CppTest +
+
+ +

Summary

+ + + + + + + + + + + + + +
TestsErrorsSuccessTime (s)
10100%10.000000
+
+ +

Test suites

+ + + + + + + + + + + + + + + +
NameTestsErrorsSuccessTime (s)
TestPerformance_RunnableGPU_A10100%10.000000
+
+ +

Suite: TestPerformance_RunnableGPU_A

+ + + + + + + + + + + + + +
NameErrorsSuccessTime (s)
performanceOnly0true10.000000
+

Back to top +

+
+ + +

+ + Valid XHTML 1.0 Strict + +

+ + diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu index 3934769..6e77a0a 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu @@ -1,4 +1,4 @@ -#include "Thread1D.cu.h" +#include "Thread2D.cu.h" #include "cudas.h" #include "ReductionAdd.cu.h" @@ -9,7 +9,9 @@ |* Private *| \*---------------------------------------------------------------------*/ -static __device__ void reductionIntraThread(int* tabSM); +static +__device__ +void reductionIntraThread(int* tabSM); /*----------------------------------------------------------------------*\ |* Implementation *| @@ -18,10 +20,16 @@ static __device__ void reductionIntraThread(int* tabSM); /** * TID partout en tabSM */ -__global__ void KAddIntProtocoleII(int* ptrSumGM) - { +__global__ +void KAddIntProtocoleII(int* ptrSumGM) { // TODO ReductionAddIntII - } + + extern __shared__ int tabSM[]; + reductionIntraThread(tabSM); + __syncthreads(); + + ReductionAdd::reduce(tabSM, ptrSumGM); +} /*--------------------------------------*\ |* Private *| @@ -30,12 +38,16 @@ __global__ void KAddIntProtocoleII(int* ptrSumGM) /** * TID partout en tabSM */ -__device__ void reductionIntraThread(int* tabSM) - { +__device__ +void reductionIntraThread(int* tabSM) { // TODO ReductionAddIntII - } + const int TID = Thread2D::tid(); + const int localTID = Thread2D::tidLocal(); + + tabSM[localTID] = TID; + +} /*----------------------------------------------------------------------*\ |* End *| \*---------------------------------------------------------------------*/ - diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu index 439c632..a6b03f9 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu @@ -14,7 +14,9 @@ using std::to_string; |* Imported *| \*---------------------------------------------------------------------*/ -extern __global__ void KAddIntProtocoleII(int* ptrSumGM); +extern +__global__ +void KAddIntProtocoleII(int* ptrSumGM); /*----------------------------------------------------------------------*\ |* Implementation *| @@ -27,25 +29,28 @@ extern __global__ void KAddIntProtocoleII(int* ptrSumGM); ReductionAddIntII::ReductionAddIntII(const Grid& grid , int* ptrSum,bool isVerbose) : //RunnableGPU(grid, "Reduce_AddInt_II_" + to_string(grid.threadCounts()),isVerbose), // classe parente RunnableGPU(grid, "Reduce_AddInt_II",isVerbose), // classe parente - ptrSum(ptrSum) - { + ptrSum(ptrSum) { // TODO ReductionAddIntII - this->sizeSM = -1; - } + this->sizeSM = grid.threadByBlock() * sizeof(int); -ReductionAddIntII::~ReductionAddIntII() - { + GM::mallocInt0(&ptrSumGM); +} + +ReductionAddIntII::~ReductionAddIntII() { // TODO ReductionAddIntII - } + GM::free(ptrSumGM); +} /*--------------------------------------*\ |* Methode *| \*-------------------------------------*/ -void ReductionAddIntII::run() - { +void ReductionAddIntII::run() { // TODO ReductionAddIntII - } + + KAddIntProtocoleII<<sizeSM>>>(ptrSumGM); + GM::memcpyDToH_int(ptrSum, ptrSumGM); +} /*----------------------------------------------------------------------*\ |* End *| diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h index 3ca8014..077246e 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h @@ -8,8 +8,7 @@ |* Declaration *| \*---------------------------------------------------------------------*/ -class ReductionAddIntII: public RunnableGPU - { +class ReductionAddIntII: public RunnableGPU { /*--------------------------------------*\ |* Constructor *| \*-------------------------------------*/ @@ -44,7 +43,7 @@ class ReductionAddIntII: public RunnableGPU int* ptrSumGM; size_t sizeSM; - }; +}; /*----------------------------------------------------------------------*\ |* End *| diff --git a/Student_Cuda_Tools_Reduction/src/main/main.cpp b/Student_Cuda_Tools_Reduction/src/main/main.cpp index 27e18d1..f6b3ddc 100755 --- a/Student_Cuda_Tools_Reduction/src/main/main.cpp +++ b/Student_Cuda_Tools_Reduction/src/main/main.cpp @@ -18,8 +18,7 @@ extern int mainTest(); |* Implementation *| \*---------------------------------------------------------------------*/ -int main(int argc , char** argv) - { +int main(int argc , char** argv) { // Limits::show(); CudaContext cudaContext; @@ -40,7 +39,7 @@ int main(int argc , char** argv) } return cudaContext.process(); - } +} /*----------------------------------------------------------------------*\ |* End *| diff --git a/Student_Cuda_Tools_Reduction/src/main/mainTest.cpp b/Student_Cuda_Tools_Reduction/src/main/mainTest.cpp index d6fb104..6e4ef0b 100755 --- a/Student_Cuda_Tools_Reduction/src/main/mainTest.cpp +++ b/Student_Cuda_Tools_Reduction/src/main/mainTest.cpp @@ -26,15 +26,14 @@ static void generic(); |* Implementation *| \*---------------------------------------------------------------------*/ -int mainTest() - { +int mainTest() { // activer ci-dessous seulement le TP voulu (pas tous) add(); //generic(); return EXIT_SUCCESS; - } +} /*--------------------------------------*\ |* private *| @@ -43,20 +42,18 @@ int mainTest() /** * activer ci-dessous la version souhaiter */ -void add() - { +void add() { VTReductionAddIntI test1; VTReductionAddIntII test2; - test1.run(); - //test2.run(); - } + // test1.run(); + test2.run(); +} /** * activer ci-dessous la version souhaiter */ -void generic() - { +void generic() { VTReductionGenericI test1; VTReductionGenericII test2; VTReductionGenericLongII test3; @@ -64,9 +61,8 @@ void generic() test1.run(); // test2.run(); // test3.run(); - } +} /*----------------------------------------------------------------------*\ |* End *| \*---------------------------------------------------------------------*/ - diff --git a/Student_Cuda_Tools_Reduction/src/main/mainUse.cpp b/Student_Cuda_Tools_Reduction/src/main/mainUse.cpp index db8dc75..d0a6053 100755 --- a/Student_Cuda_Tools_Reduction/src/main/mainUse.cpp +++ b/Student_Cuda_Tools_Reduction/src/main/mainUse.cpp @@ -31,19 +31,18 @@ static void print(bool isSuccess); static const int IS_VERBOSE = true; -int mainUse() - { +int mainUse() { // activer ci-dessous seulement le TP voulu (pas tous) bool isOk = true; reduction_add(isOk); // voir code ci-dessous pour activer la version voulue - //reduction_generic(isOk); // voir code ci-dessous pour activer la version voulue + // reduction_generic(isOk); // voir code ci-dessous pour activer la version voulue print(isOk); return isOk ? EXIT_SUCCESS : EXIT_FAILURE; - } +} /*----------------------------------------------------------------------*\ |* TP *| @@ -52,26 +51,24 @@ int mainUse() /** * activer ci-dessous la version souhaiter */ -void reduction_add(bool& isOk) - { +void reduction_add(bool& isOk) { // InbI - { - UseReductionAddIntI algo(IS_VERBOSE); - isOk &= algo.isOk(IS_VERBOSE); - } + // { + // UseReductionAddIntI algo(IS_VERBOSE); + // isOk &= algo.isOk(IS_VERBOSE); + // } // IntII -// { -// UseReductionAddIntII algo(IS_VERBOSE); -// isOk &= algo.isOk(IS_VERBOSE); -// } - } + { + UseReductionAddIntII algo(IS_VERBOSE); + isOk &= algo.isOk(IS_VERBOSE); + } +} /** * activer ci-dessous la version souhaiter */ -void reduction_generic(bool& isOk) - { +void reduction_generic(bool& isOk) { // InbI { UseReductionIntI algo(IS_VERBOSE); @@ -89,22 +86,20 @@ void reduction_generic(bool& isOk) UseReductionLongII algo(IS_VERBOSE); isOk &= algo.isOk(IS_VERBOSE); } - } +} /*----------------------------------------------------------------------*\ |* Tools *| \*---------------------------------------------------------------------*/ -void print(bool isSuccess) - { +void print(bool isSuccess) { cout << endl << Couts::REVERSE; Couts::status(isSuccess, "Success, Congratulations !", "Failed, sorry!"); cout << endl << Couts::RESET; - } +} /*----------------------------------------------------------------------*\ |* End *| \*---------------------------------------------------------------------*/ -