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.

1 files changed, 8 insertions, 3 deletions

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<float4, 1, cudaReadModeElementType> vertices;
-texture<uint4, 1, cudaReadModeElementType> triangles;
+__device__ uint4 *triangles;
 
 /* material/surface index lookup for each triangle */
 texture<int, 1, cudaReadModeElementType> 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));