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
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 0 |
+ 0 |
+ 100% |
+ 1.000000 |
+
+
+
+
+Test suites
+
+
+ | Name |
+ Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | TestReductionAddII |
+ 10 |
+ 0 |
+ 100% |
+ 1.000000 |
+
+
+
+
+Suite: TestReductionAddII
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | testDB2 |
+ 0 |
+ true |
+ 1.000000 |
+
+
+ | testDB4 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testDB8 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testDB16 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testDB32 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testDB64 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testDB128 |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testGrid |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testMonoBlock |
+ 0 |
+ true |
+ 0.000000 |
+
+
+ | testSpecialeMax |
+ 0 |
+ true |
+ 0.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
+
+
+
+
+Summary
+
+
+ | Tests |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | 1 |
+ 0 |
+ 100% |
+ 10.000000 |
+
+
+
+
+Test suites
+
+
+
+Suite: TestPerformance_RunnableGPU_A
+
+
+ | Name |
+ Errors |
+ Success |
+ Time (s) |
+
+
+ | performanceOnly |
+ 0 |
+ true |
+ 10.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 *|
\*---------------------------------------------------------------------*/
-