diff options
Diffstat (limited to 'boltzgen/kernel/template/momenta_boundary.cl.mako')
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cl.mako | 47 |
1 files changed, 7 insertions, 40 deletions
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 |