aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/collide_and_stream.cl.mako
diff options
context:
space:
mode:
Diffstat (limited to 'boltzgen/kernel/template/collide_and_stream.cl.mako')
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako33
1 files changed, 4 insertions, 29 deletions
diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
index bfc9435..e0c1c7f 100644
--- a/boltzgen/kernel/template/collide_and_stream.cl.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -1,22 +1,10 @@
+<%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 collide_and_stream_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('collide_and_stream')">
% for i, expr in enumerate(subexpr):
const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
@@ -24,17 +12,4 @@ __kernel void collide_and_stream_gid(__global ${float_type}* f_next,
% for i, expr in enumerate(assignment):
const ${float_type} ${sympy.ccode(expr)}
% endfor
-
-% for i in range(0,descriptor.q):
- preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-% if 'cell_list_dispatch' in extras:
-__kernel void collide_and_stream_cells(__global ${float_type}* f_next,
- __global ${float_type}* f_prev,
- __global unsigned int* cells)
-{
- collide_and_stream_gid(f_next, f_prev, cells[get_global_id(0)]);
-}
-% endif
+</%call>