From 5828235f806c3e87a5b1eed34ef69ef317a110bd Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Mon, 4 Nov 2019 23:38:36 +0100
Subject: Extract streaming pattern into Mako definitions

This should allow for plugging in e.g. a AA pattern implementation
without without touching any file but `AA.$target.mako`.

OpenCL and C++ target templates now look basically the same and could
potentially be merged. However this would decrease flexibility should
more differences appear in the future.  Maintaining separate template
files is an acceptable overhead to preserve flexibility.
---
 boltzgen/kernel/template/momenta_boundary.cl.mako | 47 ++++-------------------
 1 file changed, 7 insertions(+), 40 deletions(-)

(limited to 'boltzgen/kernel/template/momenta_boundary.cl.mako')

diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
index b0b4c9e..9b41a70 100644
--- a/boltzgen/kernel/template/momenta_boundary.cl.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -1,23 +1,12 @@
+<%namespace name="pattern" file="${'/pattern/%s.cl.mako' % context['streaming']}"/>
 <%
 import sympy
-
 moments_subexpr, moments_assignment = model.moments()
 collision_subexpr, collision_assignment = model.collision(f_eq = model.equilibrium(resolve_moments = False))
 %>
 
-<%def name="momenta_boundary(name, param)">
-__kernel void ${name}_momenta_boundary_gid(
-   __global ${float_type}* f_next,
-   __global ${float_type}* f_prev,
-   unsigned int gid, ${param})
-{
-    __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
-    __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('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
-
+<%def name="momenta_boundary(name, params)">
+<%call expr="pattern.operator_ab('%s_momenta_boundary' % name, params)">
 % for i, expr in enumerate(moments_subexpr):
     const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
 % endfor
@@ -31,41 +20,19 @@ __kernel void ${name}_momenta_boundary_gid(
 % for i, expr in enumerate(collision_assignment):
     const ${float_type} ${sympy.ccode(expr)}
 % endfor
-
-% for i, expr in enumerate(collision_assignment):
-    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
+</%call>
 </%def>
 
-<%call expr="momenta_boundary('velocity', '%s%d velocity' % (float_type, descriptor.d))">
+<%call expr="momenta_boundary('velocity', [('%s%d' % (float_type, descriptor.d), 'velocity')])">
     ${float_type} ${sympy.ccode(moments_assignment[0])}
 % for i, expr in enumerate(moments_assignment[1:]):
-    ${float_type} ${expr.lhs} = velocity.${['x', 'y', 'z'][i]};
+    ${float_type} ${expr.lhs} = velocity[${i}];
 % endfor
 </%call>
 
-<%call expr="momenta_boundary('density', '%s density' % float_type)">
+<%call expr="momenta_boundary('density', [(float_type, 'density')])">
     ${float_type} ${moments_assignment[0].lhs} = density;
 % for i, expr in enumerate(moments_assignment[1:]):
     ${float_type} ${sympy.ccode(expr)}
 % endfor
 </%call>
-
-% if 'cell_list_dispatch' in extras:
-__kernel void velocity_momenta_boundary_cells(__global ${float_type}* f_next,
-                                              __global ${float_type}* f_prev,
-                                              __global unsigned int*  cells,
-                                              ${float_type}${descriptor.d} velocity)
-{
-    velocity_momenta_boundary_gid(f_next, f_prev, cells[get_global_id(0)], velocity);
-}
-
-__kernel void density_momenta_boundary_cells(__global ${float_type}* f_next,
-                                             __global ${float_type}* f_prev,
-                                             __global unsigned int*  cells,
-                                             ${float_type} density)
-{
-    density_momenta_boundary_gid(f_next, f_prev, cells[get_global_id(0)], density);
-}
-% endif
-- 
cgit v1.2.3