From ebdad7dc732d4742d09fd72d312175cb07d27a59 Mon Sep 17 00:00:00 2001 From: gburri Date: Thu, 18 Dec 2014 08:53:05 +0100 Subject: [PATCH] Ray tracing (pas termine). --- .../moo/device/RayTracingDevice.cu | 43 +++++++++ .../moo/device/RayTracingDevice.h | 9 ++ .../core/04_RayTracing/moo/device/Sphere.h | 96 +++++++++++++++++++ .../moo/device/math/RayTracingMath.h | 9 ++ .../core/04_RayTracing/moo/host/RayTracing.cu | 79 +++++++++++++++ .../core/04_RayTracing/moo/host/RayTracing.h | 39 ++++++++ .../provider/RayTracingProvider.cpp | 15 +++ .../provider/RayTracingProvider.h | 14 +++ 8 files changed, 304 insertions(+) create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/math/RayTracingMath.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp create mode 100644 WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h 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 new file mode 100644 index 0000000..368e469 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.cu @@ -0,0 +1,43 @@ +#include +#include +using namespace std; + +#include "Indice2D.h" +#include "IndiceTools.h" +#include "cudaTools.h" +#include "Device.h" + +#include "RayTracingDevice.h" +#include "RayTracingMath.h" + +__global__ +void rayTracing(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, float t) + { + const int TID = Indice2D::tid(); + const int NB_THREAD = Indice2D::nbThread(); + const int WH = w * h; + + RayTracingMath rayTracing; // TODO: prendre + + uchar4 color; + color.w = 255; // Par défaut, l'image est opaque. + + double x, y; + int pixelI, pixelJ; + + int s = TID; + while (s < WH) + { + IndiceTools::toIJ(s, w, &pixelI, &pixelJ); // update (pixelI, pixelJ) + + // (i,j) domaine écran. + // (x,y) domaine math. + // domaineMath.toXY(pixelI, pixelJ, &x, &y); // (i,j) -> (x,y). + + // newtonMath.colorXY(&color, x, y); + + ptrDevPixels[s] = color; + + s += NB_THREAD; + } + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h new file mode 100644 index 0000000..47e93ae --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/RayTracingDevice.h @@ -0,0 +1,9 @@ +#ifndef RAY_TRACING_DEVICE_H +#define RAY_TRACING_DEVICE_H + +#include "DomaineMath.h" + +__global__ +void rayTracing(uchar4* ptrDevPixels, int w, int h, DomaineMath domaineMath, float t); + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h new file mode 100644 index 0000000..566b4ab --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/Sphere.h @@ -0,0 +1,96 @@ +#ifndef SPHERE_H +#define SPHERE_H + +#include "cudaTools.h" +#include "cudaType_CPU.h" +#include "mathTools.h" + +#endif +class Sphere + { + public: + __host__ + Sphere(float3 centre, float r, float hue) : + centre(centre), + hueInitial(hue) + { + this->setR(r); + } + + __host__ + Sphere() + { + } + + __host__ + void setCentre(float3 centre) + { + this->centre = centre; + } + + __host__ + void setR(float r) + { + this->r = r; + this->rCarre = r*r; + } + + __host__ + void setHueInitial(float hue) + { + this->hueInitial = hue; + } + + __device__ + float hCarre(float2 xySol) + { + float a = (centre.x - xySol.x); + float b = (centre.y - xySol.y); + return a * a + b * b; + } + + __device__ + bool isEnDessous(float hCarre) + { + return hCarre < rCarre; + } + + __device__ + float dz(float hCarre) + { + return sqrtf(rCarre - hCarre); + } + + __device__ + float brightness(float dz) + { + return dz / r; + } + + __device__ + float distance(float dz) + { + return centre.z - dz; + } + + __device__ + float getHueInitial() + { + return this->hueInitial; + } + + __device__ + float hue(float t) // usefull for animation + { + return 0.5 + 0.5 * sin(t + T + 3 * PI / 2); + } + + private: + float r; + float3 centre; + float hueInitial; + float rCarre; + + float T; // Utilisé pour l'animation. + }; +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/math/RayTracingMath.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/math/RayTracingMath.h new file mode 100644 index 0000000..91b32da --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/device/math/RayTracingMath.h @@ -0,0 +1,9 @@ +#ifndef RAY_TRACING_MATH_H +#define RAY_TRACING_MATH_H + +class RayTracingMath + { + + }; + +#endif 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 new file mode 100644 index 0000000..36e3862 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.cu @@ -0,0 +1,79 @@ +#include +#include +#include +using namespace std; + +#include "AleaTools.h" + +#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") + { + Device::assertDim(dg, db); + + const int nbSpheres = 10; + Sphere* = this->createSpheres(10); + + + // Copie les spheres dans la constant memory. + } + +RayTracing::~RayTracing() + { + delete this->ptrDomaineMathInit; + } + +void RayTracing::runGPU(uchar4* ptrDevPixels, const DomaineMath& domaineMath) + { + rayTracing<<>>(ptrDevPixels, this->w, this->h, domaineMath, this->t); + + HANDLE_ERROR(cudaDeviceSynchronize()); // Pour flusher les 'printf' (pour le DEBUG). + } + +void RayTracing::animationStep() + { + this->t += 0.1; // TODO + } + +int RayTracing::getW() + { + return this->w; + } + +int RayTracing::getH() + { + return this->h; + } + +float RayTracing::getT() + { + return this->t; + } + +string RayTracing::getTitle() + { + return this->title; + } + +Sphere* createSpheres(int n) + { + Sphere* spheres = new Sphere[n]; + const float bord = 200; + + 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)) + } + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h new file mode 100644 index 0000000..80729b8 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/moo/host/RayTracing.h @@ -0,0 +1,39 @@ +#ifndef RAY_TRACING_H +#define RAY_TRACING_H + +#include "cudaTools.h" +#include "Animable_I.h" +#include "MathTools.h" +#include "Sphere.h" + +class RayTracing : public Animable_I + { + public: + RayTracing(int w, int h); + ~RayTracing(); + + void runGPU(uchar4* ptrDevPixels) /*override*/; + void animationStep() /*override*/; + + int getW() /*override*/; + int getH() /*override*/; + + float getT() /*override*/; + + string getTitle(void) /*override*/; + + private: + Sphere* createSpheres(int n); + + float t; + + const int w; + const int h; + + const dim3 dg; + const dim3 db; + + const string title; + }; + +#endif diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp new file mode 100644 index 0000000..0486785 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.cpp @@ -0,0 +1,15 @@ +#include "RayTracingProvider.h" + +RayTracing* RayTracingProvider::create() + { + int dw = 16 * 50; + int dh = 16 * 50; + + return new RayTracing(dw, dh); + } + +ImageFonctionel* RayTracingProvider::createGL() + { + ColorRGB_01* ptrColorTitre = new ColorRGB_01(0, 0, 0); + return new ImageFonctionel(create(), ptrColorTitre); // both ptr destroy by destructor of ImageFonctionel + } diff --git a/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h new file mode 100644 index 0000000..f78ffc5 --- /dev/null +++ b/WCudaMSE/Student_Cuda_Image/src/cpp/core/04_RayTracing/provider/RayTracingProvider.h @@ -0,0 +1,14 @@ +#ifndef RAYTRACING_PROVIDER_H +#define RAYTRACING_PROVIDER_H + +#include "RayTracing.h" +#include "ImageFonctionel.h" + +class RayTracingProvider + { + public: + static RayTracing* create(); + static ImageFonctionel* createGL(); + }; + +#endif -- 2.43.0