diff options
Diffstat (limited to 'src/kernel.cu')
| -rw-r--r-- | src/kernel.cu | 23 |
1 files changed, 12 insertions, 11 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 405f06c..2d579be 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -22,7 +22,7 @@ enum }; /* flattened triangle mesh */ -texture<float4, 1, cudaReadModeElementType> vertices; +__device__ float4 *vertices; __device__ uint4 *triangles; /* lower/upper bounds for the bounding box associated with each node/leaf */ @@ -110,9 +110,9 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con uint4 triangle_data = triangles[index+i]; - 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)); + float3 v0 = make_float3(vertices[triangle_data.x]); + float3 v1 = make_float3(vertices[triangle_data.y]); + float3 v2 = make_float3(vertices[triangle_data.z]); if (intersect_triangle(origin, direction, v0, v1, v2, distance)) { @@ -146,9 +146,10 @@ __device__ curandState rng_states[100000]; extern "C" { -__global__ void set_pointer(uint4 *triangle_ptr) + __global__ void set_pointer(uint4 *triangle_ptr, float4 *vertex_ptr) { triangles = triangle_ptr; + vertices = vertex_ptr; } /* Initialize random number states */ @@ -213,9 +214,9 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i { uint4 triangle_data = triangles[triangle_index]; - 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)); + float3 v0 = make_float3(vertices[triangle_data.x]); + float3 v1 = make_float3(vertices[triangle_data.y]); + float3 v2 = make_float3(vertices[triangle_data.z]); pixels[id] = get_color(direction, v0, v1, v2, triangle_data.w); } @@ -262,9 +263,9 @@ __global__ void propagate(int nthreads, float3 *positions, float3 *directions, f uint4 triangle_data = triangles[last_hit_triangle]; - 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)); + float3 v0 = make_float3(vertices[triangle_data.x]); + float3 v1 = make_float3(vertices[triangle_data.y]); + float3 v2 = make_float3(vertices[triangle_data.z]); int material_in_index = convert(0xFF & (triangle_data.w >> 24)); int material_out_index = convert(0xFF & (triangle_data.w >> 16)); |