diff options
Diffstat (limited to 'src/kernel.cu')
| -rw-r--r-- | src/kernel.cu | 11 |
1 files changed, 8 insertions, 3 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 4d85f6e..796de54 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -12,7 +12,7 @@ /* flattened triangle mesh */ texture<float4, 1, cudaReadModeElementType> vertices; -texture<uint4, 1, cudaReadModeElementType> triangles; +__device__ uint4 *triangles; /* material/surface index lookup for each triangle */ texture<int, 1, cudaReadModeElementType> material1_index; @@ -87,7 +87,7 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con { for (i=0; i < length; i++) { - uint4 triangle_data = tex1Dfetch(triangles, index+i); + uint4 triangle_data = triangles[index+i]; float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); @@ -125,6 +125,11 @@ __device__ curandState rng_states[256*512]; extern "C" { +__global__ void set_pointer(uint4 *triangle_ptr) +{ + triangles = triangle_ptr; +} + /* Initialize random number states */ __global__ void init_rng(unsigned long long seed, unsigned long long offset) { @@ -191,7 +196,7 @@ __global__ void ray_trace(int nthreads, float3 *origins, float3 *directions, int } else { - uint4 triangle_data = tex1Dfetch(triangles, intersection_idx); + uint4 triangle_data = triangles[intersection_idx]; float3 v0 = make_float3(tex1Dfetch(vertices, triangle_data.x)); float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y)); |