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 1623820..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
@@ -9,19 +9,39 @@ using namespace std;
 #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()
@@ -30,14 +50,18 @@ 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()
@@ -63,20 +87,25 @@ string RayTracing::getTitle()
 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;
     }