From 69f2911e04ffb1b19eef1fafb8c040af271f656e Mon Sep 17 00:00:00 2001 From: Tor Aamodt Date: Thu, 15 Jul 2010 18:09:46 -0800 Subject: creating branch for adding support for CUDA 3.x and Fermi [git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 6829] --- benchmarks/CUDA/RAY/rayTracing_kernel.cu | 658 +++++++++++++++++++++++++++++++ 1 file changed, 658 insertions(+) create mode 100644 benchmarks/CUDA/RAY/rayTracing_kernel.cu (limited to 'benchmarks/CUDA/RAY/rayTracing_kernel.cu') 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 +__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__ -- cgit v1.3