diff options
| author | Anthony LaTorre <telatorre@gmail.com> | 2011-06-02 15:00:26 -0400 |
|---|---|---|
| committer | Anthony LaTorre <telatorre@gmail.com> | 2011-06-02 15:00:26 -0400 |
| commit | 8bbdf7f53b918857a09a9bee4a158f13834bfce6 (patch) | |
| tree | 0af80a167c25af46eb21f285f5bca4647c0b7e5c /src | |
| parent | d0825a136ff65b36069ff8b078b9fd97adeed0df (diff) | |
| download | chroma-8bbdf7f53b918857a09a9bee4a158f13834bfce6.tar.gz chroma-8bbdf7f53b918857a09a9bee4a158f13834bfce6.tar.bz2 chroma-8bbdf7f53b918857a09a9bee4a158f13834bfce6.zip | |
triangle mesh is now stored everywhere as a split list of vertices and triangles
Diffstat (limited to 'src')
| -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); } |