Move CUDA source inside chroma package, rename tests directory to test

1 files changed, 0 insertions, 171 deletions

diff --git a/src/render.cu b/src/render.cu
deleted file mode 100644
index d4ed443..0000000
--- a/src/render.cu
+++ /dev/null
@@ -1,171 +0,0 @@
-//-*-c-*-
-
-#include "linalg.h"
-#include "intersect.h"
-#include "mesh.h"
-#include "sorting.h"
-#include "geometry.h"
-
-#include "stdio.h"
-
-__device__ float4
-get_color(const float3 &direction, const Triangle &t, unsigned int rgba)
-{
-    float3 v01 = t.v1 - t.v0;
-    float3 v12 = t.v2 - t.v1;
-    
-    float3 surface_normal = normalize(cross(v01,v12));
-
-    float cos_theta = dot(surface_normal,-direction);
-
-    if (cos_theta < 0.0f)
-	cos_theta = -cos_theta;
-
-    unsigned int a0 = 0xff & (rgba >> 24);
-    unsigned int r0 = 0xff & (rgba >> 16);
-    unsigned int g0 = 0xff & (rgba >> 8);
-    unsigned int b0 = 0xff & rgba;
-
-    float alpha = (255 - a0)/255.0f;
-
-    return make_float4(r0*cos_theta, g0*cos_theta, b0*cos_theta, alpha);
-}
-
-extern "C"
-{
-
-__global__ void
-render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g,
-       unsigned int alpha_depth, unsigned int *pixels, float *_dx,
-       unsigned int *dxlen, float4 *_color)
-{
-    __shared__ Geometry sg;
-
-    if (threadIdx.x == 0)
-	sg = *g;
-
-    __syncthreads();
-
-    int id = blockIdx.x*blockDim.x + threadIdx.x;
-	
-    if (id >= nthreads)
-	return;
-
-    g = &sg;
-
-    float3 origin = _origin[id];
-    float3 direction = _direction[id];
-    unsigned int n = dxlen[id];
-
-    float distance;
-
-    if (n < 1 && !intersect_node(origin, direction, g, g->start_node)) {
-	pixels[id] = 0;
-	return;
-    }
-
-    unsigned int stack[STACK_SIZE];
-
-    unsigned int *head = &stack[0];
-    unsigned int *node = &stack[1];
-    unsigned int *tail = &stack[STACK_SIZE-1];
-    *node = g->start_node;
-
-    float *dx = _dx + id*alpha_depth;
-    float4 *color_a = _color + id*alpha_depth;
-
-    unsigned int i;
-
-    do {
-	unsigned int first_child = g->node_map[*node];
-	unsigned int stop = g->node_map_end[*node];
-	
-	while (*node >= g->first_node && stop == first_child+1) {
-	    *node = first_child;
-	    first_child = g->node_map[*node];
-	    stop = g->node_map_end[*node];
-	}
-		
-	if (*node >= g->first_node) {
-	    for (i=first_child; i < stop; i++) {
-		if (intersect_node(origin, direction, g, i)) {
-		    *node = i;
-		    node++;
-		}
-	    }
-
-	    node--;
-	}
-	else {
-	    // node is a leaf
-	    for (i=first_child; i < stop; i++) {
-		Triangle t = get_triangle(g, i);
-		
-		if (intersect_triangle(origin, direction, t, distance)) {
-		    if (n < 1) {
-			dx[0] = distance;
-			    
-			unsigned int rgba = g->colors[i];
-			float4 color = get_color(direction, t, rgba);
-
-			color_a[0] = color;
-		    }
-		    else {
-			unsigned long j = searchsorted(n, dx, distance);
-
-			if (j <= alpha_depth-1) {
-			    insert(alpha_depth, dx, j, distance);
-
-			    unsigned int rgba = g->colors[i];
-			    float4 color = get_color(direction, t, rgba);
-
-			    insert(alpha_depth, color_a, j, color);
-			}
-		    }
-					
-		    if (n < alpha_depth)
-			n++;
-		}
-				
-	    } // triangle loop
-	    
-	    node--;
-	    
-	} // node is a leaf
-	
-    } // while loop
-    while (node != head);
-
-    if (n < 1) {
-	pixels[id] = 0;
-	return;
-    }
-
-    dxlen[id] = n;
-
-    float scale = 1.0f;
-    float fr = 0.0f;
-    float fg = 0.0f;
-    float fb = 0.0f;
-    for (i=0; i < n; i++) {
-	float alpha = color_a[i].w;
-
-	fr += scale*color_a[i].x*alpha;
-	fg += scale*color_a[i].y*alpha;
-	fb += scale*color_a[i].z*alpha;
-	
-	scale *= (1.0f-alpha);
-    }
-    unsigned int a;
-    if (n < alpha_depth)
-	a = floorf(255*(1.0f-scale));
-    else
-    	a = 255;
-    unsigned int red = floorf(fr/(1.0f-scale));
-    unsigned int green = floorf(fg/(1.0f-scale));
-    unsigned int blue = floorf(fb/(1.0f-scale));
-
-    pixels[id] = a << 24 | red << 16 | green << 8 | blue;
-}
-
-} // extern "C"