Implement basic diffusive lighting in volumetric raytracer

Approximated surface normals are encoded into the moments texture.
These normals are restored when the ray marcher encounters an
impermeable cell. Note that normals are shifted to be >= 0 as negative
values are not representable using OpenGL textures.

If this prototype works out it might be preferable to calculate better
surface normals during preprocessing to be stored separately.

Compare to 6123c8a.

1 files changed, 27 insertions, 4 deletions

diff --git a/template/opengl.mako b/template/opengl.mako
index 6b09f69..50cdbbc 100644
--- a/template/opengl.mako
+++ b/template/opengl.mako
@@ -55,6 +55,15 @@ __kernel void collect_gl_moments(__global __read_only  ${float_type}* f,
     moments[gid] = data;
 }
 
+<%
+def neighbor_offset(c_i):
+    return {
+        2: lambda:                                      c_i[1]*memory.size_x + c_i[0],
+        3: lambda: c_i[2]*memory.size_x*memory.size_y + c_i[1]*memory.size_x + c_i[0]
+    }.get(descriptor.d)()
+
+%>
+
 __kernel void collect_gl_moments_to_texture(__global __read_only  ${float_type}* f,
                                             __global __read_only  int* material,
 % if descriptor.d == 2:
@@ -90,10 +99,24 @@ __kernel void collect_gl_moments_to_texture(__global __read_only  ${float_type}*
       data.w = ${ccode(moments_assignment[3].rhs)};
 % endif
     } else {
-      data.x = 0.0;
-      data.y = 0.0;
-      data.z = 0.0;
-      data.w = 1.0;
+      const int material_west  = material[gid + ${neighbor_offset((-1,0,0))}];
+      const int material_east  = material[gid + ${neighbor_offset((1,0,0))}];
+      const int material_north = material[gid + ${neighbor_offset((0,1,0))}];
+      const int material_south = material[gid + ${neighbor_offset((0,-1,0))}];
+      const int material_up    = material[gid + ${neighbor_offset((0,0, 1))}];
+      const int material_down  = material[gid + ${neighbor_offset((0,0,-1))}];
+
+      // recover surface normal approximation using surrounding materials
+      float3 n;
+      if (material_west  != 5) { n.x =  1; }
+      if (material_east  != 5) { n.x = -1; }
+      if (material_north != 5) { n.y = -1; }
+      if (material_south != 5) { n.y =  1; }
+      if (material_up    != 5) { n.z = -1; }
+      if (material_down  != 5) { n.z =  1; }
+
+      data.xyz = 0.5 + 0.5*n; // pack surface normal into texture
+      data.w   = 1.0; // signal impermeable material to raytracer
     }
 
     write_imagef(moments, ${moments_cell()}, data);