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

feat(reduction): add reduction generic PII long

Browse Source
This commit is contained in:
Rémi Heredero
2025-11-17 19:13:08 +01:00
parent 1ec4ef198e
commit cb6856ccde
7 changed files with 415 additions and 37 deletions
Download Patch File Download Diff File Expand all files Collapse all files

224
Student_Cuda_Tools_Reduction/out/Reduce_Generic_LongII_justesse.html Normal file
View File

@@ -0,0 +1,224 @@
<!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_Generic_LongII_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_Generic_LongII_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">80.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="#TestReductionGenericLongII">TestReductionGenericLongII</a></td>
<td style="width:10%" class="tablecell_success">15</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">80.000000</td>
</tr>
</table>
<hr />
<h3><a name="TestReductionGenericLongII"></a>Suite: TestReductionGenericLongII</h3>
<table summary="Details for suite TestReductionGenericLongII" 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">0.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">testDB256</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">testDB512</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">testDB1024</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">testspecialGridDGXMAX</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">39.000000</td>
</tr>
<tr>
<td class="tablecell_success">testSpecialGrid2</td>
<td class="tablecell_success">0</td>
<td class="tablecell_success">true</td>
<td class="tablecell_success">40.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">1.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_Generic_LongII_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_Generic_LongII_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_Generic_LongII_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">11.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">11.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">11.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>

35
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/device/reductionLongII_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D_long.cu.h" #include "Thread2D_long.cu.h"
#include "cudas.h" #include "cudas.h"
#include "Reduction.cu.h" #include "Reduction.cu.h"
@@ -23,9 +23,15 @@ static __device__ void addAtomic(long* ptrX , long y);
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__global__ void KLongProtocoleII(long* ptrSumGM) __global__
{ void KLongProtocoleII(long* ptrSumGM) {
// TODO ReductionLongII // TODO ReductionLongII
//
extern __shared__ long tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
Reduction::reduce(add, addAtomic, tabSM, ptrSumGM);
} }
/*--------------------------------------*\ /*--------------------------------------*\
@@ -35,8 +41,9 @@ __global__ void KLongProtocoleII(long* ptrSumGM)
/** /**
* TID partout en tabSM * TID partout en tabSM
*/ */
__device__ void reductionIntraThread(long* tabSM) __device__
{ // Rappel : Dans le protocoleII on cherche a calculer void reductionIntraThread(long* tabSM) {
// Rappel : Dans le protocoleII on cherche a calculer
// //
// x=x+i avec i in [0,N] // x=x+i avec i in [0,N]
// //
@@ -70,6 +77,11 @@ __device__ void reductionIntraThread(long* tabSM)
// pour TID utiliser const long TID=Thread2D_long::tid(); // (nouvelle methode) // pour TID utiliser const long TID=Thread2D_long::tid(); // (nouvelle methode)
// pour TID_LOCAL utiliser const int TID_LOCAL=Thread2D::tidLocal(); // (methode habituelle) // pour TID_LOCAL utiliser const int TID_LOCAL=Thread2D::tidLocal(); // (methode habituelle)
//
const long TID=Thread2D_long::tid();
const int TID_LOCAL=Thread2D::tidLocal();
tabSM[TID_LOCAL] = TID;
} }
@@ -78,9 +90,10 @@ __device__ void reductionIntraThread(long* tabSM)
\*---------------------------*/ \*---------------------------*/
__inline__ __inline__
__device__ long add(long x , long y) __device__
{ long add(long x , long y) {
// TODO ReductionLongII // TODO ReductionLongII
return x + y;
} }
/** /**
@@ -90,15 +103,16 @@ __device__ long add(long x , long y)
* *
* ou la technique du lock vu precedement! * ou la technique du lock vu precedement!
* *
* Question : atomicAdd pour les long existe? * Question : atomicAdd pour les long existe? non ;(
*/ */
__device__ int volatile mutex = 0; //variable global __device__ int volatile mutex = 0; //variable global
__device__ void addAtomic(long* ptrX , long y) __device__
{ void addAtomic(long* ptrX , long y) {
Lock locker(&mutex); Lock locker(&mutex);
locker.lock(); locker.lock();
// TODO ReductionLongII // TODO ReductionLongII
*ptrX = *ptrX + y;
locker.unlock(); locker.unlock();
} }
@@ -106,4 +120,3 @@ __device__ void addAtomic(long* ptrX , long y)
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\
|* End *| |* End *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/

16
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.cu
View File

@@ -27,24 +27,26 @@ extern __global__ void KLongProtocoleII(long* ptrSumGM);
ReductionLongII::ReductionLongII(const Grid& grid , long* ptrSum,bool isVerbose) : ReductionLongII::ReductionLongII(const Grid& grid , long* ptrSum,bool isVerbose) :
//RunnableGPU(grid, "Reduce_Generic_LongII_" + to_string(grid.threadCounts()),isVerbose), // classe parente //RunnableGPU(grid, "Reduce_Generic_LongII_" + to_string(grid.threadCounts()),isVerbose), // classe parente
RunnableGPU(grid, "Reduce_Generic_LongII",isVerbose), // classe parente RunnableGPU(grid, "Reduce_Generic_LongII",isVerbose), // classe parente
ptrSum(ptrSum) ptrSum(ptrSum) {
{
// TODO ReductionLongII // TODO ReductionLongII
this->sizeSM = -1; this->sizeSM = grid.threadByBlock() * sizeof(long);
GM::mallocLong0(&ptrSumGM);
} }
ReductionLongII::~ReductionLongII() ReductionLongII::~ReductionLongII() {
{
// TODO ReductionLongII // TODO ReductionLongII
GM::free(ptrSumGM);
} }
/*--------------------------------------*\ /*--------------------------------------*\
|* Methode *| |* Methode *|
\*-------------------------------------*/ \*-------------------------------------*/
void ReductionLongII::run() void ReductionLongII::run() {
{
// TODO ReductionLongII // TODO ReductionLongII
KLongProtocoleII<<<dg, db, this->sizeSM>>>(ptrSumGM);
GM::memcpyDToH_long(ptrSum, ptrSumGM);
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\

3
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/long/PII/host/ReductionLongII.h
View File

@@ -8,8 +8,7 @@
|* Declaration *| |* Declaration *|
\*---------------------------------------------------------------------*/ \*---------------------------------------------------------------------*/
class ReductionLongII: public RunnableGPU class ReductionLongII: public RunnableGPU {
{
/*--------------------------------------*\ /*--------------------------------------*\
|* Constructor *| |* Constructor *|
\*-------------------------------------*/ \*-------------------------------------*/

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

@@ -59,8 +59,8 @@ void generic() {
VTReductionGenericLongII test3; VTReductionGenericLongII test3;
// test1.run(); // test1.run();
test2.run(); // test2.run();
// test3.run(); test3.run();
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\

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

@@ -76,16 +76,16 @@ void reduction_generic(bool& isOk) {
// } // }
// IntII // IntII
{
UseReductionIntII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE);
}
// LongII
// { // {
// UseReductionLongII algo(IS_VERBOSE); // UseReductionIntII algo(IS_VERBOSE);
// isOk &= algo.isOk(IS_VERBOSE); // isOk &= algo.isOk(IS_VERBOSE);
// } // }
// LongII
{
UseReductionLongII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE);
}
} }
/*----------------------------------------------------------------------*\ /*----------------------------------------------------------------------*\