From 1c4b2276e7157acde9a3014b68d5d1667a7d6a44 Mon Sep 17 00:00:00 2001 From: gburri Date: Thu, 22 Jan 2015 11:51:51 +0100 Subject: [PATCH] RayTracing. --- .../core/04_RayTracing/moo/RayTracingCommon.h | 13 ----- .../moo/device/RayTracingDevice.cu | 5 +- .../moo/device/RayTracingDevice.h | 2 +- .../core/04_RayTracing/moo/device/Sphere.h | 21 +------ .../core/04_RayTracing/moo/host/RayTracing.cu | 19 +++--- .../core/04_RayTracing/moo/host/RayTracing.h | 2 +- .../provider/RayTracingProvider.cpp | 15 +++-- .../provider/RayTracingProvider.h | 2 +- .../src/cpp/core/mainFreeGL.cpp | 58 +++++++++++++++++-- .../src/cpp/core/mainGL.cpp | 5 +- WCudaMSE/Student_Cuda_Image/src/cpp/main.cpp | 7 ++- 11 files changed, 87 insertions(+), 62 deletions(-) delete mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/RayTracingCommon.h diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/RayTracingCommon.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/RayTracingCommon.h deleted file mode 100644 index 0f2d3ae..0000000 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/RayTracingCommon.h +++ /dev/null @@ -1,13 +0,0 @@ -#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 diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu index 97fc839..22d05d1 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu @@ -68,7 +68,7 @@ ConstantMemoryLink constantMemorySpheresLink() float brightness = 0; float hue = 0; - // Sélectionne la sphere la plus proche, il peut ne pas y avoir de sphère. + // Sélectionne la sphère la plus proche, il peut ne pas y avoir de sphère trouvée. for (int i = 0; i < NB_SPHERES; ++i) { #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL @@ -78,6 +78,7 @@ ConstantMemoryLink constantMemorySpheresLink() #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED const Sphere& sphere = tabSpheresSM[i]; #endif + const float hSquare = sphere.hSquare(pixel); if (sphere.isBelow(hSquare)) @@ -94,7 +95,7 @@ ConstantMemoryLink constantMemorySpheresLink() } } - // Si une sphere a pu être trouvée. + // Si une sphère a pu être trouvée. if (selectedSphereDistance != FLT_MAX) { ColorTools::HSB_TO_RVB(hue, 1, brightness, &ptrDevPixels[s]); diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h index 350ba9c..0c207d4 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h @@ -4,7 +4,7 @@ #include "ConstantMemoryLink.h" #include "Sphere.h" -#include "RayTracingCommon.h" +#include "RayTracingParams.h" #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h index 3c8df6e..a4c49e9 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h @@ -7,28 +7,12 @@ class Sphere { public: - /*__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) { @@ -85,7 +69,8 @@ class Sphere } /** - * Renvoie la le B de HSB compris entre 0 et 1. + * Renvoie la le B de HSB compris entre 0 et 1 dépendant du paramètre 'dz' + * lui même compris entre 0 et le rayon 'r'. */ __device__ float brightness(float dz) const @@ -105,9 +90,7 @@ class Sphere private: float r; // Rayon. float3 centre; // Position. - float rCarre; // Précalul de r². - float T; // Utilisé pour l'animation. }; #endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu index 8742692..fbbd160 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu @@ -9,12 +9,12 @@ using namespace std; #include "Device.h" #include "RayTracingDevice.h" #include "Sphere.h" -#include "RayTracingCommon.h" +#include "RayTracingParams.h" -RayTracing::RayTracing(int w, int h) : +RayTracing::RayTracing(int w, int h, int dg, int db) : t(0), w(w), h(h), - dg(8, 8, 1), - db(32, 32, 1), + dg(dg, dg, 1), + db(db, db, 1), alea(42), #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL title("Ray Tracing (Global memory)") @@ -27,6 +27,11 @@ RayTracing::RayTracing(int w, int h) : Device::assertDim(dg, db); Sphere* spheres = this->createSpheres(NB_SPHERES); + cout << "sizeof(Sphere): " << sizeof(Sphere) << endl; + cout << "Nb of spheres: " << NB_SPHERES << endl; + cout << "sizeof(spheres): " << NB_SPHERES * sizeof(Sphere) << endl; + cout << "dg: " << this->dg.x << " x " << this->dg.y << endl; + cout << "db: " << this->db.x << " x " << this->db.y << endl; #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED // Copie des spheres dans la global memory. @@ -51,12 +56,12 @@ RayTracing::~RayTracing() void RayTracing::runGPU(uchar4* ptrDevPixels) { #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED - rayTracing<<>>(ptrDevPixels, this->ptrDevSpheres, this->w, this->h, this->t); + rayTracing<<>>(ptrDevPixels, this->ptrDevSpheres, this->w, this->h, this->t); #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT - rayTracing<<>>(ptrDevPixels, this->w, this->h, this->t); + rayTracing<<>>(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() diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h index 1f9536d..128d7e0 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h @@ -11,7 +11,7 @@ class RayTracing : public Animable_I { public: - RayTracing(int w, int h); + RayTracing(int w, int h, int dg, int db); ~RayTracing(); void runGPU(uchar4* ptrDevPixels) /*override*/; diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp index 2d82571..aa2e287 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp @@ -1,15 +1,18 @@ #include "RayTracingProvider.h" -RayTracing* RayTracingProvider::create() +RayTracing* RayTracingProvider::create(int dw, int dh, int dg, int db) { - int dw = 16 * 50; - int dh = 16 * 50; - - return new RayTracing(dw, dh); + return new RayTracing(dw, dh, dg, db); } Image* RayTracingProvider::createGL() { + + const int dw = 16 * 50; + const int dh = 16 * 50; + const int dg = 32; // 32x32. + const int db = 8; // 8x8. + ColorRGB_01* ptrColorTitre = new ColorRGB_01(255, 255, 255); - return new Image(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + return new Image(create(dw, dh, dg, db), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h index c4e560f..70315b4 100644 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h @@ -7,7 +7,7 @@ class RayTracingProvider { public: - static RayTracing* create(); + static RayTracing* create(int dw, int dh, int dg, int db); static Image* createGL(); }; diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainFreeGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainFreeGL.cpp index 2e54bef..e5d126d 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainFreeGL.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainFreeGL.cpp @@ -1,13 +1,14 @@ #include #include #include +using namespace std; #include "Device.h" #include "cudaTools.h" +#include "Chronos.h" -using std::cout; -using std::endl; -using std::string; +#include "RayTracingProvider.h" +#include "RayTracingParams.h" /*----------------------------------------------------------------------*\ |* Declaration *| @@ -33,17 +34,62 @@ int mainFreeGL(void); |* Implementation *| \*---------------------------------------------------------------------*/ +void run(uchar4* ptrDevPixels, int dw, int dh, int dg, int db, int nbFrame) + { + cout << "Raytracing Without GL "; +#if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL + cout << "(Global memory)" << endl; +#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT + cout << "(Constant memory)" << endl; +#elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED + cout << "(Shared memory)" << endl; +#endif + + cout << "Nb frames: " << nbFrame << endl; + + RayTracing* rayTracing = RayTracingProvider::create(dw, dh, dg, db); + + Chronos chronos; + chronos.start(); + + for (int i = 0; i < nbFrame; i++) + rayTracing->runGPU(ptrDevPixels); + + HANDLE_ERROR(cudaDeviceSynchronize()); + + chronos.stop(); + + cout << "Time: " << chronos << endl; + cout << "-------------------" << endl; + } + /*--------------------------------------*\ |* Public *| \*-------------------------------------*/ int mainFreeGL(void) { - cout << "\n[FPS] : Free GL, please wait ..." << endl; + const int NB_FRAME_TO_RENDER = 100; + const int dw = 16 * 50; + const int dh = 16 * 50; + + // dg : 8 -> 128. + const int dgStart = 8; + const int dgStop = 128; - rippling0FreeGL(1000); // bad technique + // db : 8 -> 32 + const int dbStart = 8; + const int dbStop = 32; + + uchar4* ptrDevPixels; + HANDLE_ERROR(cudaMalloc(&ptrDevPixels, dw * dh * sizeof(uchar4))); + + run(ptrDevPixels, dw, dh, 32, 32, NB_FRAME_TO_RENDER); + return EXIT_SUCCESS; - // TODO : add other tp here ... + for (int dg = dgStart; dg <= dgStop; dg *= 2) + for (int db = dbStart; db <= dbStop; db *= 2) + run(ptrDevPixels, dw, dh, dg, db, NB_FRAME_TO_RENDER); return EXIT_SUCCESS; } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp index 13dd2db..e353be7 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/mainGL.cpp @@ -20,15 +20,13 @@ using namespace std; int mainGL(const vector& args) { - const string defaultCommand = "raytracing"; + const string defaultCommand = "demo"; const string command = args.size() > 0 ? args[0] : defaultCommand; // AutoViewer rippling0(true, true, 10, 10); // Warmup. if (command == "rippling") - { AutoViewer rippling(true, true, 10, 10); - } else if (command == "mandelbrot") { const bool multiGPU = args.size() >= 2 && args[1] == "--mp"; @@ -55,6 +53,7 @@ int mainGL(const vector& args) AutoViewer fractalJulia(true, true, 180, 80, false); AutoViewer newtown(true, true, 260, 120, false); AutoViewer heatTransfert(true, false, 1200, 300, false); + AutoViewer rayTracing(true, true, 200, 80); GLUTImageViewers::runALL(); } else diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/main.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/main.cpp index 5e56e8a..46fbbce 100755 --- a/WCudaMSE/Student_Cuda_Image/src/cpp/main.cpp +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/main.cpp @@ -46,7 +46,7 @@ int main(int argc, char** argv) cout << "main" << endl; if (Device::isCuda()) - { + { GLUTImageViewers::init(argc, argv); //Device::printAll(); @@ -54,7 +54,8 @@ int main(int argc, char** argv) // Server Cuda1: in [0,5] // Server Cuda2: in [0,2] - int deviceId = 0; + const int deviceId = 5; + initCuda(deviceId); vector args; for (int i = 1; i < argc; ++i) @@ -106,7 +107,7 @@ int start(const vector& args) if (IS_GL) { - return mainGL(args); // Bloquant, Tant qu'une fenetre est ouverte + return mainGL(args); // Bloquant, Tant qu'une fenêtre est ouverte. } else { -- 2.43.0