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