diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/kernel.cu | 2 | ||||
| -rw-r--r-- | src/mesh.h | 12 |
2 files changed, 8 insertions, 6 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 17b829c..fe518f6 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -17,7 +17,7 @@ __device__ void fAtomicAdd(float *addr, float data) } } -__device__ __noinline__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max_steps) +__device__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max_steps) { int steps = 0; while (steps < max_steps) @@ -13,8 +13,8 @@ __device__ unsigned int g_start_node; __device__ unsigned int g_first_node; /* lower/upper bounds for the bounding box associated with each node/leaf */ -texture<float4, 1, cudaReadModeElementType> upper_bounds; -texture<float4, 1, cudaReadModeElementType> lower_bounds; +__device__ float3 *g_lower_bounds; +__device__ float3 *g_upper_bounds; /* map to child node/triangle indices */ texture<unsigned int, 1, cudaReadModeElementType> node_map; @@ -38,8 +38,8 @@ __device__ int convert(int c) intersects the bounding box return true, else return false. */ __device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i) { - float3 lower_bound = make_float3(tex1Dfetch(lower_bounds, i)); - float3 upper_bound = make_float3(tex1Dfetch(upper_bounds, i)); + float3 lower_bound = g_lower_bounds[i]; + float3 upper_bound = g_upper_bounds[i]; return intersect_box(origin, direction, lower_bound, upper_bound); } @@ -134,13 +134,15 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo extern "C" { -__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node) + __global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node, float3 *lower_bounds, float3 *upper_bounds) { g_triangles = triangles; g_vertices = vertices; g_colors = colors; g_start_node = start_node; g_first_node = first_node; + g_lower_bounds = lower_bounds; + g_upper_bounds = upper_bounds; } __global__ void set_colors(unsigned int *colors) |