summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/mesh.h40
-rw-r--r--src/render.cu51
2 files changed, 42 insertions, 49 deletions
diff --git a/src/mesh.h b/src/mesh.h
index c6ef8f6..734ab92 100644
--- a/src/mesh.h
+++ b/src/mesh.h
@@ -12,11 +12,11 @@
less than `min_distance`, return true, else return false. */
__device__ bool
intersect_node(const float3 &origin, const float3 &direction,
- Geometry *geometry, const int &i, float min_distance=-1.0f)
+ Geometry *g, int i, float min_distance=-1.0f)
{
/* assigning these to local variables is faster for some reason */
- float3 lower_bound = geometry->lower_bounds[i];
- float3 upper_bound = geometry->upper_bounds[i];
+ float3 lower_bound = g->lower_bounds[i];
+ float3 upper_bound = g->upper_bounds[i];
float distance_to_box;
@@ -42,17 +42,15 @@ intersect_node(const float3 &origin, const float3 &direction,
the intersection and return the index of the triangle which the ray
intersected, else return -1. */
__device__ int
-intersect_mesh(const float3 &origin, const float3& direction,
- Geometry *geometry, float &min_distance,
- int last_hit_triangle = -1)
+intersect_mesh(const float3 &origin, const float3& direction, Geometry *g,
+ float &min_distance, int last_hit_triangle = -1)
{
int triangle_index = -1;
float distance;
min_distance = -1.0f;
- if (!intersect_node(origin, direction, geometry, geometry->start_node,
- min_distance))
+ if (!intersect_node(origin, direction, g, g->start_node, min_distance))
return -1;
unsigned int stack[STACK_SIZE];
@@ -60,25 +58,24 @@ intersect_mesh(const float3 &origin, const float3& direction,
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;
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,
- min_distance)) {
+ if (intersect_node(origin, direction, g, i, min_distance)) {
*node = i;
node++;
}
@@ -92,10 +89,9 @@ intersect_mesh(const float3 &origin, const float3& direction,
if (last_hit_triangle == i)
continue;
- 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 (triangle_index == -1) {
triangle_index = i;
min_distance = distance;
@@ -145,7 +141,7 @@ distance_to_mesh(int nthreads, float3 *_origin, float3 *_direction,
__global__ void
color_solids(int first_triangle, int nthreads, int *solid_id_map,
- bool *solid_hit, unsigned int *solid_colors, Geometry *geometry)
+ bool *solid_hit, unsigned int *solid_colors, Geometry *g)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
@@ -155,7 +151,7 @@ color_solids(int first_triangle, int nthreads, int *solid_id_map,
int triangle_id = first_triangle + id;
int solid_id = solid_id_map[triangle_id];
if (solid_hit[solid_id])
- geometry->colors[triangle_id] = solid_colors[solid_id];
+ g->colors[triangle_id] = solid_colors[solid_id];
}
} // extern "C"
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"
generated by cgit (git 2.34.1) at 2025-05-08 23:00:19 +0000