WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu

   1 #include <iostream>
   2 #include <cstdio>
   3 #include <cfloat>
   4 using namespace std;
   5
   6 #include "Indice1D.h"
   7 #include "Indice2D.h"
   8 #include "IndiceTools.h"
   9 #include "ColorTools.h"
  10 #include "cudaTools.h"
  11 #include "Device.h"
  12
  13 #include "RayTracingDevice.h"
  14
  15 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
  16 // Les spheres dans la constant memory.
  17 __constant__ Sphere CONST_SPHERES[NB_SPHERES];
  18
  19 ConstantMemoryLink constantMemorySpheresLink()
  20     {
  21     float* ptrDevTabData;
  22     size_t sizeAll = NB_SPHERES * sizeof(Sphere);
  23     HANDLE_ERROR(cudaGetSymbolAddress((void**)&ptrDevTabData, CONST_SPHERES));
  24     ConstantMemoryLink cmLink =
  25         {
  26           (void**)ptrDevTabData, NB_SPHERES, sizeAll
  27         };
  28     return cmLink;
  29     }
  30 #endif
  31
  32 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
  33     __global__
  34     void rayTracing(uchar4* ptrDevPixels, const Sphere* ptrDevSpheres, int w, int h, float t)
  35 #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
  36     __global__
  37     void rayTracing(uchar4* ptrDevPixels, int w, int h, float t)
  38 #endif
  39     {
  40
  41     const int TID = Indice2D::tid();
  42     const int NB_THREAD = Indice2D::nbThread();
  43     const int WH = w * h;
  44
  45 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
  46     __shared__ Sphere tabSpheresSM[NB_SPHERES];
  47
  48     // Copy the global memory to shared memory.
  49     const int NB_THREAD_LOCAL = Indice1D::nbThreadBlock();
  50     const int TID_LOCAL = Indice1D::tidLocal();
  51     int sLocal = TID_LOCAL;
  52     while (sLocal < NB_SPHERES)
  53         {
  54         tabSpheresSM[sLocal] = ptrDevSpheres[sLocal];
  55         sLocal += NB_THREAD_LOCAL;
  56         }
  57
  58     __syncthreads();
  59 #endif
  60
  61     int s = TID;
  62     while (s < WH)
  63         {
  64         const float2 pixel = IndiceTools::toPixelFloat(s, w);
  65
  66         // Les données de la sphère sélectionnée.
  67         float selectedSphereDistance = FLT_MAX;
  68         float brightness = 0;
  69         float hue = 0;
  70
  71         // Sélectionne la sphere la plus proche, il peut ne pas y avoir de sphère.
  72         for (int i = 0; i < NB_SPHERES; ++i)
  73             {
  74 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL
  75             const Sphere& sphere = ptrDevSpheres[i];
  76 #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
  77             const Sphere& sphere = CONST_SPHERES[i];
  78 #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
  79             const Sphere& sphere = tabSpheresSM[i];
  80 #endif
  81             const float hSquare = sphere.hSquare(pixel);
  82
  83             if (sphere.isBelow(hSquare))
  84                 {
  85                 const float dz = sphere.dz(hSquare);
  86                 const float sphereDistance = sphere.distance(dz);
  87
  88                 if (sphereDistance < selectedSphereDistance)
  89                     {
  90                     selectedSphereDistance = sphereDistance;
  91                     brightness = sphere.brightness(dz);
  92                     hue = sphere.hue(t);
  93                     }
  94                 }
  95             }
  96
  97         // Si une sphere a pu être trouvée.
  98         if (selectedSphereDistance != FLT_MAX)
  99             {
 100             ColorTools::HSB_TO_RVB(hue, 1, brightness, &ptrDevPixels[s]);
 101             }
 102
 103         ptrDevPixels[s].w = 255;
 104
 105         s += NB_THREAD;
 106         }
 107     }