#ifndef __MESH_H__
#define __MESH_H__

#include "intersect.h"

#define STACK_SIZE 500

/* flattened triangle mesh */
__device__ float3 *g_vertices;
__device__ uint4 *g_triangles;
__device__ unsigned int *g_colors;
__device__ unsigned int g_start_node;
__device__ unsigned int g_first_node;

/* lower/upper bounds for the bounding box associated with each node/leaf */
__device__ float3 *g_lower_bounds;
__device__ float3 *g_upper_bounds;

/* map to child node/triangle indices */
texture<unsigned int, 1, cudaReadModeElementType> node_map;
texture<unsigned int, 1, cudaReadModeElementType> node_map_end;

__device__ float3 make_float3(const float4 &a)
{
	return make_float3(a.x, a.y, a.z);
}

__device__ int convert(int c)
{
	if (c & 0x80)
		return (0xFFFFFF00 | c);
	else
		return c;
}

/* Test the intersection between a ray starting at `origin` traveling in the
   direction `direction` and the bounding box around node `i`. If the ray
   intersects the bounding box return true, else return false. */
__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i, const float &min_distance)
{
	float3 lower_bound = g_lower_bounds[i];
	float3 upper_bound = g_upper_bounds[i];

	float distance_to_box;

	if (intersect_box(origin, direction, lower_bound, upper_bound, distance_to_box))
	{
		if (min_distance < 0.0f)
			return true;

		if (distance_to_box > min_distance)
			return false;

		return true;
	}
	else
	{
		return false;
	}

} // intersect_node

/* Find the intersection between a ray starting at `origin` traveling in the
   direction `direction` and the global mesh texture. If the ray intersects
   the texture return the index of the triangle which the ray intersected,
   else return -1. */
__device__ int intersect_mesh(const float3 &origin, const float3& direction, float &min_distance, int last_hit_triangle = -1)
{
	int triangle_index = -1;

	float distance;
	min_distance = -1.0f;

	if (!intersect_node(origin, direction, g_start_node, min_distance))
		return -1;

	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;

	unsigned int i;

	do
	{
		unsigned int first_child = tex1Dfetch(node_map, *node);
		unsigned int stop = tex1Dfetch(node_map_end, *node);

		while (*node >= g_first_node && stop == first_child+1)
		{
			*node = first_child;
			first_child = tex1Dfetch(node_map, *node);
			stop = tex1Dfetch(node_map_end, *node);
		}
		
		if (*node >= g_first_node)
		{
			for (i=first_child; i < stop; i++)
			{
				if (intersect_node(origin, direction, i, min_distance))
				{
					*node = i;
					node++;
				}
			}

			node--;
		}
		else // node is a leaf
		{
			for (i=first_child; i < stop; i++)
			{
				if (last_hit_triangle == i)
					continue;

				uint4 triangle_data = g_triangles[i];

				float3 v0 = g_vertices[triangle_data.x];
				float3 v1 = g_vertices[triangle_data.y];
				float3 v2 = g_vertices[triangle_data.z];

				if (intersect_triangle(origin, direction, v0, v1, v2, distance))
				{
					if (triangle_index == -1)
					{
						triangle_index = i;
						min_distance = distance;
						continue;
					}

					if (distance < min_distance)
					{
						triangle_index = i;
						min_distance = distance;
					}
				}
			} // triangle loop

			node--;

		} // node is a leaf

	} // while loop
	while (node != head);

	return triangle_index;
}

extern "C"
{

__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node, float3 *lower_bounds, float3 *upper_bounds)
{
	g_triangles = triangles;
	g_vertices = vertices;
	g_colors = colors;
	g_start_node = start_node;
	g_first_node = first_node;
	g_lower_bounds = lower_bounds;
	g_upper_bounds = upper_bounds;
}

__global__ void set_colors(unsigned int *colors)
{
	g_colors = colors;
}

__global__ void color_solids(int nthreads, unsigned int ntriangles, int *solid_id_map, int *solid_ids, unsigned int *solid_colors)
{
	int id = blockIdx.x*blockDim.x + threadIdx.x;

	if (id >= nthreads)
		return;

	int solid_id = solid_ids[id];
	unsigned int color = solid_colors[id];

	for (int i=0; i < ntriangles; i++)
	{
		if (solid_id_map[i] == solid_id)
			g_colors[i] = color;
	}
}

} // extern "c"

#endif