Task gpu012: Laser show in CUDA/OpenCL (C++/C#)

The task is to implement reflection of laser rays on planar mirror surfaces in simple 3D scene. You have to use CUDA or OpenCL environment for computing reflections, CPU version is optional (you can add it for reference).
Scene consists of many mirror faces (triangles) and a set of laser generators. Scene and/or laser generators should be animated. Your GPU kernel computes all ray-mirror intersections, generated reflected rays, etc. Finally the scene is rendered using OpenGL (interop). User is able to explore the 3D scene interactively using trackball..

laser show

Details

Rays: rays will be simple line segments, drawn using GL_LINE primitives. You are advised to use a few different primary colors (laser gun colors) and to dim every ray a little after bouncing on the scene face. You can define a limit for number of bounces (4 to 10).

Animation: ray generators and/or scene components (disco-ball..) should be animated. Animation can be done on CPU (modification of transform matrices) and passed to GPU kernel using constant values. Different approach is to do animation in GPU kernel entirely - as well as intersection computations.

Ray-triangle intersection: you can write your own code (perhaps based on Geometry.RayTriangleIntersection() from the C# library) or use the following CUDA code:

struct STriangleCUDA
{
  float3 a, b, c;
  float3 edge1, edge2, normal;
};

__device__ bool IntersectTriangle ( STriangleCUDA *tri, float3 orig, float3 dir, float *t )
{
  const float3 pvec = make_float3( dir.y*tri->edge2.z - dir.z*tri->edge2.y,          // dir ^ tri.edge2;
                                   dir.z*tri->edge2.x - dir.x*tri->edge2.z,
                                   dir.x*tri->edge2.y - dir.y*tri->edge2.x );
  const float det = tri->edge1.x*pvec.x + tri->edge1.y*pvec.y + tri->edge1.z*pvec.z; // tri.edge1 * pvec;
  if ( fabs( det ) < EPSILON )
    return false;

  const float inv_det = 1.0f / det;

  const float3 tvec = make_float3( orig.x - tri->a.x, orig.y - tri->a.y, orig.z - tri->a.z ); // orig - tri.a;
  float lambda = tvec.x*pvec.x + tvec.y*pvec.y + tvec.z*pvec.z; //tvec * pvec;
  lambda *= inv_det;

  if ( lambda < -MINIEPSILON || lambda > 1.0f + MINIEPSILON )
    return false;

  const float3 qvec = make_float3( tvec.y*tri->edge1.z - tvec.z*tri->edge1.y,        // tvec ^ tri.edge1;
                                   tvec.z*tri->edge1.x - tvec.x*tri->edge1.z,
                                   tvec.x*tri->edge1.y - tvec.y*tri->edge1.x );
  float mue = dir.x*qvec.x + dir.y*qvec.y + dir.z*qvec.z;  //dir * qvec;
  mue *= inv_det;
  if ( mue < MINIEPSILON || mue+lambda > 1.0f + MINIEPSILON )
    return false;

  float f = tri->edge2.x*qvec.x + tri->edge2.y*qvec.y + tri->edge2.z*qvec.z;         // tri.edge2 * qvec;
  f *= inv_det;
  if ( *t <= f || f < EPSILON )
    return false;

  *t = f;
  return true;
}

Template projects

C#: You can use the 090opencl project as a template. Mandelbrot set is computed on GPU (OpenCL) or CPU, OpenGL interop is optionally used for rendering the result (bitmap texture).

C++: You can use template sample14 which implements almost exact analogy in C++ and CUDA. OpenGL interop is used here as well.

Your solution

Send all your modifications in complete Visual Studio project, including kernel files and optional scene files (I don't expect explicit scene data will be very useful, but it is your choice). Write a brief document about your solution, don't forget to update the F1 key handle in the C++ case.
Define which parts of the assignments you had solved. The specific point assignments are listed below.

Deadline

Hand in before: 1. 7. 2018

Points

Complete CPU implementation: 8 points (valid only if you finish GPGPU solution as well),
14 points for functional CUDA/OpenCL solution,
6 points for basic realtime animation.

Bonuses: up to 5 additional points for multiple colors, attenuation of the rays in reflections,
up to 6 points for more rich animations, multiple variants of movements..
up to 3 points for using OpenGL interop (keeping data in GPU memory all the time),
5-8 points for more shapes in the scene, other than triangles.

Project

Recommended starting point: Visual Studio project 090opencl from the grcis repository (C#) or sample14 from the ogl repository (C++).