srand(time(NULL));\r
}\r
\r
+AleaTools::AleaTools(uint seed)\r
+ {\r
+ srand(seed);\r
+ }\r
+\r
AleaTools::~AleaTools()\r
{\r
// rien\r
*ptrJ = s - w * (*ptrI);\r
}\r
\r
+ float2 IndiceTools::toPixelFloat(int s, int w)\r
+ {\r
+ const float y = (float)s / (float)w;\r
+ const float x = (float)s - (float)w * y;\r
+ return { x, y };\r
+ }\r
+\r
/**\r
* i[0,H[ j[0,W[ --> s[0,W*H[\r
*/\r
- int IndiceTools::toS(int W, int i, int j)\r
+ int IndiceTools::toS(int w, int i, int j)\r
{\r
- return (i * W) + j;\r
+ return (i * w) + j;\r
}\r
\r
}\r
*/\r
AleaTools();\r
\r
+ AleaTools(uint seed);\r
+\r
/*--------------------------------------*\\r
|* Destructor *|\r
\*-------------------------------------*/\r
#ifndef INDICE_TOOLS_CPU_H_\r
#define INDICE_TOOLS_CPU_H_\r
\r
+#include "cudaType_CPU.h"\r
+\r
/*----------------------------------------------------------------------*\\r
|* Declaration *|\r
\*---------------------------------------------------------------------*/\r
/**\r
* s[0,W*H[ --> i[0,H[ j[0,W[\r
*/\r
- static void toIJ(int s, int W, int* ptri, int* ptrj);\r
+ static void toIJ(int s, int w, int* ptri, int* ptrj);\r
+\r
+ /**\r
+ * s[0,W*H[ --> x[0,W[ y[0,H[\r
+ */\r
+ static float2 toPixelFloat(int s, int w);\r
\r
/**\r
* i[0,H[ j[0,W[ --> s[0,W*H[\r
*/\r
- static int toS(int W, int i, int j);\r
+ static int toS(int w, int i, int j);\r
\r
/*-------------------------------------*\\r
|* Attributs *|\r
#define INDICES_TOOLS_GPU_H_\r
\r
#include "both_define.h"\r
+#include "cudaType_CPU.h"\r
\r
/*----------------------------------------------------------------------*\\r
|* Declaration *|\r
* h = hauteur\r
*/\r
__BOTH__\r
- static void toIJ( int s, int w, int* ptrI, int* ptrJ)\r
+ static void toIJ( int s, int w, int* ptrI, int* ptrJ)\r
{\r
*ptrI = s / w;\r
*ptrJ = s - w * (*ptrI);\r
}\r
\r
+ /**\r
+ * s[0,W*H[ --> x[0,W[ y[0,H[\r
+ */\r
+ __BOTH__\r
+ static float2 toPixelFloat(int s, int w)\r
+ {\r
+ const int y = s / w;\r
+ const int x = s - w * y;\r
+ const float2 pixel = { (float)x, (float)y };\r
+ return pixel;\r
+ }\r
+\r
/**\r
* i[0,H[ j[0,W[ --> s[0,W*H[\r
* w = largeur\r
};\r
}\r
\r
-#endif \r
+#endif\r
\r
/*----------------------------------------------------------------------*\\r
|* End *|\r
// cublas\r
#include "cublas_v2.h" // Pour specifier la version qu'on veut! (definition CUBLASAPI)\r
//#include "cublas.h" // Pour specifier la version qu'on veut!(definition CUBLASAPI)\r
-#include "cublas_api.h" // Après include "cublas_v2.h" ou "cublas.h"\r
+#include "cublas_api.h" // Apr�s include "cublas_v2.h" ou "cublas.h"\r
\r
-\r
-\r
-\r
-typedef unsigned char uchar;\r
+#ifndef uchar\r
+ typedef unsigned char uchar;\r
+#endif\r
\r
/*----------------------------------------------------------------------*\\r
|* Declaration *|\r
--- /dev/null
+#ifndef RAY_TRACING_COMMON
+#define RAY_TRACING_COMMON
+
+#define NB_SPHERES 2000
+
+
+#define MEMORY_MODEL_GLOBAL 1
+#define MEMORY_MODEL_CONSTANT 2
+#define MEMORY_MODEL_SHARED 3
+
+#define CURRENT_MEMORY_MODEL MEMORY_MODEL_SHARED
+
+#endif
#include <iostream>
-#include <stdio.h>
+#include <cstdio>
+#include <cfloat>
using namespace std;
+#include "Indice1D.h"
#include "Indice2D.h"
#include "IndiceTools.h"
+#include "ColorTools.h"
#include "cudaTools.h"
#include "Device.h"
#include "RayTracingDevice.h"
-__global__
-void rayTracing(uchar4* ptrDevPixels, int w, int h, float t)
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+// Les spheres dans la constant memory.
+__constant__ Sphere CONST_SPHERES[NB_SPHERES];
+
+ConstantMemoryLink constantMemorySpheresLink()
+ {
+ float* ptrDevTabData;
+ size_t sizeAll = NB_SPHERES * sizeof(Sphere);
+ HANDLE_ERROR(cudaGetSymbolAddress((void**)&ptrDevTabData, CONST_SPHERES));
+ ConstantMemoryLink cmLink =
+ {
+ (void**)ptrDevTabData, NB_SPHERES, sizeAll
+ };
+ return cmLink;
+ }
+#endif
+
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ __global__
+ void rayTracing(uchar4* ptrDevPixels, const Sphere* ptrDevSpheres, int w, int h, float t)
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+ __global__
+ void rayTracing(uchar4* ptrDevPixels, int w, int h, float t)
+#endif
{
+
const int TID = Indice2D::tid();
const int NB_THREAD = Indice2D::nbThread();
const int WH = w * h;
- uchar4 color;
- color.w = 255; // Par défaut, l'image est opaque.
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ __shared__ Sphere tabSpheresSM[NB_SPHERES];
- double x, y;
- int pixelI, pixelJ;
+ // Copy the global memory to shared memory.
+ const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
+ const int TID_LOCAL = Indice1D::tidLocal();
+ int sLocal = TID_LOCAL;
+ while (sLocal < NB_SPHERES)
+ {
+ tabSpheresSM[sLocal] = ptrDevSpheres[sLocal];
+ sLocal += NB_THREAD_LOCAL;
+ }
+
+ __syncthreads();
+#endif
int s = TID;
while (s < WH)
{
- IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ)
+ const float2 pixel = IndiceTools::toPixelFloat(s, w);
+
+ // Les données de la sphère sélectionnée.
+ float selectedSphereDistance = FLT_MAX;
+ float brightness = 0;
+ float hue = 0;
+
+ // Sélectionne la sphere la plus proche, il peut ne pas y avoir de sphère.
+ for (int i = 0; i < NB_SPHERES; ++i)
+ {
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL
+ const Sphere& sphere = ptrDevSpheres[i];
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+ const Sphere& sphere = CONST_SPHERES[i];
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ const Sphere& sphere = tabSpheresSM[i];
+#endif
+ const float hSquare = sphere.hSquare(pixel);
+
+ if (sphere.isBelow(hSquare))
+ {
+ const float dz = sphere.dz(hSquare);
+ const float sphereDistance = sphere.distance(dz);
- // (i,j) domaine écran.
- // (x,y) domaine math.
- // domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y).
+ if (sphereDistance < selectedSphereDistance)
+ {
+ selectedSphereDistance = sphereDistance;
+ brightness = sphere.brightness(dz);
+ hue = sphere.hue(t);
+ }
+ }
+ }
- // newtonMath.colorXY(&color, x, y);
+ // Si une sphere a pu être trouvée.
+ if (selectedSphereDistance != FLT_MAX)
+ {
+ ColorTools::HSB_TO_RVB(hue, 1, brightness, &ptrDevPixels[s]);
+ }
- ptrDevPixels[s] = color;
+ ptrDevPixels[s].w = 255;
s += NB_THREAD;
}
#ifndef RAY_TRACING_DEVICE_H
#define RAY_TRACING_DEVICE_H
-__global__
-void rayTracing(uchar4* ptrDevPixels, int w, int h, float t);
+#include "ConstantMemoryLink.h"
+
+#include "Sphere.h"
+#include "RayTracingCommon.h"
+
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+
+ __global__
+ void rayTracing(uchar4* ptrDevPixels, const Sphere* ptrDevSpheres, int w, int h, float t);
+
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+
+ /**
+ * Accès à la mémoire constante contenant les sphères.
+ */
+ ConstantMemoryLink constantMemorySpheresLink();
+
+ __global__
+ void rayTracing(uchar4* ptrDevPixels, int w, int h, float t);
+
+#endif
#endif
class Sphere
{
public:
- __host__
+ /*__host__
Sphere(float3 centre, float r, float hue) :
centre(centre),
hueInitial(hue)
{
this->setR(r);
- }
+ }*/
__host__
+ __device__
Sphere()
{
}
+ /*Sphere(const Sphere& other) :
+ r(other.r),
+ centre(other.centre),
+ rCarre(other.rCarre),
+ T(other.T)
+ {
+ }*/
+
__host__
void setCentre(const float3& centre)
{
void setR(float r)
{
this->r = r;
- this->rCarre = r*r;
+ this->rCarre = r * r;
}
__host__
void setHueInitial(float hue)
{
- this->hueInitial = hue;
+ this->T = asinf(2 * hue - 1) - 3 * PI / 2;
}
+ /*
+ * Calcul de h².
+ */
__device__
- float hCarre(float2 xySol)
+ float hSquare(float2 xySol) const
{
- float a = (centre.x - xySol.x);
- float b = (centre.y - xySol.y);
+ const float a = (this->centre.x - xySol.x);
+ const float b = (this->centre.y - xySol.y);
return a * a + b * b;
}
+ /*
+ * Est-ce que la sphre se trouve en dessous d'un point (x, y) en fonction de h².
+ * Voir la méthode 'hCarre(..)'.
+ */
__device__
- bool isEnDessous(float hCarre)
+ bool isBelow(float hCarre) const
{
- return hCarre < rCarre;
+ return hCarre < this->rCarre;
}
+ /*
+ * Hauteur du point de collision du rayon par rapport au centre de la sphère.
+ */
__device__
- float dz(float hCarre)
+ float dz(float hCarre) const
{
return sqrtf(rCarre - hCarre);
}
__device__
- float brightness(float dz)
+ float distance(float dz) const
{
- return dz / r;
+ return this->centre.z - dz;
}
+ /**
+ * Renvoie la le B de HSB compris entre 0 et 1.
+ */
__device__
- float distance(float dz)
+ float brightness(float dz) const
{
- return centre.z - dz;
+ return dz / this->r;
}
+ /*
+ * La couleur
+ */
__device__
- float getHueInitial()
+ float hue(float t) const
{
- return this->hueInitial;
- }
-
- __device__
- float hue(float t) // usefull for animation
- {
- return 0.5 + 0.5 * sin(t + T + 3 * PI / 2);
+ return 0.5 + 0.5 * sinf(t + this->T + 3 * PI / 2);
}
private:
- float r;
- float3 centre;
- float hueInitial;
- float rCarre;
+ float r; // Rayon.
+ float3 centre; // Position.
+
+ float rCarre; // Précalul de r².
float T; // Utilisé pour l'animation.
};
#include "Device.h"
#include "RayTracingDevice.h"
#include "Sphere.h"
-
-RayTracing::RayTracing(int w, int h)
- : w(w), h(h),
- dg(8, 8, 1),
- db(32, 32, 1),
- title("Ray Tracing")
+#include "RayTracingCommon.h"
+
+RayTracing::RayTracing(int w, int h) :
+ t(0), w(w), h(h),
+ dg(8, 8, 1),
+ db(32, 32, 1),
+ alea(42),
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL
+ title("Ray Tracing (Global memory)")
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+ title("Ray Tracing (Constant memory)")
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ title("Ray Tracing (Shared memory)")
+#endif
{
Device::assertDim(dg, db);
- const int nbSpheres = 10;
- Sphere* shperes = this->createSpheres(10);
-
- // Copie les spheres dans la constant memory.
+ Sphere* spheres = this->createSpheres(NB_SPHERES);
+
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ // Copie des spheres dans la global memory.
+ const size_t sizeSpheres = NB_SPHERES * sizeof(Sphere);
+ HANDLE_ERROR(cudaMalloc(&this->ptrDevSpheres, sizeSpheres));
+ HANDLE_ERROR(cudaMemcpy(this->ptrDevSpheres, spheres, sizeSpheres, cudaMemcpyHostToDevice));
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
+ // Copie des spheres en constant memory.
+ ConstantMemoryLink cmSpheresLink = constantMemorySpheresLink();
+ Sphere* ptrDevConstSperes = (Sphere*)cmSpheresLink.ptrDevTab;
+ size_t sizeALL = cmSpheresLink.sizeAll;
+ HANDLE_ERROR(cudaMemcpy(ptrDevConstSperes, spheres, sizeALL, cudaMemcpyHostToDevice));
+#endif
+
+ cout << "sizeof(Sphere): " << sizeof(Sphere) << endl;
}
RayTracing::~RayTracing()
void RayTracing::runGPU(uchar4* ptrDevPixels)
{
+#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
+ rayTracing<<<dg,db>>>(ptrDevPixels, this->ptrDevSpheres, this->w, this->h, this->t);
+#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
rayTracing<<<dg,db>>>(ptrDevPixels, this->w, this->h, this->t);
+#endif
- // HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
+ HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG).
}
void RayTracing::animationStep()
{
- this->t += 0.1; // TODO.
+ this->t += 0.005;
}
int RayTracing::getW()
Sphere* RayTracing::createSpheres(int n)
{
Sphere* spheres = new Sphere[n];
- const float bord = 200;
+ const int RAYON_MIN = 10;
+ const int RAYON_MAX = this->w / 10 - 1;
+
+ const float bord = RAYON_MAX + 1;
for (int i = 0; i < n; ++i)
{
- spheres[i].setR(float(this->alea.uniformeAB(20, this->w / 10 - 1)));
+ spheres[i].setR(float(this->alea.uniformeAB(RAYON_MIN, RAYON_MAX)));
- float3 centre
+ const float3 centre =
{
- float(this->alea.uniformeAB(double(bord), double(this->w - bord))),
- float(this->alea.uniformeAB(double(bord), double(this->h - bord))),
- float(this->alea.uniformeAB(10.0, 2.0 * this->w))
+ float(this->alea.uniformeAB(double(bord), double(this->w - bord))), // x.
+ float(this->alea.uniformeAB(double(bord), double(this->h - bord))), // y.
+ float(this->alea.uniformeAB(10.0, 2.0 * this->w)) // z.
};
spheres[i].setCentre(centre);
spheres[i].setHueInitial(float(this->alea.uniformeAB(0.0, 1.0)));
}
+
+ return spheres;
}
std::string getTitle(void) /*override*/;
private:
+ /**
+ * Crée un tablean de 'n' sphères dont le rayon, la couleur et la position sont générés aléatoirement.
+ */
Sphere* createSpheres(int n);
+ Sphere* ptrDevSpheres; // Pointeur sur la mémoire du GPU.
+
AleaTools alea;
float t;
Image* RayTracingProvider::createGL()
{
- ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0);
+ ColorRGB_01* ptrColorTitre = new ColorRGB_01(255, 255, 255);
return new Image(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel
}
\r
int mainGL(const vector<string>& args)\r
{\r
- const string defaultCommand = "demo";\r
+ const string defaultCommand = "raytracing";\r
const string command = args.size() > 0 ? args[0] : defaultCommand;\r
\r
// AutoViewer<Rippling0Image, Rippling0Provider> rippling0(true, true, 10, 10); // Warmup.\r
AutoViewer<ImageFonctionel, NewtonProvider> newtown(true, true, 10, 10);\r
else if (command == "heat-transfert")\r
AutoViewer<Image, HeatTransfertProvider> heatTransfert(true, false, 10, 10);\r
+ else if (command == "raytracing")\r
+ AutoViewer<Image, RayTracingProvider> rayTracing(true, true, 20, 20);\r
else if (command == "convolution")\r
{\r
const string videoPath = args.size() >= 2 ? args[1] : "/media/Data/Video/nasaFHD_short.avi"; // Vidéo par défaut si pas donnée en paramètre.\r
cout << "Command unknown: " << command << endl;\r
}\r
\r
- // AutoViewer<ImageFonctionel, RayTracingProvider> rayTracing(true, true, 20, 20); // Commenté car projet approfondit.\r
-\r
return EXIT_SUCCESS;\r
}\r
projects / GPU.git / commitdiff