diff options
| author | Stan Seibert <stan@mtrr.org> | 2011-06-21 15:46:24 -0400 |
|---|---|---|
| committer | Stan Seibert <stan@mtrr.org> | 2011-06-21 15:46:24 -0400 |
| commit | dc3b5bbf2cb52ae9d0fe2dc5583c29bfdf1b419f (patch) | |
| tree | be40b95e830b2f384ac754ddf6393f2fda5fdd80 /src | |
| parent | d41ad45a2d1c4ffd32375c6b8e24f09d34585913 (diff) | |
| download | chroma-dc3b5bbf2cb52ae9d0fe2dc5583c29bfdf1b419f.tar.gz chroma-dc3b5bbf2cb52ae9d0fe2dc5583c29bfdf1b419f.tar.bz2 chroma-dc3b5bbf2cb52ae9d0fe2dc5583c29bfdf1b419f.zip | |
Switch vertex list on GPU from a texture to a standard device array.
This avoids the texture size limitation.
Diffstat (limited to 'src')
| -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)); |