diff options
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.
Diffstat (limited to 'boltzgen/kernel/template/bounce_back_boundary.cl.mako')
-rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cl.mako | 35 |
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> |