diff options
Diffstat (limited to 'src/kernel.cu')
| -rw-r--r-- | src/kernel.cu | 19 |
1 files changed, 10 insertions, 9 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 5794057..a020180 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -10,7 +10,8 @@ #define STACK_SIZE 500 /* flattened triangle mesh */ -texture<float4, 1, cudaReadModeElementType> mesh; +texture<float4, 1, cudaReadModeElementType> vertices; +texture<uint4, 1, cudaReadModeElementType> triangles; /* lower/upper bounds for the bounding box associated with each node/leaf */ texture<float4, 1, cudaReadModeElementType> upper_bounds; @@ -79,11 +80,11 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con { for (i=0; i < length; i++) { - int mesh_idx = 3*(index + i); + uint4 triangle_data = tex1Dfetch(triangles, index+i); - float3 v0 = make_float3(tex1Dfetch(mesh, mesh_idx)); - float3 v1 = make_float3(tex1Dfetch(mesh, mesh_idx+1)); - float3 v2 = make_float3(tex1Dfetch(mesh, mesh_idx+2)); + float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); + float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); + float3 v2 = make_float3(tex1Dfetch(vertices, triangle_data.z)); if (intersect_triangle(origin, direction, v0, v1, v2, distance)) { @@ -183,11 +184,11 @@ __global__ void ray_trace(int nthreads, float3 *origins, float3 *directions, int } else { - int mesh_idx = 3*intersection_idx; + uint4 triangle_data = tex1Dfetch(triangles, intersection_idx); - float3 v0 = make_float3(tex1Dfetch(mesh, mesh_idx)); - float3 v1 = make_float3(tex1Dfetch(mesh, mesh_idx+1)); - float3 v2 = make_float3(tex1Dfetch(mesh, mesh_idx+2)); + float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); + float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); + float3 v2 = make_float3(tex1Dfetch(vertices, triangle_data.z)); *(pixels+idx) = get_color(direction, v0, v1, v2); } |