Extract GL moments, particle buffers and add texture buffer

3 files changed, 138 insertions, 79 deletions

diff --git a/template/kernel.mako b/template/kernel.mako
index a57c2ec..025baf5 100644
--- a/template/kernel.mako
+++ b/template/kernel.mako
@@ -98,82 +98,3 @@ __kernel void collect_moments(__global __read_only  ${float_type}* f,
     moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
 % endfor
 }
-
-__kernel void collect_gl_moments(__global __read_only  ${float_type}* f,
-                                 __global __read_only  int* material,
-                                 __global __write_only float4* moments)
-{
-    const unsigned int gid = ${gid()};
-
-    __global __read_only ${float_type}* preshifted_f = f + gid;
-
-% for i in range(0,descriptor.q):
-    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-    float4 data;
-
-    if (material[gid] == 1) {
-% if descriptor.d == 2:
-      data.x = ${ccode(moments_assignment[0].rhs)};
-      data.y = ${ccode(moments_assignment[1].rhs)};
-      data.z = ${ccode(moments_assignment[2].rhs)};
-      data.w = sqrt(data.y*data.y + data.z*data.z);
-% elif descriptor.d == 3:
-      data.x = ${ccode(moments_assignment[0].rhs)};
-      data.y = ${ccode(moments_assignment[1].rhs)};
-      data.z = ${ccode(moments_assignment[2].rhs)};
-      data.w = ${ccode(moments_assignment[3].rhs)};
-% endif
-    } else {
-      data.x = 0.0;
-      data.y = 0.0;
-      data.z = 0.0;
-      data.w = -material[gid];
-    }
-
-    moments[gid] = data;
-}
-
-__kernel void update_particles(__global __read_only  float4* moments,
-                               __global __read_only  int*    material,
-                               __global __write_only float4* particles,
-                               __global __read_only  float4* init_particles,
-                               float aging)
-{
-  const unsigned int pid = get_global_id(0);
-
-  float4 particle = particles[pid];
-
-% if descriptor.d == 2:
-  const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x);
-% elif descriptor.d == 3:
-  const unsigned int gid = floor(particle.z)*${memory.size_x*memory.size_y} + floor(particle.y)*${memory.size_x} + floor(particle.x);
-% endif
-
-  const float4 moment = moments[gid];
-
-  if (material[gid] == 1) {
-    particle.x += moment.y;
-    particle.y += moment.z;
-% if descriptor.d == 2:
-    particle.w += min(particle.x, particle.y) * aging;
-% elif descriptor.d == 3:
-    particle.z += moment.w;
-    float dy = (particle.y-${geometry.size_y/2.0});
-    float dz = (particle.z-${geometry.size_z/2.0});
-    dy *= dy;
-    dz *= dz;
-    particle.w = 10.0*sqrt(moment.y*moment.y+moment.z*moment.z+moment.w*moment.w);
-% endif
-  } else {
-    particle.xyz = init_particles[pid].xyz;
-    particle.w   = 0.0;
-  }
-
-  particles[pid] = particle;
-}
diff --git a/template/opengl.mako b/template/opengl.mako
new file mode 100644
index 0000000..6b09f69
--- /dev/null
+++ b/template/opengl.mako
@@ -0,0 +1,100 @@
+<%
+def gid():
+    return {
+        2: 'get_global_id(1)*%d + get_global_id(0)' % memory.size_x,
+        3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (memory.size_x*memory.size_y, memory.size_x)
+    }.get(descriptor.d)
+
+def pop_offset(i):
+    return i * memory.volume
+
+def moments_cell():
+    return {
+        2: '(int2)(get_global_id(0), get_global_id(1))',
+        3: '(int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0)'
+    }.get(descriptor.d)
+%>
+
+__kernel void collect_gl_moments(__global __read_only  ${float_type}* f,
+                                 __global __read_only  int* material,
+                                 __global __write_only float4* moments)
+{
+    const unsigned int gid = ${gid()};
+
+    __global __read_only ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+    float4 data;
+
+    if (material[gid] == 1) {
+% if descriptor.d == 2:
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      data.w = sqrt(data.y*data.y + data.z*data.z);
+% elif descriptor.d == 3:
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      data.w = ${ccode(moments_assignment[3].rhs)};
+% endif
+    } else {
+      data.x = 0.0;
+      data.y = 0.0;
+      data.z = 0.0;
+      data.w = -material[gid];
+    }
+
+    moments[gid] = data;
+}
+
+__kernel void collect_gl_moments_to_texture(__global __read_only  ${float_type}* f,
+                                            __global __read_only  int* material,
+% if descriptor.d == 2:
+                                            __write_only image2d_t moments)
+% elif descriptor.d == 3:
+                                            __write_only image3d_t moments)
+% endif
+{
+    const unsigned int gid = ${gid()};
+
+    __global __read_only ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+    float4 data;
+
+    if (material[gid] != 5) {
+% if descriptor.d == 2:
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      data.w = sqrt(data.y*data.y + data.z*data.z);
+% elif descriptor.d == 3:
+      data.x = ${ccode(moments_assignment[0].rhs)};
+      data.y = ${ccode(moments_assignment[1].rhs)};
+      data.z = ${ccode(moments_assignment[2].rhs)};
+      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;
+    }
+
+    write_imagef(moments, ${moments_cell()}, data);
+}
diff --git a/template/particles.mako b/template/particles.mako
new file mode 100644
index 0000000..6235307
--- /dev/null
+++ b/template/particles.mako
@@ -0,0 +1,38 @@
+__kernel void update_particles(__global __read_only  float4* moments,
+                               __global __read_only  int*    material,
+                               __global __write_only float4* particles,
+                               __global __read_only  float4* init_particles,
+                               float aging)
+{
+  const unsigned int pid = get_global_id(0);
+
+  float4 particle = particles[pid];
+
+% if descriptor.d == 2:
+  const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x);
+% elif descriptor.d == 3:
+  const unsigned int gid = floor(particle.z)*${memory.size_x*memory.size_y} + floor(particle.y)*${memory.size_x} + floor(particle.x);
+% endif
+
+  const float4 moment = moments[gid];
+
+  if (material[gid] == 1 && particle.w < 1.0) {
+    particle.x += moment.y;
+    particle.y += moment.z;
+% if descriptor.d == 2:
+    particle.w += min(particle.x, particle.y) * aging;
+% elif descriptor.d == 3:
+    particle.z += moment.w;
+    float dy = (particle.y-${geometry.size_y/2.0});
+    float dz = (particle.z-${geometry.size_z/2.0});
+    dy *= dy;
+    dz *= dz;
+    particle.w = 10.0*sqrt(moment.y*moment.y+moment.z*moment.z+moment.w*moment.w);
+% endif
+  } else {
+    particle.xyz = init_particles[pid].xyz;
+    particle.w   = particle.w-1.0;
+  }
+
+  particles[pid] = particle;
+}