move photon initialization to the gpu. it's unclear if this is a speed improvement.

1 files changed, 17 insertions, 10 deletions

diff --git a/src/kernel.cu b/src/kernel.cu
index 584c9e4..55fde70 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -141,26 +141,33 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con
 	return triangle_index;
 }
 
-__device__ curandState rng_states[100000];
-
 extern "C"
 {
 
-  __global__ void set_pointer(uint4 *triangle_ptr, float3 *vertex_ptr)
+__global__ void fill_float(int nthreads, float *a, float value)
 {
-	triangles = triangle_ptr;
-	vertices = vertex_ptr;
+	int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (id >= nthreads)
+		return;
+
+	a[id] = value;
 }
 
-/* Initialize random number states */
-__global__ void init_rng(int nthreads, unsigned long long seed, unsigned long long offset)
+__global__ void fill_float3(int nthreads, float3 *a, float3 value)
 {
 	int id = blockIdx.x*blockDim.x + threadIdx.x;
 
 	if (id >= nthreads)
 		return;
 
-	curand_init(seed, id, offset, rng_states+id);
+	a[id] = value;
+}
+
+__global__ void set_pointer(uint4 *triangle_ptr, float3 *vertex_ptr)
+{
+	triangles = triangle_ptr;
+	vertices = vertex_ptr;
 }
 
 /* Translate `points` by the vector `v` */
@@ -223,7 +230,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i
 
 } // ray_trace
 
-__global__ void propagate(int nthreads, float3 *positions, float3 *directions, float *wavelengths, float3 *polarizations, float *times, int *states, int *last_hit_triangles, int start_node, int first_node, int max_steps)
+__global__ void propagate(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, float *wavelengths, float3 *polarizations, float *times, int *states, int *last_hit_triangles, int start_node, int first_node, int max_steps)
 {
 	int id = blockIdx.x*blockDim.x + threadIdx.x;
 
@@ -445,8 +452,8 @@ __global__ void propagate(int nthreads, float3 *positions, float3 *directions, f
 
 	} // while(nsteps < max_steps)
 
-	states[id] = state;
 	rng_states[id] = rng;
+	states[id] = state;
 	positions[id] = position;
 	directions[id] = direction;
 	polarizations[id] = polarization;