Implémentation du raytracing pour Global Memory/Shared Memory/Constant Memory

[GPU.git] / WCudaMSE / Student_Cuda_Image / src / cpp / core / 04_RayTracing / moo / host / RayTracing.cu

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 36e3862..8742692 100644 (file)
--- 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
@@ -8,37 +8,60 @@ using namespace std;
 #include "RayTracing.h"
 #include "Device.h"
 #include "RayTracingDevice.h"
-
-RayTracing::RayTracing(int w, int h)
-    : w(w), h(h),
-      dg(8, 8, 1),
-      db(32, 32, 1),
-      title("Ray Tracing")
+#include "Sphere.h"
+#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* = 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()
     {
-    delete this->ptrDomaineMathInit;
     }
 
-void RayTracing::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath)
+void RayTracing::runGPU(uchar4* ptrDevPixels)
     {
-    rayTracing<<<dg,db>>>(ptrDevPixels, this->w, this->h, domaineMath, this->t);
+#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).
     }
 
 void RayTracing::animationStep()
     {
-    this->t += 0.1; // TODO
+    this->t += 0.005;
     }
 
 int RayTracing::getW()
@@ -61,19 +84,28 @@ string RayTracing::getTitle()
     return this->title;
     }
 
-Sphere* createSpheres(int n)
+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(AleaTools::uniformeAB(20, this->w / 10 - 1)));
-        spheres[i].setCentre(float3 {
-                float(AleaTools::uniformeAB(bord, this->w - bord)),
-                float(AleaTools::uniformeAB(bord, this->h - bord)),
-                float(AleaTools::uniformeAB(10, 2.0 * this->w))
-            });
-        spheres[i].setHue(float(AleaTools::uniformeAB(0, 1))
+        spheres[i].setR(float(this->alea.uniformeAB(RAYON_MIN, RAYON_MAX)));
+
+        const float3 centre =
+            {
+            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;
     }