aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/bounce_back_boundary.cl.mako
diff options
context:
space:
mode:
Diffstat (limited to 'boltzgen/kernel/template/bounce_back_boundary.cl.mako')
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako35
1 files changed, 5 insertions, 30 deletions
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
index e26cbe1..7e87a94 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
@@ -1,40 +1,15 @@
+<%namespace name="pattern" file="${'/pattern/%s.cl.mako' % context['streaming']}"/>
<%
import sympy
+subexpr, assignment = model.collision(f_eq = model.equilibrium(resolve_moments = True))
%>
-__kernel void bounce_back_boundary_gid(__global ${float_type}* f_next,
- __global ${float_type}* f_prev,
- unsigned int gid)
-{
- __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
-
-<%
- subexpr, assignment = model.collision(f_eq = model.equilibrium(resolve_moments = True))
-%>
-
+<%call expr="pattern.operator_ab('bounce_back_boundary')">
% for i, expr in enumerate(subexpr):
const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
-% for i, expr in enumerate(assignment):
- const ${float_type} ${sympy.ccode(expr)}
-% endfor
-
% for i, c_i in enumerate(descriptor.c):
- preshifted_f_next[${layout.pop_offset(i)}] = f_next_${descriptor.c.index(-c_i)};
+ const ${float_type} ${assignment[i].lhs} = ${sympy.ccode(assignment[descriptor.c.index(-c_i)].rhs)};
% endfor
-}
-
-% if 'cell_list_dispatch' in extras:
-__kernel void bounce_back_boundary_cells(__global ${float_type}* f_next,
- __global ${float_type}* f_prev,
- __global unsigned int* cells)
-{
- bounce_back_boundary_gid(f_next, f_prev, cells[get_global_id(0)]);
-}
-% endif
+</%call>