diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/Reduction.cu.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/Reduction.cu.h index 0f06ea7..57b0a1a 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/Reduction.cu.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/Reduction.cu.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAdd.cu.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAdd.cu.h index 937c17b..c8fc940 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAdd.cu.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAdd.cu.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntI.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntI.h index 7cae763..47d8783 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntI.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntI.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntII.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntII.h index fcbddde..2e2d818 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntII.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionAddIntII.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntI.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntI.h index e67c39a..0e809a6 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntI.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntI.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntII.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntII.h index 6159104..1de9a56 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntII.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionIntII.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PII/host/ReductionIntII.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionLongII.h b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionLongII.h index 1d0fae5..e11fe48 120000 --- a/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionLongII.h +++ b/Student_Cuda_Tools_Reduction/INC_SYMLINK/PROJECT/ReductionLongII.h @@ -1 +1 @@ -/home/bilat/CBI/Dropbox/02_CBI_LINUX/CoursCuda/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h \ No newline at end of file +/home/mse15/CUDA/toStudent/code/WCudaStudent/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h \ No newline at end of file diff --git a/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntI_justesse.html b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntI_justesse.html index d155666..a345cca 100755 --- a/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntI_justesse.html +++ b/Student_Cuda_Tools_Reduction/out/Reduce_Add_IntI_justesse.html @@ -84,10 +84,10 @@ Designed by CppTest Time (s) - 0 - 13 - 100% - 0.000000 + 0 + 0 + 100% + 0.000000
@@ -102,11 +102,11 @@ Designed by CppTest Time (s) - TestReductionAddI - 13 - 13 - 0% - 0.000000 + TestReductionAddI + 13 + 0 + 100% + 0.000000
@@ -120,1194 +120,88 @@ Designed by CppTest Time (s) - testDB2 - 1 - false - 0.000000 + testDB2 + 0 + true + 0.000000 - testDB4 - 1 - false - 0.000000 + testDB4 + 0 + true + 0.000000 - testDB8 - 1 - false - 0.000000 + testDB8 + 0 + true + 0.000000 - testDB16 - 1 - false - 0.000000 + testDB16 + 0 + true + 0.000000 - testDB32 - 1 - false - 0.000000 + testDB32 + 0 + true + 0.000000 - testDB64 - 1 - false - 0.000000 + testDB64 + 0 + true + 0.000000 - testDB128 - 1 - false - 0.000000 + testDB128 + 0 + true + 0.000000 - testDB256 - 1 - false - 0.000000 + testDB256 + 0 + true + 0.000000 - testDB512 - 1 - false - 0.000000 + testDB512 + 0 + true + 0.000000 - testDB1024 - 1 - false - 0.000000 + testDB1024 + 0 + true + 0.000000 - testGrid - 60 - false - 0.000000 + testGrid + 0 + true + 0.000000 - testMonoBlock - 5 - false - 0.000000 + testMonoBlock + 0 + true + 0.000000 - testSpecialeMax - 1 - false - 0.000000 + testSpecialeMax + 0 + true + 0.000000

Back to top

-

Test results

-

TestReductionAddI::testDB2

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB2
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB4

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB4
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB8

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB8
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB16

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB16
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB32

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB32
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB64

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB64
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB128

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB128
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB256

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB256
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB512

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB512
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testDB1024

- - - - - - - - - - - - - -
TestTestReductionAddI::testDB1024
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testGrid

- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testGrid
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testMonoBlock

- - - - - - - - - - - - - -
TestTestReductionAddI::testMonoBlock
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testMonoBlock
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testMonoBlock
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testMonoBlock
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
- - - - - - - - - - - - - -
TestTestReductionAddI::testMonoBlock
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-

TestReductionAddI::testSpecialeMax

- - - - - - - - - - - - - -
TestTestReductionAddI::testSpecialeMax
Filesrc/core/cudatools/09_tests/TestCuda_A.cpp:144
MessageisOk
-

Back to TestReductionAddI -

-
-

diff --git a/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h b/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h index f5c9adc..8bd6a88 100755 --- a/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h +++ b/Student_Cuda_Tools_Reduction/src/core/01_algorithme/add/ReductionAdd.cu.h @@ -1,13 +1,12 @@ #pragma once -#include "Thread1D.cu.h" +#include "Thread2D.cu.h" /*----------------------------------------------------------------------*\ |* Implementation *| \*---------------------------------------------------------------------*/ -class ReductionAdd - { +class ReductionAdd { public: /** @@ -46,8 +45,9 @@ class ReductionAdd * */ template - static __device__ void reduce(T* tabSM, T* ptrResultGM) - { + static + __device__ + void reduce(T* tabSM, T* ptrResultGM) { // Rappel : // |ThreadByBlock|=|tabSM| . // Il y autant de case en SM que de thread par block. @@ -55,11 +55,11 @@ class ReductionAdd // 1 thread <---> 1 armoire // TODO ReductionAdd - // reductionIntraBlock - // reductionInterblock + reductionIntraBlock(tabSM); + reductionInterBlock(tabSM, ptrResultGM); // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? pas a le fin en tous les cas - } + } private: @@ -71,8 +71,9 @@ class ReductionAdd * used by reductionIntraBlock */ template - static __device__ void ecrasement(T* tabSM, int middle) - { + static + __device__ + void ecrasement(T* tabSM, int middle) { // Indications : // (I1) je suis un thread, je dois faire quoi ? // (I2) Tous les threads doivent-ils faire quelquechose? @@ -80,39 +81,59 @@ class ReductionAdd // TODO ReductionAdd + const int localTID = Thread2D::tidLocal(); + + if(localTID < middle) { + tabSM[localTID] = tabSM[localTID] + tabSM[localTID + middle]; + } + // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? - } + } /** * Sur place, le resultat est dans tabSM[0] */ template - static __device__ void reductionIntraBlock(T* tabSM) - { + static + __device__ + void reductionIntraBlock(T* tabSM) { // Reduce tab SM (all in [0]) // Ecrasement sucessifs dans une boucle (utiliser la methode ecrasement ci-dessus) // TODO ReductionAdd + const int NB_THREAD_LOCAL = Thread2D::nbThreadLocal(); + int middle = NB_THREAD_LOCAL>>1; + + while (middle > 0) { + ecrasement(tabSM, middle); + __syncthreads(); + middle = middle >> 1; + } + // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? - } + } /*--------------------------------------*\ |* reductionInterblock *| \*-------------------------------------*/ template - static __device__ void reductionInterBlock(T* tabSM, T* ptrResultGM) - { + static + __device__ + void reductionInterBlock(T* tabSM, T* ptrResultGM) { // SM -> GM // Indication: // (I1) Utiliser atomicAdd(pointeurDestination, valeurSource); // (i2) Travailler sous l hypothese d'une grid2d,avec Thread2D // TODO ReductionAdd + if(Thread2D::tidLocal() == 0) { + atomicAdd(ptrResultGM, tabSM[0]); + } // __syncthreads();// pour touts les threads d'un meme block, necessaires? ou? - } + } - }; +}; /*----------------------------------------------------------------------*\ |* End *| diff --git a/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h b/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h index 4a2f32e..23eed8a 100755 --- a/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h +++ b/Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h @@ -1,7 +1,7 @@ #pragma once #include "Lock.cu.h" -#include "Thread1D.cu.h" +#include "Thread2D.cu.h" /*----------------------------------------------------------------------*\ |* prt fonction / reduction *| @@ -14,8 +14,7 @@ |* Implementation *| \*---------------------------------------------------------------------*/ -class Reduction - { +class Reduction { public: /** @@ -50,14 +49,16 @@ class Reduction * ReductionGeneric::reduce(add,addAtomic,tabSm,ptrResultGM); */ template - static __device__ void reduce(BinaryOperator(OP) ,AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) + static + __device__ + void reduce(BinaryOperator(OP) ,AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) //static __device__ void reduce(T (*OP)(T, T) ,void (*ATOMIC_OP)(T*, T), T* tabSM, T* ptrResultGM) // idem ci-dessus mais sans define - { + { // Meme principe que ReductionAdd // TODO ReductionGeneric // Meme principe que ReductionAdd - } + } private: @@ -69,37 +70,40 @@ class Reduction * used by reductionIntraBlock */ template - static __device__ void ecrasement(BinaryOperator(OP),T* tabSM, int middle) - { + static + __device__ + void ecrasement(BinaryOperator(OP),T* tabSM, int middle) { // TODO ReductionGeneric // Meme principe que ReductionAdd // OP est la variable representant l'operateur binaire - } + } /** * Sur place, le resultat est dans tabSM[0] */ template - static __device__ void reductionIntraBlock(BinaryOperator(OP),T* tabSM) - { + static + __device__ + void reductionIntraBlock(BinaryOperator(OP),T* tabSM) { // TODO ReductionGeneric // Meme principe que ReductionAdd // OP est la variable representant l'operateur binaire - } + } /*--------------------------------------*\ |* reductionInterblock *| \*-------------------------------------*/ template - static __device__ void reductionInterBlock(AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) - { + static + __device__ + void reductionInterBlock(AtomicOp(ATOMIC_OP), T* tabSM, T* ptrResultGM) { // TODO ReductionGeneric // Meme principe que ReductionAdd // ATOMIC_OP est la variable representant l'operateur binaire atomic - } + } - }; +}; /*----------------------------------------------------------------------*\ |* End *| diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/device/reductionAddIntI_device.cu b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/device/reductionAddIntI_device.cu index dfa6425..7300970 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/device/reductionAddIntI_device.cu +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/device/reductionAddIntI_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,13 +20,19 @@ static __device__ void reductionIntraThread(int* tabSM); /** * 1 partout en tabSM */ -__global__ void KAddIntProtocoleI(int* ptrSumGM) - { +__global__ +void KAddIntProtocoleI(int* ptrSumGM) { // TODO ReductionAddIntI // Reception tabSM + extern __shared__ int tabSM[]; + // ReductionIntraThread + reductionIntraThread(tabSM); + __syncthreads(); + // ReductionAdd + ReductionAdd::reduce(tabSM, ptrSumGM); // __syncthreads(); // des threads de meme block!// Question : utile? ou? } @@ -36,12 +44,15 @@ __global__ void KAddIntProtocoleI(int* ptrSumGM) /** * 1 partout en tabSM */ -__device__ void reductionIntraThread(int* tabSM) - { +__device__ +void reductionIntraThread(int* tabSM) { // TODO ReductionAddIntI - } + const int localTID = Thread2D::tidLocal(); + + tabSM[localTID] = 1; + +} /*----------------------------------------------------------------------*\ |* End *| \*---------------------------------------------------------------------*/ - diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.cu b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.cu index b37c56f..5456b1e 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.cu +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.cu @@ -27,32 +27,33 @@ extern __global__ void KAddIntProtocoleI(int* ptrSumGM); ReductionAddIntI::ReductionAddIntI(const Grid& grid , int* ptrSum , bool isVerbose) : //RunnableGPU(grid, "Reduce_Add_IntI_" + to_string(grid.threadCounts()),isVerbose), // classe parente RunnableGPU(grid, "Reduce_Add_IntI", isVerbose), // classe parente - ptrSum(ptrSum) - { + ptrSum(ptrSum) { // TODO ReductionAddIntI // MM pour ptrSumGM (oubliez pas initialisation) - this->sizeSM = -1; + this->sizeSM = grid.threadByBlock() * sizeof(int); // Tip: Il y a une methode dedier pour malloquer un int cote device et l'initialiser a zero // - // GM::mallocInt0(&ptrSumGM); - } + GM::mallocInt0(&ptrSumGM); +} -ReductionAddIntI::~ReductionAddIntI() - { +ReductionAddIntI::~ReductionAddIntI() { // TODO ReductionAddIntI - } + GM::free(ptrSumGM); +} /*--------------------------------------*\ |* Methode *| \*-------------------------------------*/ -void ReductionAddIntI::run() - { +void ReductionAddIntI::run() { // TODO ReductionAddIntI // appeler le kernel // recuperer le resulat coter host + KAddIntProtocoleI<<sizeSM>>>(ptrSumGM); + GM::memcpyDToH_int(ptrSum, ptrSumGM); + // Tip: Il y a une methode dedier ramener coter host un int // // GM::memcpyDtoH_int(ptrDestination, ptrSourceGM);); diff --git a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h index 954f75a..d84dd86 100755 --- a/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h +++ b/Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PI/host/ReductionAddIntI.h @@ -8,8 +8,7 @@ |* Declaration *| \*---------------------------------------------------------------------*/ -class ReductionAddIntI: public RunnableGPU - { +class ReductionAddIntI: public RunnableGPU { /*--------------------------------------*\ |* Constructor *| \*-------------------------------------*/ diff --git a/Student_Cuda_Tools_Reduction/src/main/main.cpp b/Student_Cuda_Tools_Reduction/src/main/main.cpp index 615354c..27e18d1 100755 --- a/Student_Cuda_Tools_Reduction/src/main/main.cpp +++ b/Student_Cuda_Tools_Reduction/src/main/main.cpp @@ -26,8 +26,8 @@ int main(int argc , char** argv) // public { - cudaContext.deviceId = 0; // in [0,2] width Server Cuda3 - cudaContext.launchMode = LaunchModeMOO::USE; // USE TEST (only) + cudaContext.deviceId = 1; // in [0,2] width Server Cuda3 + cudaContext.launchMode = LaunchModeMOO::TEST; // USE TEST (only) cudaContext.deviceDriver = DeviceDriver::LOAD_ALL; // LOAD_CURRENT LOAD_ALL cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT @@ -45,4 +45,3 @@ int main(int argc , char** argv) /*----------------------------------------------------------------------*\ |* End *| \*---------------------------------------------------------------------*/ -