From 11b308e7363e937966b035b4891db32b4eece3bf Mon Sep 17 00:00:00 2001
From: Tor Aamodt <aamodt@ece.ubc.ca>
Date: Fri, 1 Oct 2010 08:55:28 -0800
Subject: integrating recent changes from fermi-test into fermi (i'll use
 "fermi" for more disruptive changes to the pipeline model such as updating
 the MSHRs and getting rid of the warp tracker, ripping out DWF, etc...)

[git-p4: depot-paths = "//depot/gpgpu_sim_research/fermi/distribution/": change = 7805]
---
 benchmarks/CUDA/RAY/rayTracing_kernel.cu | 658 -------------------------------
 1 file changed, 658 deletions(-)
 delete 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
deleted file mode 100644
index c524a9c..0000000
--- a/benchmarks/CUDA/RAY/rayTracing_kernel.cu
+++ /dev/null
@@ -1,658 +0,0 @@
-/*
- * 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__
-- 
cgit v1.3