Klagarge/MSE-GPU-Cuda
1
0
Fork 0
Code Issues Pull Requests Releases Activity

feat(reduction): add reduction PII add

Browse Source
This commit is contained in:
Rémi Heredero
2025-11-17 18:00:54 +01:00
parent dbaa9440b4
commit 135f425fd7
8 changed files with 400 additions and 60 deletions
Download Patch File Download Diff File Expand all files Collapse all files

194
Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_justesse.html Normal file
View File

@@ -0,0 +1,194 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Add_IntII_justesse </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Add_IntII_justesse </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestReductionAddII">TestReductionAddII</a></td>
<td style="width:10%" class="tablecell_success">10</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">1.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionAddII"></a>Suite: TestReductionAddII</h3>
<table summary="Details for suite TestReductionAddII" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">testDB2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">1.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB4</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB8</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB16</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB32</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB64</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testDB128</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testGrid</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testMonoBlock</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialeMax</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">0.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

140
Student_Cuda_Tools_Reduction/out/Reduce_Add_IntII_performance.html Normal file
View File

@@ -0,0 +1,140 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Strict//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-strict.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" xml:lang="en" lang="en">
<head>
<meta http-equiv="content-type" content="text/html; charset=utf-8" />
<meta name="generator" content="CppTest - https://github.com/cpptest/cpptest" />
<title>Reduce_Add_IntII_performance </title>
<style type="text/css" media="screen">
<!--
hr {
width: 100%;
border-width: 0px;
height: 1px;
color: #cccccc;
background-color: #cccccc;
padding: 0px;
}
table {
width:100%;
border-collapse:separate;
border-spacing: 2px;
border:0px;
}
tr {
margin:0px;
padding:0px;
}
td {
margin:0px;
padding:1px;
}
.table_summary {
}
.table_suites {
}
.table_suite {
}
.table_result {
margin: 0px 0px 1em 0px;
}
.tablecell_title {
background-color: #a5cef7;
font-weight: bold;
}
.tablecell_success {
background-color: #efefe7;
}
.tablecell_error {
color: #ff0808;
background-color: #efefe7;
font-weight: bold;
}
p.spaced {
margin: 0px;
padding: 1em 0px 2em 0px;
}
p.unspaced {
margin: 0px;
padding: 0px 0px 2em 0px;
}
-->
</style>
</head>
<body>
<h1><a name="top"></a>Reduce_Add_IntII_performance </h1>
<div style="text-align:right">
Designed by <a href="https://github.com/cpptest/cpptest">CppTest</a>
</div>
<hr />
<h2>Summary</h2>
<table summary="Summary of test results" class="table_summary">
<tr>
<td style="width:30%" class="tablecell_title">Tests</td>
<td style="width:30%" class="tablecell_title">Errors</td>
<td style="width:30%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td style="width:30%" class="tablecell_success">1</td>
<td style="width:30%" class="tablecell_success">0</td>
<td style="width:30%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td>
</tr>
</table>
<hr />
<h2>Test suites</h2>
<table summary="Test Suites" class="table_suites">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Tests</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success"><a href="#TestPerformance_RunnableGPU_A">TestPerformance_RunnableGPU_A</a></td>
<td style="width:10%" class="tablecell_success">1</td>
<td style="width:10%" class="tablecell_success">0</td>
<td style="width:10%" class="tablecell_success">100%</td>
<td style="width:10%" class="tablecell_success">10.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestPerformance_RunnableGPU_A"></a>Suite: TestPerformance_RunnableGPU_A</h3>
<table summary="Details for suite TestPerformance_RunnableGPU_A" class="table_suite">
<tr>
<td class="tablecell_title">Name</td>
<td style="width:10%" class="tablecell_title">Errors</td>
<td style="width:10%" class="tablecell_title">Success</td>
<td style="width:10%" class="tablecell_title">Time (s)</td>
</tr>
<tr>
<td class="tablecell_success">performanceOnly</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">10.000000</td>
</tr>
</table>
<p class="spaced"><a href="#top">Back to top</a>
</p>
<hr />
<p>
<a href="http://validator.w3.org/#validate-by-upload">
Valid XHTML 1.0 Strict
</a>
</p>
</body>
</html>

30
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/device/reductionAddIntII_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h" #include "Thread2D.cu.h"
#include "cudas.h" #include "cudas.h"
#include "ReductionAdd.cu.h" #include "ReductionAdd.cu.h"
@@ -9,7 +9,9 @@
|* Private *| |* Private *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
static __device__ void reductionIntraThread(int* tabSM); static
__device__
void reductionIntraThread(int* tabSM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -18,10 +20,16 @@ static __device__ void reductionIntraThread(int* tabSM);
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__global__ void KAddIntProtocoleII(int* ptrSumGM) __global__
{ void KAddIntProtocoleII(int* ptrSumGM) {
// TODO ReductionAddIntII // TODO ReductionAddIntII
}
extern __shared__ int tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
ReductionAdd::reduce(tabSM, ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Private *| |* Private *|
@@ -30,12 +38,16 @@ __global__ void KAddIntProtocoleII(int* ptrSumGM)
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__device__ void reductionIntraThread(int* tabSM) __device__
{ void reductionIntraThread(int* tabSM) {
// TODO ReductionAddIntII // TODO ReductionAddIntII
} const int TID = Thread2D::tid();
const int localTID = Thread2D::tidLocal();
tabSM[localTID] = TID;
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

27
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.cu
View File

@@ -14,7 +14,9 @@ using std::to_string;
|* Imported *| |* Imported *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
extern __global__ void KAddIntProtocoleII(int* ptrSumGM); extern
__global__
void KAddIntProtocoleII(int* ptrSumGM);
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Implementation *| |* Implementation *|
@@ -27,25 +29,28 @@ extern __global__ void KAddIntProtocoleII(int* ptrSumGM);
ReductionAddIntII::ReductionAddIntII(const Grid& grid , int* ptrSum,bool isVerbose) : 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_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_AddInt_II",isVerbose), // classe parente RunnableGPU(grid, "Reduce_AddInt_II",isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionAddIntII // TODO ReductionAddIntII
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(int);
}
ReductionAddIntII::~ReductionAddIntII() GM::mallocInt0(&ptrSumGM);
{ }
ReductionAddIntII::~ReductionAddIntII() {
// TODO ReductionAddIntII // TODO ReductionAddIntII
} GM::free(ptrSumGM);
}
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionAddIntII::run() void ReductionAddIntII::run() {
{
// TODO ReductionAddIntII // TODO ReductionAddIntII
}
KAddIntProtocoleII<<<dg,db, this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_int(ptrSum, ptrSumGM);
}
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/add/int/PII/host/ReductionAddIntII.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionAddIntII: public RunnableGPU class ReductionAddIntII: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/
@@ -44,7 +43,7 @@ class ReductionAddIntII: public RunnableGPU
int* ptrSumGM; int* ptrSumGM;
size_t sizeSM; size_t sizeSM;
}; };
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

5
Student_Cuda_Tools_Reduction/src/main/main.cpp
View File

@@ -18,8 +18,7 @@ extern int mainTest();
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
int main(int argc , char** argv) int main(int argc , char** argv) {
{
// Limits::show(); // Limits::show();
CudaContext cudaContext; CudaContext cudaContext;
@@ -40,7 +39,7 @@ int main(int argc , char** argv)
} }
return cudaContext.process(); return cudaContext.process();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|

20
Student_Cuda_Tools_Reduction/src/main/mainTest.cpp
View File

@@ -26,15 +26,14 @@ static void generic();
|* Implementation *| |* Implementation *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
int mainTest() int mainTest() {
{
// activer ci-dessous seulement le TP voulu (pas tous) // activer ci-dessous seulement le TP voulu (pas tous)
add(); add();
//generic(); //generic();
return EXIT_SUCCESS; return EXIT_SUCCESS;
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* private *| |* private *|
@@ -43,20 +42,18 @@ int mainTest()
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void add() void add() {
{
VTReductionAddIntI test1; VTReductionAddIntI test1;
VTReductionAddIntII test2; VTReductionAddIntII test2;
test1.run(); // test1.run();
//test2.run(); test2.run();
} }
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void generic() void generic() {
{
VTReductionGenericI test1; VTReductionGenericI test1;
VTReductionGenericII test2; VTReductionGenericII test2;
VTReductionGenericLongII test3; VTReductionGenericLongII test3;
@@ -64,9 +61,8 @@ void generic()
test1.run(); test1.run();
// test2.run(); // test2.run();
// test3.run(); // test3.run();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

39
Student_Cuda_Tools_Reduction/src/main/mainUse.cpp
View File

@@ -31,19 +31,18 @@ static void print(bool isSuccess);
static const int IS_VERBOSE = true; static const int IS_VERBOSE = true;
int mainUse() int mainUse() {
{
// activer ci-dessous seulement le TP voulu (pas tous) // activer ci-dessous seulement le TP voulu (pas tous)
bool isOk = true; bool isOk = true;
reduction_add(isOk); // voir code ci-dessous pour activer la version voulue 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); print(isOk);
return isOk ? EXIT_SUCCESS : EXIT_FAILURE; return isOk ? EXIT_SUCCESS : EXIT_FAILURE;
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* TP *| |* TP *|
@@ -52,26 +51,24 @@ int mainUse()
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void reduction_add(bool& isOk) void reduction_add(bool& isOk) {
{
// InbI // InbI
{ // {
UseReductionAddIntI algo(IS_VERBOSE); // UseReductionAddIntI algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE); // isOk &= algo.isOk(IS_VERBOSE);
} // }
// IntII // IntII
// { {
// UseReductionAddIntII algo(IS_VERBOSE); UseReductionAddIntII algo(IS_VERBOSE);
// isOk &= algo.isOk(IS_VERBOSE); isOk &= algo.isOk(IS_VERBOSE);
// } }
} }
/** /**
* activer ci-dessous la version souhaiter * activer ci-dessous la version souhaiter
*/ */
void reduction_generic(bool& isOk) void reduction_generic(bool& isOk) {
{
// InbI // InbI
{ {
UseReductionIntI algo(IS_VERBOSE); UseReductionIntI algo(IS_VERBOSE);
@@ -89,22 +86,20 @@ void reduction_generic(bool& isOk)
UseReductionLongII algo(IS_VERBOSE); UseReductionLongII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE); isOk &= algo.isOk(IS_VERBOSE);
} }
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* Tools *| |* Tools *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
void print(bool isSuccess) void print(bool isSuccess) {
{
cout << endl << Couts::REVERSE; cout << endl << Couts::REVERSE;
Couts::status(isSuccess, "Success, Congratulations !", "Failed, sorry!"); Couts::status(isSuccess, "Success, Congratulations !", "Failed, sorry!");
cout << endl << Couts::RESET; cout << endl << Couts::RESET;
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/