diff options
Diffstat (limited to 'benchmarks/CUDA/RAY/rayTracing_kernel.cu')
| -rw-r--r-- | benchmarks/CUDA/RAY/rayTracing_kernel.cu | 658 |
1 files changed, 658 insertions, 0 deletions
diff --git a/benchmarks/CUDA/RAY/rayTracing_kernel.cu b/benchmarks/CUDA/RAY/rayTracing_kernel.cu new file mode 100644 index 0000000..c524a9c --- /dev/null +++ b/benchmarks/CUDA/RAY/rayTracing_kernel.cu @@ -0,0 +1,658 @@ +/*
+ * Copyright 2008 BOROUJERDI Maxime. Tous droits reserves.
+ */
+
+//#define FIXED_CONST_PARSE
+#ifndef __RAYTRACING_KERNEL_H__
+#define __RAYTRACING_KERNEL_H__
+
+#include "cutil_math.h"
+
+typedef struct
+{
+ float4 m[3];
+} matrice3x4;
+
+typedef struct {
+ float4 m[4];
+} matrice4x4;
+
+typedef struct{
+ float3 A; // origine
+ float3 u; // direction
+} Rayon;
+
+typedef struct Sphere{
+ float3 C; // centre
+ float r; // rayon
+ float R,V,B,A;
+ /*Sphere() : C(make_float3(0.0f,0.0f,0.0f)), r(0.5f), rvba(make_float4(1.0f,0.0f,0.0f,1.0f)) { }
+ Sphere(const float3 _C, float _r, const float4 _rvba) : C(_C), r(_r), rvba(_rvba) { }
+ Sphere(const float3 _C, float _r) : C(_C), r(_r), rvba(make_float4(1.0f,0.0f,0.0f,1.0f)) { }*/
+} Sphere;
+
+typedef struct Node {
+ Sphere s;
+ uint fg, fd;
+} Node;
+
+/*__host__ __device__ void createNode(Node * n, Node * fg, Node * fd, const Sphere & s)
+{
+ n->fg = fg;
+ n->fd = fd;
+ n->C = s.C;
+ n->r = s.r;
+}
+
+__host__ __device__ Node * filsGauche(Node * n) { return n->fg; }
+
+__host__ __device__ Node * filsDroite(Node * n) { return n->fd; }*/
+
+//__host__ __device__ Sphere sphere(Node * n) { return n->s; }
+
+__constant__ matrice3x4 MView; // matrice inverse de la matrice de vue
+
+__constant__ Node cnode[numObj];
+
+template <class T>
+__device__ void swap(T & v1, T & v2)
+{
+ T tmp(v1);
+ v1 = v2;
+ v2 = tmp;
+}
+
+__device__ float intersectionSphere(Rayon R, float3 C, float r)
+{
+ float3 L(C-R.A);
+ float d(dot(L,R.u)), l2(dot(L,L)), r2(r*r), m2, q, res;
+
+ if( d < 0.0f && l2 > r2 ) {
+ res = 0.0f;
+ }
+ else
+ {
+ m2 = l2 - d*d;
+ if( m2 > r2 ) {
+ res = 0.0f;
+ }
+ else
+ {
+ q = sqrt(r2-m2);
+ if( l2 > r2 ) res = d - q;
+ else res = d + q;
+ }
+ }
+
+ return res;
+}
+
+__device__ float intersectionPlan( Rayon R, float3 C, float3 N2 )
+{
+ float res;
+ float3 N = normalize(make_float3(0.0f,1.0f,0.0f));
+ float m(dot(N,R.u)), d, t;
+ float3 L;
+
+ if( fabs(m) < 0.0001f ) {
+ res = 0.0f;
+ }
+ else {
+ L = R.A - C;
+ d = dot(N,L);
+ t = -d/m;
+ if( t > 0 ) {
+ res = t;
+ }
+ else {
+ res = 0.0f;
+ }
+ }
+
+ return res;
+}
+
+__device__ float3 getNormale(float3 P, float3 C)
+{
+ return normalize(P-C);
+}
+
+__device__ float3 getNormaleP(float3 P)
+{
+ return normalize(make_float3(0.0f,1.0f,0.0f));
+}
+
+// multiplication d'un vecteur par une matrice (sans translation)
+__device__ float3 mul(matrice3x4 M, float3 v)
+{
+ float3 r;
+ r.x = dot(v, make_float3(M.m[0]));
+ r.y = dot(v, make_float3(M.m[1]));
+ r.z = dot(v, make_float3(M.m[2]));
+ return r;
+}
+
+// multiplication d'un vecteur par une matrice avec translation
+__device__ float4 mul(matrice3x4 M, float4 v)
+{
+ float4 r;
+ r.x = dot(v, M.m[0]);
+ r.y = dot(v, M.m[1]);
+ r.z = dot(v, M.m[2]);
+ r.w = 1.0f;
+ return r;
+}
+
+__device__ uint rgbaFloatToInt(float4 rgba)
+{
+ #ifdef DEVICE_EMU
+ printf("%d: rgba = %f %f %f %f\n", threadIdx.x, rgba.x, rgba.y, rgba.z, rgba.w);
+ #endif
+ rgba.x = __saturatef(rgba.x); // clamp entre [0.0, 1.0]
+ rgba.y = __saturatef(rgba.y);
+ rgba.z = __saturatef(rgba.z);
+ rgba.w = __saturatef(rgba.w);
+#ifdef DEVICE_EMU
+ printf("%d: rgba = %x %x %x %x\n", threadIdx.x, uint(rgba.x*255), uint(rgba.y*255), uint(rgba.z*255), uint(rgba.w*255));
+#endif
+ return (uint(rgba.w*255)<<24)
+ | (uint(rgba.z*255)<<16)
+ | (uint(rgba.y*255)<<8 )
+ | (uint(rgba.x*255) );
+}
+
+/*__device__ void myswap(Sphere &x, Sphere &y)
+{
+Sphere t = x;
+x = y;
+y = t;
+}*/
+/*__global__ void d_render(uint * d_output, uint imageW, uint imageH, float pas, float df, float tPixel)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ //float tPixel = 2.0f/(float)min(imageW,imageH);
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ Sphere s(cnode[1].s), s2(cnode[2].s), st(cnode[2].s);
+ float t, t2, tt;
+ s.C.x += pas, s2.C.x += pas;
+ t = intersectionSphere(R,s.C,s.r);
+ t2 = intersectionSphere(R,s2.C,s2.r);
+ if( !t ) {
+ //myswap(s,s2);
+ //swap(t,t2);
+ tt = t;
+ t = t2;
+ t2 = tt;
+ st = s;
+ s = s2;
+ s2 = st;
+ }
+ else if( t2 && t2 < t ) {
+ //myswap(s,s2);
+ //swap(t,t2);
+ tt = t;
+ t = t2;
+ t2 = tt;
+ st = s;
+ s = s2;
+ s2 = st;
+ }
+ float4 f = make_float4(0,1,0,1)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ uint n = rgbaFloatToInt(f);
+ //printf("%f\n",d_node[0].s.r);
+ if( t > 0.0f )
+ d_output[id] = n;
+ //else d_output[id] = 0;
+ }
+ __syncthreads();
+}
+*/
+/*__global__ void rayCast (uint * d_output, uint * d_temp, uint imageW, uint imageH, float pas, float df)
+//(uint * result, uint * temp, uint imageW, uint imageH, float pas, float df)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y);
+ uint id = x + y * gridDim.x;
+ //float tmp= float(imageW)/float(gridDim.x);
+ float t;
+
+ //if( x < gridDim.x && y < gridDim.y )
+ if( d_temp[id] == 0 )
+ {
+ float tPixel = 2.0f/float(imageW);
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ Sphere s(cnode[1].s);
+ s.C.x += pas;
+ t = intersectionSphere(R,s.C,s.r/(imageW/gridDim.x));
+
+ if( t > 0.0f ) {
+ //float4 f = make_float4(0,1,0,1)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ d_output[id] = rgbaFloatToInt(make_float4(0,1,0,1));
+ //printf("%d %d\n",int(x*tmp),int((y*tmp)/2));
+ }
+ else {
+// float tmp= float(imageW)/gridDim.x;
+// d_temp[int(x*tmp+(y*tmp)*imageW)] = 1;
+// d_temp[int(x*tmp+(tmp*(float(y)+0.5f)*imageW))] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(y*tmp)*imageW)] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(tmp*(float(y)+0.5f)*imageW))] = 1;
+ //if(gridDim.x==16) printf("hep %d %f\n",gridDim.x,t);
+ }
+ }
+ else {
+// float tmp= float(imageW)/gridDim.x;
+// d_temp[int(x*tmp+(y*tmp)*imageW)] = 1;
+// d_temp[int(x*tmp+(tmp*(float(y)+0.5f)*imageW))] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(y*tmp)*imageW)] = 1;
+// d_temp[int(tmp*(float(x)+0.5f)+(tmp*(float(y)+0.5f)*imageW))] = 1;
+ //if(gridDim.x==16) printf("hep %d %f\n",gridDim.x,t);
+ }
+ //__syncthreads();
+}*/
+
+/*__global__ __device__ void rayCalc(float3 * A, float3 * u, float * prof, uint imageW, uint imageH, float df, float tPixel)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+ A[id] = R.A;
+ u[id] = R.u;
+ prof[id] = 1000.0f;
+ }
+}*/
+
+
+/*__global__ __device__ void rayTrace(uint * Obj, float * prof, float3 * A, float3 * u, uint imageW, uint imageH, float pas, float df, uint nObj)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ Sphere s(cnode[nObj].s);
+ float t;
+ s.C.x += pas;
+ Rayon R;
+ R.A = A[id];
+ R.u = u[id];
+ t = intersectionSphere(R,s.C,s.r);
+
+ if( t > 0.0f && t < prof[id] ) {
+ prof[id] = t;
+ Obj[id] = nObj;
+ }
+ }
+}*/
+/*
+__global__ __device__ void color(uint * result, uint * Obj, float * prof, float3 * A, float3 * u, uint imageW, uint imageH, float pas)
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint id = x + y * imageW;
+
+ if( x < imageW && y < imageH )
+ {
+ float t(prof[id]);
+ if( t > 0.0f && t < 1000.0f ) {
+ Rayon R;
+ R.A = A[id];
+ R.u = u[id];
+ Sphere s(cnode[Obj[id]].s);
+ s.C.x += pas;
+ float4 f = make_float4(s.R,s.V,s.B,s.A)*(dot(getNormale(R.A+R.u*t,s.C),(-1.0f)*R.u));
+ result[id] = rgbaFloatToInt(f);
+ }
+ else {
+ result[id] = 0;
+ }
+ prof[id] = 100000.0f;
+ }
+}*/
+#ifdef DEBUG_RT_CUDA
+__device__ bool notShadowRay( float4* d_debug_float4, uint* d_debug_uint, int i, Node * node, float3 A, float3 u, float pas ) {
+
+#else
+__device__ bool notShadowRay( Node * node, float3 A, float3 u, float pas ) {
+#endif
+ float t(0.0f);
+ Node n;
+ Rayon ray;
+ float3 L(make_float3(10.0f,10.0f,10.0f)), tmp;
+ float dst(dot(tmp=(L-A),tmp));
+ ray.A = A+u*0.0001f;
+ ray.u = u;
+ for( int j(0); j < numObj && !t; j++ ) {
+ n = cnode[j];
+ n.s.C.x += pas;
+ if( n.fg ){
+ t = intersectionPlan(ray,n.s.C,n.s.C);
+ #ifdef DEVICE_EMU
+// printf("%d: j=%d, intersectionPlan t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_uint4[threadIdx.x*16+4*j+0]=10;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+0].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+0].y = 99999.9f;
+ #endif
+ }
+ else{
+#ifdef DEVICE_EMU
+ printf("%d: i=%d, n.s.C = %e %e %e\n", threadIdx.x, i, n.s.C.x, n.s.C.y, n.s.C.z);
+ printf("%d: i=%d, n.s.r = %e\n", threadIdx.x, i, n.s.r);
+#endif
+ #ifdef DEBUG_RT_CUDA
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].x = n.s.C.x;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].y = n.s.C.y;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].z = n.s.C.z;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+0].w = n.s.r;
+ #endif
+ t = intersectionSphere(ray,n.s.C,n.s.r);
+ #ifdef DEVICE_EMU
+ printf("%d: j=%d, intersectionSphere t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ d_debug_float4[threadIdx.x*32+16*i+3*j+1].x = t;
+ d_debug_float4[threadIdx.x*32+16*i+3*j+1].y = 99999.9f;
+ #endif
+ }
+ if( t > 0.0f && dot(tmp=(A+u*t),tmp) > dst ){
+ t = 0.0f;
+ #ifdef DEVICE_EMU
+// printf("%d: j=%d, && dot t=%e\n", threadIdx.x, j, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*32+16*i+3*j+2].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+3*j+2].y = 99999.9f;
+ #endif
+ }
+ }
+ #ifdef DEVICE_EMU
+// printf("%d: t=%e\n", threadIdx.x, t);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*32+16*i+13].x = t;
+// d_debug_float4[threadIdx.x*32+16*i+13].y = 99999.9f;
+ d_debug_float4[threadIdx.x*32+16*i+15].x = 88888.8f;
+ #endif
+ return t == 0.0f;
+}
+
+__device__ float float2int_pow20(float a)
+{
+ return a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a;
+}
+
+__device__ float float2int_pow50(float a)
+{
+ return a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a* \
+ a*a*a*a*a* a*a*a*a*a;
+
+}
+#ifdef DEBUG_RT_CUDA
+__global__ __device__ void render(float4* d_debug_float4, uint* d_debug_uint, uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+#else
+__global__ __device__ void render(uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+#endif
+{
+ uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
+ uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
+ uint tid(__umul24(threadIdx.y, blockDim.x) + threadIdx.x);
+
+ uint id(x + y * imageW);
+ float4 pile[5];
+ uint Obj, nRec(5), n(0);
+ //__shared__ Node node[numObj];
+ float prof, tmp;
+
+ //if( tid < numObj ) node[tid] = cnode[tid];
+
+ for( int i(0); i < nRec; ++i )
+ pile[i] = make_float4(0.0f,0.0f,0.0f,1.0f);
+
+ if( x < imageW && y < imageH )
+ {
+ prof = 10000.0f;
+ result[id] = 0;
+ float tPixel(2.0f/float(min(imageW,imageH)));
+ float4 f(make_float4(0.0f,0.0f,0.0f,1.0f));
+ matrice3x4 M(MView);
+ Rayon R;
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+#ifdef DEVICE_EMU
+// printf("%d: R.A = %e %e %e\n", threadIdx.x, R.A.x, R.A.y, R.A.z);
+// printf("%d: R.u = %e %e %e\n", threadIdx.x, R.u.x, R.u.y, R.u.z);
+#endif
+#ifdef DEBUG_RT_CUDA
+// d_debug_float4[threadIdx.x*2+0].x= R.A.x;
+// d_debug_float4[threadIdx.x*2+0].y= R.A.y;
+// d_debug_float4[threadIdx.x*2+0].z= R.A.z;
+// d_debug_float4[threadIdx.x*2+1].x= R.u.x;
+// d_debug_float4[threadIdx.x*2+1].y= R.u.y;
+// d_debug_float4[threadIdx.x*2+1].z= R.u.z;
+#endif
+ __syncthreads();
+
+ for( int i(0); i < nRec && n == i; i++ ) {
+
+ for( int j(0); j < numObj; j++ ) {
+ Node nod(cnode[j]);
+ Sphere s(nod.s);
+ float t;
+ s.C.x += pas;
+ if( nod.fg )
+ t = intersectionPlan(R,s.C,s.C);
+ else
+ t = intersectionSphere(R,s.C,s.r);
+
+ if( t > 0.0f && t < prof ) {
+ prof = t;
+ Obj = j;
+ }
+ }
+#ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*5+i].x= prof;
+#endif
+#ifdef DEVICE_EMU
+// printf("%d: i=%d, t=%e\n", threadIdx.x, i, prof);
+#endif
+ float t = prof;
+ if( t > 0.0f && t < 10000.0f ) {
+ n++;
+ Node nod(cnode[Obj]);
+ Sphere s(nod.s);
+ s.C.x += pas;
+ float4 color(make_float4(s.R,s.V,s.B,s.A));
+ float3 P(R.A+R.u*t), L(normalize(make_float3(10.0f,10.0f,10.0f)-P)), V(normalize(R.A-P));
+ float3 N(nod.fg?getNormaleP(P):getNormale(P,s.C));
+ float3 Np(dot(V,N)<0.0f?(-1*N):N);
+ pile[i] = 0.05f * color;
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+// printf("%d: i=%d, color = %e %e %e %e\n", threadIdx.x, i, color.x, color.y, color.z, color.w);
+// printf("%d: i=%d, P = %e %e %e\n", threadIdx.x, i, P.x, P.y, P.z);
+// printf("%d: i=%d, L = %e %e %e\n", threadIdx.x, i, L.x, L.y, L.z);
+// printf("%d: i=%d, V = %e %e %e\n", threadIdx.x, i, V.x, V.y, V.z);
+// printf("%d: i=%d, N = %e %e %e\n", threadIdx.x, i, N.x, N.y, N.z);
+// printf("%d: i=%d, Np = %e %e %e\n", threadIdx.x, i, Np.x, Np.y, Np.z);
+// printf("%d: i=%d, dot(Np,L) = %e\n", threadIdx.x, i, dot(Np,L));
+ //printf("%d: i=%d, notShadowRay(cnode,P,L,pas) = %d\n", threadIdx.x, i, (int) notShadowRay(cnode,P,L,pas));
+
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+0]= pile[i];
+// d_debug_float4[threadIdx.x*16+i*8+0]= color;
+// d_debug_float4[threadIdx.x*16+i*8+1].x= P.x;d_debug_float4[threadIdx.x*16+i*8+1].y= P.y;d_debug_float4[threadIdx.x*16+i*8+1].z= P.z;
+// d_debug_float4[threadIdx.x*16+i*8+2].x= L.x;d_debug_float4[threadIdx.x*16+i*8+2].y= L.y;d_debug_float4[threadIdx.x*16+i*8+2].z= L.z;
+// d_debug_float4[threadIdx.x*16+i*8+3].x= V.x;d_debug_float4[threadIdx.x*16+i*8+3].y= V.y;d_debug_float4[threadIdx.x*16+i*8+3].z= V.z;
+// d_debug_float4[threadIdx.x*16+i*8+4].x= N.x;d_debug_float4[threadIdx.x*16+i*8+4].y= N.y;d_debug_float4[threadIdx.x*16+i*8+4].z= N.z;
+// d_debug_float4[threadIdx.x*16+i*8+5].x= Np.x;d_debug_float4[threadIdx.x*16+i*8+5].y= Np.y;d_debug_float4[threadIdx.x*16+i*8+5].z= Np.z;
+// d_debug_float4[threadIdx.x*16+i*8+6].x= dot(Np,L);
+ //d_debug_float4[threadIdx.x*16+i*8+7].x= (float) notShadowRay(cnode,P,L,pas);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ if( dot(Np,L) > 0.0f && notShadowRay(d_debug_float4, d_debug_uint, i, cnode,P,L,pas) ) {
+ #else
+ if( dot(Np,L) > 0.0f && notShadowRay(cnode,P,L,pas) ) {
+ #endif
+ //float3 Ri(2.0f*Np*dot(Np,L) - L);
+ float3 Ri(normalize(L+V));
+ //Ri = (L+V)/normalize(L+V);
+ pile[i] += 0.3f * color* (min(1.0f,dot(Np,L)));
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+1]= pile[i];
+ #endif
+ #ifdef FIXED_CONST_PARSE
+ tmp = 0.8f * pow(max(0.0f,min(1.0f,dot(Np,Ri))),50.0f);
+ #else
+ tmp = 0.8f * float2int_pow50(max(0.0f,min(1.0f,dot(Np,Ri))));
+ #endif
+ pile[i].x += tmp;
+ pile[i].y += tmp;
+ pile[i].z += tmp;
+ #ifdef DEVICE_EMU
+// printf("%d: i=%d, pile[i] = %e %e %e %e\n", threadIdx.x, i, pile[i].x, pile[i].y, pile[i].z, pile[i].w);
+ #endif
+ #ifdef DEBUG_RT_CUDA
+ //d_debug_float4[threadIdx.x*16+i*3+2]= pile[i];
+ #endif
+ }
+
+ R.u = 2.0f*N*dot(N,V) - V;
+ R.u = normalize(R.u);
+ R.A = P+R.u*0.0001f;
+ }
+ prof = 10000.0f;
+ }
+ #ifdef DEBUG_RT_CUDA
+ /*d_debug_float4[threadIdx.x*5+0]= pile[0];
+ d_debug_float4[threadIdx.x*5+1]= pile[1];
+ d_debug_float4[threadIdx.x*5+2]= pile[2];
+ d_debug_float4[threadIdx.x*5+3]= pile[3];
+ d_debug_float4[threadIdx.x*5+4]= pile[4];*/
+ #endif
+#ifdef DEVICE_EMU
+// printf("%d: pile[0] = %e %e %e %e\n", threadIdx.x, pile[0].x, pile[0].y, pile[0].z, pile[0].w);
+// printf("%d: pile[1] = %e %e %e %e\n", threadIdx.x, pile[1].x, pile[1].y, pile[1].z, pile[1].w);
+// printf("%d: pile[2] = %e %e %e %e\n", threadIdx.x, pile[2].x, pile[2].y, pile[2].z, pile[2].w);
+// printf("%d: pile[3] = %e %e %e %e\n", threadIdx.x, pile[3].x, pile[3].y, pile[3].z, pile[3].w);
+// printf("%d: pile[4] = %e %e %e %e\n", threadIdx.x, pile[4].x, pile[4].y, pile[4].z, pile[4].w);
+#endif
+ for( int i(n-1); i > 0; i-- )
+ pile[i-1] = pile[i-1] + 0.8f*pile[i];
+#ifdef DEVICE_EMU
+// printf("%d: pile[0] = %e %e %e %e\n", threadIdx.x, pile[0].x, pile[0].y, pile[0].z, pile[0].w);
+#endif
+ result[id] += rgbaFloatToInt(pile[0]);
+ }
+}
+
+/*__global__ __device__ void renderPixel(uint * result, Node * dnode, uint imageW, uint imageH, float pas, float df)
+{
+ uint id(blockIdx.x + __umul24(blockIdx.y, imageW));
+ uint tid(threadIdx.x), x(blockIdx.x), y(blockIdx.y);
+ Node node;
+ float t(0.0f), tPixel;
+ float4 Color(make_float4(0.0f,0.0f,0.0f,1.0f));
+ matrice3x4 M(MView);
+ Rayon R;
+ Sphere s;
+ __shared__ float T[numObj];
+ __shared__ uint Obj;
+
+ T[tid] = 10000.0f;
+
+ if( x < imageW && y < imageH && tid < numObj ) {
+ node = dnode[tid];
+ if( tid == 0 ) result[id] = 0;
+ tPixel = 2.0f/float(min(imageW,imageH));
+ R.A = make_float3(M.m[0].w,M.m[1].w,M.m[2].w);
+ R.u = make_float3(M.m[0])*df
+ + make_float3(M.m[2])*(float(x)-float(imageW)*0.5f)*tPixel
+ + make_float3(M.m[1])*(float(y)-float(imageH)*0.5f)*tPixel;
+ R.u = normalize(R.u);
+
+ s = node.s;
+ s.C.x += pas;
+
+ if( node.fg )
+ t = intersectionPlan(R,s.C,s.C);
+ else
+ t = intersectionSphere(R,s.C,s.r);
+
+ T[tid] = t;
+
+ __syncthreads();
+
+ if( tid == 0 ) {
+ float tmp(t);
+ Obj = 0;
+ for( int i(1); i < numObj; i++ ) {
+ if( T[i] > 0.0f && ( tmp == 0.0f || T[i] < tmp ) ) {
+ tmp = T[i];
+ Obj = i;
+ }
+ }
+ }
+
+ __syncthreads();
+
+ if( tid == Obj && t > 0.0f ) {
+ s = node.s;
+ s.C.x += pas;
+ float3 P(R.A+R.u*t), L(normalize(make_float3(0,1,2)-P)), V(-1*R.u);
+ float3 N(node.fg?getNormaleP(P):getNormale(P,s.C));
+ if( dot(N,L) > 0.0f ) {
+ Color = 0.5f*make_float4(s.R,s.V,s.B,s.A)*(max(0.0f,dot(N,L)));
+ #ifdef FIXED_CONST_PARSE
+ Color += 0.8f*make_float4(1.0f,1.0f,1.0f,1.0f)*pow(max(0.0f,min(1.0f,dot(2.0f*N*dot(N,L)-L,V))),20.0f);
+ #else
+ Color += 0.8f*make_float4(1.0f,1.0f,1.0f,1.0f)*float2int_pow20(max(0.0f,min(1.0f,dot(2.0f*N*dot(N,L)-L,V))));
+ #endif
+ }
+ result[id] = rgbaFloatToInt(Color);
+ }
+ }
+
+}
+*/
+
+#endif // __RAYTRACING_KERNEL_H__
|