move triangles to mesh

2 files changed, 11 insertions, 5 deletions

diff --git a/src/intersect.cu b/src/intersect.cu
index 0bb73a0..c78350e 100644
--- a/src/intersect.cu
+++ b/src/intersect.cu
@@ -1,5 +1,7 @@
 //-*-c-*-
 
+texture<float4, 1, cudaReadModeElementType> mesh;
+
 __device__ bool intersect_triangle(const float3 &x, const float3 &p, const float3 &v0, const float3 &v1, const float3 &v2, float3 &intersection)
 {
 	Matrix m = make_matrix(v1-v0, v2-v0, -p);
@@ -45,6 +47,11 @@ __device__ int get_color(const float3 &p, const float3 &v0, const float3& v1, co
 	return rgb*65536 + rgb*256 + rgb;
 }
 
+__device__ float3 make_float3(const float4 &a)
+{
+	return make_float3(a.x, a.y, a.z);
+}
+
 extern "C"
 {
 
@@ -68,7 +75,7 @@ __global__ void rotate(int max_idx, float3 *x, float phi, float3 axis)
 	x[idx] = rotate(x[idx], phi, axis);
 }
 
-__global__ void intersect_triangle_mesh(int max_idx, float3 *xarr, float3 *parr, int n, float3* mesh, int *pixelarr)
+__global__ void intersect_triangle_mesh(int max_idx, float3 *xarr, float3 *parr, int n, int *pixelarr)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 
@@ -87,9 +94,9 @@ __global__ void intersect_triangle_mesh(int max_idx, float3 *xarr, float3 *parr,
 	int i;
 	for (i=0; i < n; i++)
 	{
-		float3 v0 = *(mesh+3*i);
-		float3 v1 = *(mesh+3*i+1);
-		float3 v2 = *(mesh+3*i+2);
+		float3 v0 = make_float3(tex1Dfetch(mesh, 3*i));
+		float3 v1 = make_float3(tex1Dfetch(mesh, 3*i+1));
+		float3 v2 = make_float3(tex1Dfetch(mesh, 3*i+2));
 			
 		if (intersect_triangle(x, p, v0, v1, v2, intersection))
 		{
diff --git a/src/rotate.h b/src/rotate.h
index 52d6d6a..fec76a8 100644
--- a/src/rotate.h
+++ b/src/rotate.h
@@ -2,7 +2,6 @@
 #define __ROTATE_H__
 
 __device__ const Matrix IDENTITY_MATRIX = {1,0,0,0,1,0,0,0,1};
-__device__ const Matrix ZERO_MATRIX = {0,0,0,0,0,0,0,0,0};
 
 __device__ __host__ Matrix make_rotation_matrix(float phi, const float3 &n)
 {