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

[GPU.git] / WCudaMSE / Student_Cuda_Image / src / cpp / core / 04_RayTracing / moo / device / RayTracingDevice.cu

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 3d80bc4..97fc839 100644 (file)
--- 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
@@ -1,39 +1,106 @@
 #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;
         }