8 #include "IndiceTools.h"
9 #include "ColorTools.h"
10 #include "cudaTools.h"
13 #include "RayTracingDevice.h"
15 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
16 // Les spheres dans la constant memory.
17 __constant__ Sphere CONST_SPHERES[NB_SPHERES];
19 ConstantMemoryLink constantMemorySpheresLink()
22 size_t sizeAll = NB_SPHERES * sizeof(Sphere);
23 HANDLE_ERROR(cudaGetSymbolAddress((void**)&ptrDevTabData, CONST_SPHERES));
24 ConstantMemoryLink cmLink =
26 (void**)ptrDevTabData, NB_SPHERES, sizeAll
32 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_GLOBAL or CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
34 void rayTracing(uchar4* ptrDevPixels, const Sphere* ptrDevSpheres, int w, int h, float t)
35 #elif CURRENT_MEMORY_MODEL == MEMORY_MODEL_CONSTANT
37 void rayTracing(uchar4* ptrDevPixels, int w, int h, float t)
41 const int TID = Indice2D::tid();
42 const int NB_THREAD = Indice2D::nbThread();
45 #if CURRENT_MEMORY_MODEL == MEMORY_MODEL_SHARED
46 __shared__ Sphere tabSpheresSM[NB_SPHERES];
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)
54 tabSpheresSM[sLocal] = ptrDevSpheres[sLocal];
55 sLocal += NB_THREAD_LOCAL;
64 const float2 pixel = IndiceTools::toPixelFloat(s, w);
66 // Les données de la sphère sélectionnée.
67 float selectedSphereDistance = FLT_MAX;
71 // Sélectionne la sphère la plus proche, il peut ne pas y avoir de sphère trouvée.
72 for (int i = 0; i < NB_SPHERES; ++i)
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];
82 const float hSquare = sphere.hSquare(pixel);
84 if (sphere.isBelow(hSquare))
86 const float dz = sphere.dz(hSquare);
87 const float sphereDistance = sphere.distance(dz);
89 if (sphereDistance < selectedSphereDistance)
91 selectedSphereDistance = sphereDistance;
92 brightness = sphere.brightness(dz);
98 // Si une sphère a pu être trouvée.
99 if (selectedSphereDistance != FLT_MAX)
101 ColorTools::HSB_TO_RVB(hue, 1, brightness, &ptrDevPixels[s]);
104 ptrDevPixels[s].w = 255;