Use Mako defines to generate momenta boundaries

2 files changed, 67 insertions, 10 deletions

diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
new file mode 100644
index 0000000..50044d2
--- /dev/null
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -0,0 +1,46 @@
+<%def name="momenta_boundary(name, param)">
+__kernel void ${name}_momenta_boundary(
+   __global ${float_type}* f_next,
+   __global ${float_type}* f_prev,
+   unsigned int gid, ${param})
+{
+    __global ${float_type}* preshifted_f_next = f_next + gid;
+    __global ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, c_i in enumerate(descriptor.c):
+    const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+    ${caller.body()}
+
+% for i, expr in enumerate(collision_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i in range(0,descriptor.q):
+    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
+</%def>
+
+<%call expr="momenta_boundary('velocity', '%s%d velocity' % (float_type, descriptor.d))">
+    ${float_type} ${ccode(moments_assignment[0])}
+% for i, expr in enumerate(moments_assignment[1:]):
+    ${float_type} ${expr.lhs} = velocity.${['x', 'y', 'z'][i]};
+% endfor
+</%call>
+
+<%call expr="momenta_boundary('density', '%s density' % float_type)">
+    ${float_type} ${moments_assignment[0].lhs} = density;
+% for i, expr in enumerate(moments_assignment[1:]):
+    ${float_type} ${ccode(expr)}
+% endfor
+</%call>
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
index d7e0dd2..5475241 100644
--- a/boltzgen/kernel/template/momenta_boundary.cpp.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -1,7 +1,8 @@
-void velocity_momenta_boundary(      ${float_type}* f_next,
-                               const ${float_type}* f_prev,
-                               std::size_t gid,
-                               ${float_type} velocity[${descriptor.d}])
+<%def name="momenta_boundary(name, param)">
+void ${name}_momenta_boundary(
+          ${float_type}* f_next,
+    const ${float_type}* f_prev,
+    std::size_t gid, ${param})
 {
           ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
     const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
@@ -14,12 +15,7 @@ void velocity_momenta_boundary(      ${float_type}* f_next,
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
 
-% for i, expr in enumerate(moments_assignment):
-    ${float_type} ${ccode(expr)}
-% endfor
-% for i, expr in enumerate(moments_assignment[1:]):
-    ${expr.lhs} = velocity[${i}];
-% endfor
+    ${caller.body()}
 
 % for i, expr in enumerate(collision_subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
@@ -33,3 +29,18 @@ void velocity_momenta_boundary(      ${float_type}* f_next,
     preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
 % endfor
 }
+</%def>
+
+<%call expr="momenta_boundary('velocity', '%s velocity[%d]' % (float_type, descriptor.d))">
+    ${float_type} ${ccode(moments_assignment[0])}
+% for i, expr in enumerate(moments_assignment[1:]):
+    ${float_type} ${expr.lhs} = velocity[${i}];
+% endfor
+</%call>
+
+<%call expr="momenta_boundary('density', '%s density' % float_type)">
+    ${float_type} ${moments_assignment[0].lhs} = density;
+% for i, expr in enumerate(moments_assignment[1:]):
+    ${float_type} ${ccode(expr)}
+% endfor
+</%call>