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

feat/reduction): add reduction generic PI int

Browse Source
This commit is contained in:
Rémi Heredero
2025-11-17 18:34:58 +01:00
parent 135f425fd7
commit 78fb87e978
8 changed files with 429 additions and 43 deletions
Download Patch File Download Diff File Expand all files Collapse all files

212
Student_Cuda_Tools_Reduction/out/Reduce_Generic_IntI_justesse.html Normal file
View File

@@ -0,0 +1,212 @@
<!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_IntI_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_IntI_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="#TestReductionGenericI">TestReductionGenericI</a></td>
<td style="width:10%" class="tablecell_success">13</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="TestReductionGenericI"></a>Suite: TestReductionGenericI</h3>
<table summary="Details for suite TestReductionGenericI" 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">1.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_Generic_IntI_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_IntI_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_IntI_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>

20
Student_Cuda_Tools_Reduction/src/core/01_algorithme/generic/Reduction.cu.h
View File

@@ -55,9 +55,11 @@ class Reduction {
//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
reductionIntraBlock<T>(OP,tabSM);
// TODO ReductionGeneric
// Meme principe que ReductionAdd
reductionInterBlock<T>(ATOMIC_OP,tabSM, ptrResultGM);
}
private:
@@ -76,6 +78,11 @@ class Reduction {
// TODO ReductionGeneric
// Meme principe que ReductionAdd
// OP est la variable representant l'operateur binaire
const int tidLocal = Thread2D::tidLocal();
if (tidLocal < middle) {
tabSM[tidLocal] = OP(tabSM[tidLocal], tabSM[tidLocal + middle]);
}
}
/**
@@ -88,6 +95,15 @@ class Reduction {
// TODO ReductionGeneric
// Meme principe que ReductionAdd
// OP est la variable representant l'operateur binaire
const int NB_THEAD_LOCAL = Thread2D::nbThreadLocal();
int middle = NB_THEAD_LOCAL >> 1;
while (middle > 0) {
ecrasement<T>(OP,tabSM, middle);
__syncthreads();
middle = middle >> 1;
}
}
/*--------------------------------------*\
@@ -101,6 +117,10 @@ class Reduction {
// TODO ReductionGeneric
// Meme principe que ReductionAdd
// ATOMIC_OP est la variable representant l'operateur binaire atomic
if (Thread2D::tidLocal() == 0) {
ATOMIC_OP(ptrResultGM, tabSM[0]);
}
}
};

34
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/device/reductionIntI_device.cu
View File

@@ -1,4 +1,4 @@
#include "Thread1D.cu.h"
#include "Thread2D.cu.h"
#include "cudas.h"
#include "Reduction.cu.h"
@@ -23,9 +23,15 @@ static __device__ void addAtomicV2(int* ptrX , int y);
/**
* 1 partout en tabSM
*/
__global__ void KIntProtocoleI(int* ptrSumGM)
{
__global__
void KIntProtocoleI(int* ptrSumGM) {
// TODO ReductionIntI
extern __shared__ int tabSM[];
reductionIntraThread(tabSM);
__syncthreads();
Reduction::reduce(add, addAtomicV1, tabSM, ptrSumGM);
}
/*--------------------------------------*\
@@ -35,9 +41,11 @@ __global__ void KIntProtocoleI(int* ptrSumGM)
/**
* 1 partout en tabSM
*/
__device__ void reductionIntraThread(int* tabSM)
{
__device__
void reductionIntraThread(int* tabSM) {
// TODO ReductionIntI
const int tidLocal = Thread2D::tidLocal();
tabSM[tidLocal] = 1;
}
/*----------------------------*\
@@ -45,18 +53,20 @@ __device__ void reductionIntraThread(int* tabSM)
\*---------------------------*/
__inline__
__device__ int add(int x , int y)
{
__device__
int add(int x , int y) {
// TODO ReductionIntI
return x + y;
}
/**
* Utiliser la methode system : atomicAdd(pointeurDestination, valeurSource);
*/
__inline__
__device__ void addAtomicV1(int* ptrX , int y)
{
__device__
void addAtomicV1(int* ptrX , int y) {
// TODO ReductionIntI
atomicAdd(ptrX, y);
}
/**
@@ -65,12 +75,13 @@ __device__ void addAtomicV1(int* ptrX , int y)
* Necessaire aussi pour des objets par exemple
*/
__device__ int volatile mutex = 0; //variable global
__device__ void addAtomicV2(int* ptrX , int y)
{
__device__
void addAtomicV2(int* ptrX , int y) {
Lock locker(&mutex);
locker.lock();
// TODO ReductionIntI
*ptrX = *ptrX + y;
locker.unlock();
}
@@ -78,4 +89,3 @@ __device__ void addAtomicV2(int* ptrX , int y)
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

21
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.cu
View File

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

3
Student_Cuda_Tools_Reduction/src/core/02_use_protocole/generic/int/PI/host/ReductionIntI.h
View File

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

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

@@ -29,8 +29,8 @@ static void generic();
int mainTest() {
// activer ci-dessous seulement le TP voulu (pas tous)
add();
//generic();
// add();
generic();
return EXIT_SUCCESS;
}

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

@@ -36,8 +36,8 @@ int mainUse() {
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_add(isOk); // voir code ci-dessous pour activer la version voulue
reduction_generic(isOk); // voir code ci-dessous pour activer la version voulue
print(isOk);
@@ -76,16 +76,16 @@ void reduction_generic(bool& isOk) {
}
// IntII
{
UseReductionIntII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE);
}
// {
// UseReductionIntII algo(IS_VERBOSE);
// isOk &= algo.isOk(IS_VERBOSE);
// }
// LongII
{
UseReductionLongII algo(IS_VERBOSE);
isOk &= algo.isOk(IS_VERBOSE);
}
// {
// UseReductionLongII algo(IS_VERBOSE);
// isOk &= algo.isOk(IS_VERBOSE);
// }
}
/*----------------------------------------------------------------------*\