From e06a8b551c730e3d1111fc571c5d48edb85f70ce Mon Sep 17 00:00:00 2001 From: Stan Seibert Date: Tue, 7 Jun 2011 21:53:10 -0400 Subject: Switch triangle texture to device array, use int32 and float32 datatypes everywhere, and build final mesh without concatenation of lists. This allows for very large detectors, like full size LBNE. --- src/kernel.cu | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) (limited to 'src') diff --git a/src/kernel.cu b/src/kernel.cu index 4d85f6e..796de54 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -12,7 +12,7 @@ /* flattened triangle mesh */ texture vertices; -texture triangles; +__device__ uint4 *triangles; /* material/surface index lookup for each triangle */ texture material1_index; @@ -87,7 +87,7 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con { for (i=0; i < length; i++) { - uint4 triangle_data = tex1Dfetch(triangles, index+i); + uint4 triangle_data = triangles[index+i]; float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); @@ -125,6 +125,11 @@ __device__ curandState rng_states[256*512]; extern "C" { +__global__ void set_pointer(uint4 *triangle_ptr) +{ + triangles = triangle_ptr; +} + /* Initialize random number states */ __global__ void init_rng(unsigned long long seed, unsigned long long offset) { @@ -191,7 +196,7 @@ __global__ void ray_trace(int nthreads, float3 *origins, float3 *directions, int } else { - uint4 triangle_data = tex1Dfetch(triangles, intersection_idx); + uint4 triangle_data = triangles[intersection_idx]; float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); -- cgit