summaryrefslogtreecommitdiff
path: root/src/mesh.h
diff options
context:
space:
mode:
authorAnthony LaTorre <tlatorre9@gmail.com>2011-08-11 16:43:09 -0400
committerAnthony LaTorre <tlatorre9@gmail.com>2011-08-11 16:43:09 -0400
commit00aa07a12f8f00e29a3a04f2850ccd4fc91cc951 (patch)
tree50c76a5661693f2335278db2d6cc924d3cc60e20 /src/mesh.h
parent9d68f8ea2b6a64e83195f16aa066caac5eeb8fae (diff)
parent1b9aef380e629c136e045412621483ebb4a4164f (diff)
downloadchroma-00aa07a12f8f00e29a3a04f2850ccd4fc91cc951.tar.gz
chroma-00aa07a12f8f00e29a3a04f2850ccd4fc91cc951.tar.bz2
chroma-00aa07a12f8f00e29a3a04f2850ccd4fc91cc951.zip
merge heads
Diffstat (limited to 'src/mesh.h')
-rw-r--r--src/mesh.h12
1 files changed, 7 insertions, 5 deletions
diff --git a/src/mesh.h b/src/mesh.h
index b4714c4..bb30bef 100644
--- a/src/mesh.h
+++ b/src/mesh.h
@@ -13,8 +13,8 @@ __device__ unsigned int g_start_node;
__device__ unsigned int g_first_node;
/* lower/upper bounds for the bounding box associated with each node/leaf */
-texture<float4, 1, cudaReadModeElementType> upper_bounds;
-texture<float4, 1, cudaReadModeElementType> lower_bounds;
+__device__ float3 *g_lower_bounds;
+__device__ float3 *g_upper_bounds;
/* map to child node/triangle indices */
texture<unsigned int, 1, cudaReadModeElementType> node_map;
@@ -38,8 +38,8 @@ __device__ int convert(int c)
intersects the bounding box return true, else return false. */
__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i)
{
- float3 lower_bound = make_float3(tex1Dfetch(lower_bounds, i));
- float3 upper_bound = make_float3(tex1Dfetch(upper_bounds, i));
+ float3 lower_bound = g_lower_bounds[i];
+ float3 upper_bound = g_upper_bounds[i];
return intersect_box(origin, direction, lower_bound, upper_bound);
}
@@ -134,13 +134,15 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo
extern "C"
{
-__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node)
+ __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)