summaryrefslogtreecommitdiff
path: root/src/render.cu
diff options
context:
space:
mode:
Diffstat (limited to 'src/render.cu')
-rw-r--r--src/render.cu51
1 files changed, 24 insertions, 27 deletions
diff --git a/src/render.cu b/src/render.cu
index d9ce1b1..005950f 100644
--- a/src/render.cu
+++ b/src/render.cu
@@ -9,17 +9,17 @@
#include "stdio.h"
__device__ float4
-get_color(const float3 &direction, const Triangle &triangle, unsigned int rgba)
+get_color(const float3 &direction, const Triangle &t, unsigned int rgba)
{
- float3 v01 = triangle.v1 - triangle.v0;
- float3 v12 = triangle.v2 - triangle.v1;
+ 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 = dot(-surface_normal,-direction);
+ cos_theta = -cos_theta;
unsigned int a0 = 0xff & (rgba >> 24);
unsigned int r0 = 0xff & (rgba >> 16);
@@ -35,7 +35,7 @@ extern "C"
{
__global__ void
-render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
+render(int nthreads, float3 *_origin, float3 *_direction, Geometry *g,
unsigned int alpha_depth, unsigned int *pixels, float *_dx,
unsigned int *dxlen, float4 *_color)
{
@@ -50,8 +50,7 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
float distance;
- if (n < 1 && !intersect_node(origin, direction, geometry,
- geometry->start_node)) {
+ if (n < 1 && !intersect_node(origin, direction, g, g->start_node)) {
pixels[id] = 0;
return;
}
@@ -61,7 +60,7 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
unsigned int *head = &stack[0];
unsigned int *node = &stack[1];
unsigned int *tail = &stack[STACK_SIZE-1];
- *node = geometry->start_node;
+ *node = g->start_node;
float *dx = _dx + id*alpha_depth;
float4 *color_a = _color + id*alpha_depth;
@@ -69,18 +68,18 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
unsigned int i;
do {
- unsigned int first_child = geometry->node_map[*node];
- unsigned int stop = geometry->node_map_end[*node];
+ unsigned int first_child = g->node_map[*node];
+ unsigned int stop = g->node_map_end[*node];
- while (*node >= geometry->first_node && stop == first_child+1) {
+ while (*node >= g->first_node && stop == first_child+1) {
*node = first_child;
- first_child = geometry->node_map[*node];
- stop = geometry->node_map_end[*node];
+ first_child = g->node_map[*node];
+ stop = g->node_map_end[*node];
}
- if (*node >= geometry->first_node) {
+ if (*node >= g->first_node) {
for (i=first_child; i < stop; i++) {
- if (intersect_node(origin, direction, geometry, i)) {
+ if (intersect_node(origin, direction, g, i)) {
*node = i;
node++;
}
@@ -91,15 +90,14 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
else {
// node is a leaf
for (i=first_child; i < stop; i++) {
- Triangle triangle = get_triangle(geometry, i);
+ Triangle t = get_triangle(g, i);
- if (intersect_triangle(origin, direction, triangle,
- distance)) {
+ if (intersect_triangle(origin, direction, t, distance)) {
if (n < 1) {
dx[0] = distance;
- unsigned int rgba = geometry->colors[i];
- float4 color = get_color(direction, triangle, rgba);
+ unsigned int rgba = g->colors[i];
+ float4 color = get_color(direction, t, rgba);
color_a[0] = color;
}
@@ -109,9 +107,8 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
if (j <= alpha_depth-1) {
insert(alpha_depth, dx, j, distance);
- unsigned int rgba = geometry->colors[i];
- float4 color = get_color(direction, triangle,
- rgba);
+ unsigned int rgba = g->colors[i];
+ float4 color = get_color(direction, t, rgba);
insert(alpha_depth, color_a, j, color);
}
@@ -155,11 +152,11 @@ render(int nthreads, float3 *_origin, float3 *_direction, Geometry *geometry,
a = floorf(255*(1.0f-scale));
else
a = 255;
- unsigned int r = floorf(fr/(1.0f-scale));
- unsigned int g = floorf(fg/(1.0f-scale));
- unsigned int b = floorf(fb/(1.0f-scale));
+ 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 | r << 16 | g << 8 | b;
+ pixels[id] = a << 24 | red << 16 | green << 8 | blue;
}
} // extern "C"