summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/kernel.cu27
1 files changed, 25 insertions, 2 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index ed53032..0d5e54b 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -1,5 +1,6 @@
//-*-c-*-
#include <math_constants.h>
+#include <curand_kernel.h>
#include "linalg.h"
#include "matrix.h"
@@ -111,9 +112,30 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con
return triangle_idx;
}
+__device__ curandState rng_states[256*512];
+
extern "C"
{
+/* Initialize random number states */
+__global__ void init_rng(unsigned long long seed, unsigned long long offset)
+{
+ int idx = blockIdx.x*blockDim.x + threadIdx.x;
+ curand_init(seed, idx, offset, rng_states+idx);
+}
+
+/* */
+__global__ void uniform_sphere(int max_idx, float3 *points)
+{
+ int idx = blockIdx.x*blockDim.x + threadIdx.x;
+
+ if (idx > max_idx)
+ return;
+
+ //curandState rng = *(rng_states+idx);
+
+}
+
/* Translate `points` by the vector `v` */
__global__ void translate(int max_idx, float3 *points, float3 v)
{
@@ -176,7 +198,7 @@ __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 first_leaf, int *hit_solids)
+__global__ void propagate(int max_idx, float3 *origins, float3 *directions, int first_leaf, int *hit_triangles)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
@@ -187,7 +209,8 @@ __global__ void propagate(int max_idx, float3 *origins, float3 *directions, int
float3 direction = *(directions+idx);
direction /= norm(direction);
- *(hit_solids+idx) = intersect_mesh(origin, direction, first_leaf);
+ *(hit_triangles+idx) = intersect_mesh(origin, direction, first_leaf);
+
} // propagate
} // extern "c"