diff options
| author | Anthony LaTorre <telatorre@gmail.com> | 2011-05-27 12:45:12 -0400 |
|---|---|---|
| committer | Anthony LaTorre <telatorre@gmail.com> | 2011-05-27 12:45:12 -0400 |
| commit | 343d4b3e726595b9f5cda34d1e714098e10f8757 (patch) | |
| tree | 3a88d2c826b129f7c89dc7dab3c78280add45788 /src | |
| parent | 6ff042998b8c93652b82e6f34d9dfc1ef40f4c56 (diff) | |
| download | chroma-343d4b3e726595b9f5cda34d1e714098e10f8757.tar.gz chroma-343d4b3e726595b9f5cda34d1e714098e10f8757.tar.bz2 chroma-343d4b3e726595b9f5cda34d1e714098e10f8757.zip | |
kernel ran off the end of the ray array
Diffstat (limited to 'src')
| -rw-r--r-- | src/kernel.cu | 20 |
1 files changed, 10 insertions, 10 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index f14ed69..5794057 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -125,11 +125,11 @@ __global__ void init_rng(unsigned long long seed, unsigned long long offset) } /* */ -__global__ void uniform_sphere(int max_idx, float3 *points) +__global__ void uniform_sphere(int nthreads, float3 *points) { int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx > max_idx) + if (idx >= nthreads) return; //curandState rng = *(rng_states+idx); @@ -137,11 +137,11 @@ __global__ void uniform_sphere(int max_idx, float3 *points) } /* Translate `points` by the vector `v` */ -__global__ void translate(int max_idx, float3 *points, float3 v) +__global__ void translate(int nthreads, float3 *points, float3 v) { int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx > max_idx) + if (idx >= nthreads) return; *(points+idx) += v; @@ -149,11 +149,11 @@ __global__ void translate(int max_idx, float3 *points, float3 v) /* Rotate `points` through an angle `phi` counter-clockwise about the axis `axis` (when looking towards +infinity). */ -__global__ void rotate(int max_idx, float3 *points, float phi, float3 axis) +__global__ void rotate(int nthreads, float3 *points, float phi, float3 axis) { int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx > max_idx) + if (idx >= nthreads) return; *(points+idx) = rotate(*(points+idx), phi, axis); @@ -164,11 +164,11 @@ __global__ void rotate(int max_idx, float3 *points, float phi, float3 axis) set the pixel associated with the ray to a 32 bit color whose brightness is determined by the cosine of the angle between the ray and the normal of the triangle it intersected, else set the pixel to 0. */ -__global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int start_node, int first_node, int *pixels) +__global__ void ray_trace(int nthreads, float3 *origins, float3 *directions, int start_node, int first_node, int *pixels) { int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx > max_idx) + if (idx >= nthreads) return; float3 origin = *(origins+idx); @@ -198,11 +198,11 @@ __global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int intersects the mesh set the hit_solid array value associated with the photon to the triangle index of the triangle the photon intersected, else set the hit_solid array value to -1. */ -__global__ void propagate(int max_idx, float3 *origins, float3 *directions, int start_node, int first_node, int *hit_triangles) +__global__ void propagate(int nthreads, float3 *origins, float3 *directions, int start_node, int first_node, int *hit_triangles) { int idx = blockIdx.x*blockDim.x + threadIdx.x; - if (idx > max_idx) + if (idx >= nthreads) return; float3 origin = *(origins+idx); |