diff options
Diffstat (limited to 'src/intersect.cu')
| -rw-r--r-- | src/intersect.cu | 15 |
1 files changed, 11 insertions, 4 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)) { |