update event structure. break gpu.GPU class into separate smaller independent classes.

2 files changed, 47 insertions, 35 deletions

diff --git a/src/kernel.cu b/src/kernel.cu
index d36d260..4418305 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -62,41 +62,6 @@ __device__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max
 extern "C"
 {
 
-/* Translate the points `a` by the vector `v` */
-__global__ void translate(int nthreads, float3 *a, float3 v)
-{
-	int id = blockIdx.x*blockDim.x + threadIdx.x;
-
-	if (id >= nthreads)
-		return;
-
-	a[id] += v;
-}
-
-/* Rotate the points `a` through an angle `phi` counter-clockwise about the
-   axis `axis` (when looking towards +infinity). */
-__global__ void rotate(int nthreads, float3 *a, float phi, float3 axis)
-{
-	int id = blockIdx.x*blockDim.x + threadIdx.x;
-
-	if (id >= nthreads)
-		return;
-
-	a[id] = rotate(a[id], phi, axis);
-}
-
-__global__ void rotate_around_point(int nthreads, float3 *a, float phi, float3 axis, float3 point)
-{
-	int id = blockIdx.x*blockDim.x + threadIdx.x;
-
-	if (id >= nthreads)
-		return;
-
-	a[id] -= point;
-	a[id] = rotate(a[id], phi, axis);
-	a[id] += point;
-}
-
 __global__ void update_xyz_lookup(int nthreads, int total_threads, int offset, float3 position, curandState *rng_states, float wavelength, float3 xyz, float3 *xyz_lookup1, float3 *xyz_lookup2, int max_steps)
 {
 	int kernel_id = blockIdx.x*blockDim.x + threadIdx.x;
diff --git a/src/transform.cu b/src/transform.cu
new file mode 100644
index 0000000..57bd509
--- /dev/null
+++ b/src/transform.cu
@@ -0,0 +1,47 @@
+//-*-c-*-
+
+#include "linalg.h"
+#include "rotate.h"
+
+extern "C"
+{
+
+/* Translate the points `a` by the vector `v` */
+__global__ void translate(int nthreads, float3 *a, float3 v)
+{
+	int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (id >= nthreads)
+		return;
+
+	a[id] += v;
+}
+
+/* Rotate the points `a` through an angle `phi` counter-clockwise about the
+   axis `axis` (when looking towards +infinity). */
+__global__ void rotate(int nthreads, float3 *a, float phi, float3 axis)
+{
+	int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (id >= nthreads)
+		return;
+
+	a[id] = rotate(a[id], phi, axis);
+}
+
+/* Rotate the points `a` through an angle `phi` counter-clockwise
+   (when looking towards +infinity along `axis`) about the axis defined
+   by the point `point` and the vector `axis` . */
+__global__ void rotate_around_point(int nthreads, float3 *a, float phi, float3 axis, float3 point)
+{
+	int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+	if (id >= nthreads)
+		return;
+
+	a[id] -= point;
+	a[id] = rotate(a[id], phi, axis);
+	a[id] += point;
+}
+
+} // extern "c"