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

feat(rippling): add lab Image Rippling 2D

Browse Source 
- Full compute - Demi compute (optimisation)
This commit is contained in:
Rémi Heredero
2025-11-14 12:34:05 +01:00
parent 3cc8fef087
commit 2daed2f9eb
27 changed files with 493 additions and 156 deletions
Download Patch File Download Diff File Expand all files Collapse all files

56
Student_Cuda_Image/src/core/01_student/01_Rippling/device/math/RipplingMath.cu.h
View File

@@ -13,26 +13,24 @@
|* Public *|
\*-------------------------------------*/
class RipplingMath
{
class RipplingMath {
/*--------------------------------------*\
|* Constructeur *|
\*-------------------------------------*/
public:
__device__ RipplingMath(uint w , uint h , float t) :
__device__
RipplingMath(uint w , uint h , float t) :
dim2(w / 2), //
t(t)
{
// rien
t(t) {
// rien
}
__device__
virtual ~RipplingMath()
{
virtual ~RipplingMath() {
// rien
}
}
/*--------------------------------------*\
|* Methode *|
@@ -41,8 +39,7 @@ class RipplingMath
public:
__device__
void colorIJ(uchar4* ptrColorIJ , int i , int j)
{
void colorIJ(uchar4* ptrColorIJ , int i , int j) {
uchar levelGris = levelGray(i, j); // update levelGris
ptrColorIJ->x = levelGris;
@@ -58,13 +55,13 @@ class RipplingMath
// Etape 2: Puis une fois que l'image grise est valider, attaquer rippling
// debug temp
// {
// ptrColorIJ->x = 128;
// ptrColorIJ->y = 128;
// ptrColorIJ->z = 128;
// ptrColorIJ->w = 255; // opacity facultatif
// }
}
// {
// ptrColorIJ->x = 128;
// ptrColorIJ->y = 128;
// ptrColorIJ->z = 128;
// ptrColorIJ->w = 255; // opacity facultatif
// }
}
private:
@@ -74,24 +71,31 @@ class RipplingMath
__inline__
__device__
uchar levelGray(int i , int j )
{
uchar levelGray(int i , int j ) {
float result;
dij(i, j, &result); // warning : dij return void. Ne peut pas etre "imbriquer dans une fonction"
dij(i, j, &result); // warning : dij return void. Ne peut pas etre "imbriquer dans une fonction"
result = result / 10.f;
result = 128.f + 127.f * cosf(result - (t / 7.f)) / (result + 1.f);
return (uchar)result;
// TODO Rippling GPU : cf formules math rippling.pdf (attribut dim2 = dim/2
}
}
__inline__
__device__
void dij(int i , int j , float* ptrResult)
{
void dij(int i , int j , float* ptrResult) {
//TODO Rippling GPU cf fonction math pdf
float fi = i - dim2;
float fj = j - dim2;
float d = sqrtf(fi * fi + fj * fj);
*ptrResult = d;
// Ne pas utiliser la fonction pow pour elever au carrer !
// Utiliser l'opérateur *
}
}
/*--------------------------------------*\
|* Attribut *|
@@ -103,7 +107,7 @@ class RipplingMath
int dim2; // dim2=dim/2
float t;
};
};
/*----------------------------------------------------------------------*\
|* End *|

61
Student_Cuda_Image/src/core/01_student/01_Rippling/device/ripplingDevice.cu
View File

@@ -18,12 +18,11 @@ static __device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , flo
|* Implementation *|
\*---------------------------------------------------------------------*/
__global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t)
{
ripplingBaseline(tabPixelsGM, w, h, t);
// ripplingDemi(tabPixelsGM, w, h, t);
__global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t) {
// ripplingBaseline(tabPixelsGM, w, h, t);
ripplingDemi(tabPixelsGM, w, h, t);
// ripplingQuart(tabPixelsGM, w, h, t);
}
}
/*----------------------------------------------------------------------*\
|* private *|
@@ -33,23 +32,36 @@ __global__ void rippling(uchar4* tabPixelsGM , uint w , uint h , float t)
* v1
*/
__inline__
__device__ void ripplingBaseline(uchar4* tabPixelsGM , uint w , uint h , float t)
{
__device__
void ripplingBaseline(uchar4* tabPixelsGM , uint w , uint h , float t) {
// TODO instacier RipplingMath
RipplingMath rgM = RipplingMath(w, h, t);
const int TID = Thread2D::tid();
const int NB_THREAD = Thread2D::nbThread();
const int WH = w * h;
// TODO Rippling GPU pattern entrelacement
// TODO Rippling GPU pattern entrelacement
int s = TID;
while(s < WH) {
int i = s / w;
int j = s % w;
uchar4 color;
rgM.colorIJ(&color, i, j);
tabPixelsGM[s] = color;
s += NB_THREAD;
}
}
/**
* v2 : optimisation
*/
__inline__
__device__ void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t)
{
__device__
void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t) {
// Indication:
// (I1) Utiliser la symetrie horizontale de l'image
// (I2) Calculer que la demi partie superieur
@@ -57,14 +69,36 @@ __device__ void ripplingDemi(uchar4* tabPixelsGM , uint w , uint h , float t)
// Partez de la fin de l'image, peut-etre
// TODO Rippling GPU
RipplingMath rgM = RipplingMath(w, h, t);
const int TID = Thread2D::tid();
const int NB_THREAD = Thread2D::nbThread();
const int WH = w * h;
const int WH2 = (WH / 2) + 1;
int s = TID;
while(s < WH2) {
int i = s / w;
int j = s % w;
uchar4 color;
rgM.colorIJ(&color, i, j);
tabPixelsGM[s] = color;
// Miroir
int si = WH - s - 1;
tabPixelsGM[si] = color;
s += NB_THREAD;
}
}
/**
* v3 : optimsation : defi (difficile) (pas necessaire pour test performance)
*/
__inline__
__device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t)
{
__device__
void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t) {
// Indication:
// (I1) Utiliser la symetrie horizontale et verticale de l'image
// (I2) Calculer que le quart en huat a gauche
@@ -76,9 +110,8 @@ __device__ void ripplingQuart(uchar4* tabPixelsGM , uint w , uint h , float t)
// (C1) Utiliser toujours le pattern d'entrelacement
// TODO Rippling GPU
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

23
Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.cu
View File

@@ -24,19 +24,17 @@ extern __global__ void rippling(uchar4* tabPixelsGM,uint w, uint h,float t);
\*-------------------------*/
Rippling::Rippling(const Grid& grid , uint w , uint h , float dt , bool isVerbose) :
Animable_I<uchar4>(grid, w, h, "Rippling-Cuda-uchar4", isVerbose) // super classe
{
Animable_I<uchar4>(grid, w, h, "Rippling-Cuda-uchar4", isVerbose) { // super classe
assert(w == h); // specific rippling
// Animation
this->dt = dt;
this->t = 0; // protected dans Animable
}
}
Rippling::~Rippling()
{
Rippling::~Rippling() {
// rien
}
}
/*-------------------------*\
|* Methode *|
@@ -48,25 +46,22 @@ Rippling::~Rippling()
*
* Note : domaineMath pas use car image pas zoomable
*/
void Rippling::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath)
{
void Rippling::process(uchar4* tabPixelsGM , uint w , uint h , const DomaineMath& domaineMath) {
// TODO Rippling
// lancer le kernel avec <<<dg,db>>>
// le kernel est importer ci-dessus (ligne 19)
assert(false); // to delete once implement
}
rippling<<<dg,db>>>(tabPixelsGM,w,h,t);
}
/**
* Override
* Call periodicly by the API
*/
void Rippling::animationStep()
{
void Rippling::animationStep() {
t += dt;
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

5
Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling.h
View File

@@ -10,8 +10,7 @@
|* Declaration *|
\*---------------------------------------------------------------------*/
class Rippling: public Animable_I<uchar4>
{
class Rippling: public Animable_I<uchar4> {
/*--------------------------------------*\
|* Constructor *|
\*-------------------------------------*/
@@ -50,7 +49,7 @@ class Rippling: public Animable_I<uchar4>
// Inputs
float dt;
};
};
/*----------------------------------------------------------------------*\
|* End *|

21
Student_Cuda_Image/src/core/01_student/01_Rippling/host/Rippling_BestGrid.h
View File

@@ -11,30 +11,27 @@
|* Impelmentation *|
\*---------------------------------------------------------------------*/
namespace rippling
{
namespace rippling {
class BestGrid
{
class BestGrid {
public:
static Grid get()
{
static Grid get() {
const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP();
// TODO Rippling
dim3 dg(MP, 1, 1);
dim3 db(CORE_MP, 4, 1);
Grid grid(dg, db);
return grid;
// to remove once coded
{
Couts::redln("aie aie aie, your best grid won t build itself");
assert(false);
}
}
};
}
}
/*----------------------------------------------------------------------*\
|* End *|

26
Student_Cuda_Image/src/main/main.cpp
View File

@@ -24,34 +24,32 @@ extern int mainTest();
|* Implementation *|
\*---------------------------------------------------------------------*/
int main(int argc , char** argv)
{
int main(int argc , char** argv) {
CudaContextImage cudaContext;
// public
{
cudaContext.deviceId = 2; // in [0,2] width Server Cuda3
cudaContext.launchMode = LaunchModeImage::IMAGE; // IMAGE BENCHMARKING BRUTFORCE TESTING
cudaContext.deviceId = 2; // in [0,2] width Server Cuda3
cudaContext.launchMode = LaunchModeImage::TESTING; // IMAGE BENCHMARKING BRUTFORCE TESTING
cudaContext.deviceDriver = DeviceDriver::LOAD_CURRENT; // LOAD_CURRENT LOAD_ALL
cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT
cudaContext.deviceDriver = DeviceDriver::LOAD_CURRENT; // LOAD_CURRENT LOAD_ALL
cudaContext.deviceInfo = DeviceInfo::ALL_SIMPLE; // NONE ALL ALL_SIMPLE CURRENT
}
// private
{
cudaContext.args.argc = argc;
cudaContext.args.argv = argv;
cudaContext.args.argc = argc;
cudaContext.args.argv = argv;
cudaContext.mainImage = mainImage;
cudaContext.mainBenchmark = mainBenchmark;
cudaContext.mainBrutforce = mainBrutforce;
cudaContext.mainTest = mainTest;
cudaContext.mainImage = mainImage;
cudaContext.mainBenchmark = mainBenchmark;
cudaContext.mainBrutforce = mainBrutforce;
cudaContext.mainTest = mainTest;
}
return cudaContext.process();
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

36
Student_Cuda_Image/src/main/mainBenchmark.cpp
View File

@@ -32,8 +32,7 @@ static void raytracingCM2SM();
|* Implementation *|
\*---------------------------------------------------------------------*/
int mainBenchmark()
{
int mainBenchmark() {
// Please, un a la fois!
rippling();
// mandelbrot(); // Conseil : use nFixe (by example nMin=nMax=80)
@@ -44,67 +43,60 @@ int mainBenchmark()
// raytracingCM2SM();
return EXIT_SUCCESS;
}
}
/*--------------------------------------*\
|* Private *|
\*-------------------------------------*/
void rippling()
{
void rippling() {
const double DURATION_MAX_S = 8;
RipplingProvider provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
void mandelbrot()
{
void mandelbrot() {
const double DURATION_MAX_S = 8;
MandelbrotProvider provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
void raytracingGM()
{
void raytracingGM() {
const double DURATION_MAX_S = 8;
RaytracingProviderGM provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
void raytracingCM()
{
void raytracingCM() {
const double DURATION_MAX_S = 8;
RaytracingProviderCM provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
void raytracingSM()
{
void raytracingSM() {
const double DURATION_MAX_S = 8;
RaytracingProviderSM provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
void raytracingCM2SM()
{
void raytracingCM2SM() {
const double DURATION_MAX_S = 8;
RaytracingProviderCM2SM provider;
BenchmarkImage<uchar4>::run(&provider, DURATION_MAX_S);
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

40
Student_Cuda_Image/src/main/mainBruteforce.cpp
View File

@@ -40,8 +40,7 @@ static void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const
|* Implementation *|
\*---------------------------------------------------------------------*/
int mainBrutforce()
{
int mainBrutforce() {
Matlab matlab;
// Please, un a la fois!
@@ -55,73 +54,66 @@ int mainBrutforce()
matlab.play();
return EXIT_SUCCESS;
}
}
/*--------------------------------------*\
|* Private *|
\*-------------------------------------*/
void rippling(Matlab* ptrMatlab)
{
void rippling(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.01; // 0.9 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RipplingProvider provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
void mandelbrot(Matlab* ptrMatlab)
{
void mandelbrot(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.4; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
MandelbrotProvider provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
void raytracingGM(Matlab* ptrMatlab)
{
void raytracingGM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderGM provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
void raytracingSM(Matlab* ptrMatlab)
{
void raytracingSM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderSM provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
void raytracingCM(Matlab* ptrMatlab)
{
void raytracingCM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderCM provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
void raytracingCM2SM(Matlab* ptrMatlab)
{
void raytracingCM2SM(Matlab* ptrMatlab) {
const double DURATION_MAX_S = 0.9; // 1 grid
const PlotType PLOT_TYPE = PlotType::ALL_GRAPHE;
RaytracingProviderCM2SM provider;
bruteForce<uchar4>((ProviderUse_I*)&provider, ptrMatlab, PLOT_TYPE, DURATION_MAX_S);
}
}
/*--------------------------------------*\
|* Tools *|
\*-------------------------------------*/
template<typename T>
void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotType& plotType , double durationMaxS)
{
void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotType& plotType , double durationMaxS) {
// Hardware
const int MP = Hardware::getMPCount();
const int CORE_MP = Hardware::getCoreCountMP();
@@ -140,7 +132,7 @@ void bruteForce(ProviderUse_I* ptrProviderUse , Matlab* ptrMatlab , const PlotTy
GridMaillage gridMaillage(iteratorDGx, iteratorDBx);
BruteForce::run(ptrProviderUse, &gridMaillage, ptrMatlab, plotType, durationMaxS);
}
}
/*----------------------------------------------------------------------*\
|* End *|

6
Student_Cuda_Image/src/main/mainImage.cpp
View File

@@ -24,8 +24,7 @@ using std::string;
|* Implementation *|
\*---------------------------------------------------------------------*/
int mainImage(const Args& args)
{
int mainImage(const Args& args) {
gpu::GLUTImageViewers::init(args.argc, args.argv); //only once
// ImageOption : (boolean,boolean) : (isSelection ,isAnimation,isOverlay,isShowHelp)
@@ -43,9 +42,8 @@ int mainImage(const Args& args)
gpu::GLUTImageViewers::runALL(); // Bloquant, Tant qu'une fenetre est ouverte
return EXIT_SUCCESS;
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/

21
Student_Cuda_Image/src/main/mainTest.cpp
View File

@@ -31,8 +31,7 @@ static void raytracing();
|* Public *|
\*-------------------------------------*/
int mainTest()
{
int mainTest() {
// activer ci-dessous seulement le TP voulu (pas tous)
rippling();
@@ -40,36 +39,33 @@ int mainTest()
//raytracing(); // voir code ci-dessous pour activer la version voulue
return EXIT_SUCCESS;
}
}
/*--------------------------------------*\
|* private *|
\*-------------------------------------*/
void rippling()
{
void rippling() {
VTRippling test1;
test1.run();
}
}
/**
* fp16 only
*/
void mandelbrot()
{
void mandelbrot() {
assert(mandelbrotReal::isFp16());
VTMandelbrot test1;
test1.run();
}
}
/**
* activer ci-dessous la version souhaiter
*/
void raytracing()
{
void raytracing() {
VTRaytracingGM test1;
VTRaytracingSM test2;
VTRaytracingCM test3;
@@ -79,9 +75,8 @@ void raytracing()
// test2.run();
// test3.run();
// test4.run();
}
}
/*----------------------------------------------------------------------*\
|* End *|
\*---------------------------------------------------------------------*/