diff options
Diffstat (limited to 'boltzgen')
| -rw-r--r-- | boltzgen/kernel/generator.py | 12 | ||||
| -rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cl.mako | 35 | ||||
| -rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cpp.mako | 27 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collect_moments.cl.mako | 26 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collect_moments.cpp.mako | 21 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collide_and_stream.cl.mako | 33 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collide_and_stream.cpp.mako | 24 | ||||
| -rw-r--r-- | boltzgen/kernel/template/equilibrilize.cl.mako | 22 | ||||
| -rw-r--r-- | boltzgen/kernel/template/equilibrilize.cpp.mako | 14 | ||||
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cl.mako | 47 | ||||
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cpp.mako | 27 | ||||
| -rw-r--r-- | boltzgen/kernel/template/pattern/AB.cl.mako | 69 | ||||
| -rw-r--r-- | boltzgen/kernel/template/pattern/AB.cpp.mako | 45 |
13 files changed, 175 insertions, 227 deletions
diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py index 8da91ba..dd44a56 100644 --- a/boltzgen/kernel/generator.py +++ b/boltzgen/kernel/generator.py @@ -1,8 +1,14 @@ from mako.template import Template +from mako.lookup import TemplateLookup + from pathlib import Path from . import memory +template_lookup = TemplateLookup(directories = [ + Path(__file__).parent/"template" +]) + class Generator: def __init__(self, model, target, precision, index, layout): self.model = model @@ -25,12 +31,13 @@ class Generator: if not template_path.exists(): raise Exception("Target '%s' doesn't provide '%s'" % (self.target, template)) - return Template(filename = str(template_path)).render( + return Template(filename = str(template_path), lookup = template_lookup).render( descriptor = self.descriptor, model = self.model, geometry = geometry, index = self.index_impl(geometry), layout = self.layout_impl(self.descriptor, self.index_impl, geometry), + streaming = 'AB', float_type = self.float_type, extras = extras ) @@ -42,12 +49,13 @@ class Generator: return "\n".join(map(lambda f: self.instantiate(f, geometry, extras), functions)) def custom(self, geometry, source, extras = []): - return Template(text = source).render( + return Template(text = source, lookup = template_lookup).render( descriptor = self.descriptor, model = self.model, geometry = geometry, index = self.index_impl(geometry), layout = self.layout_impl(self.descriptor, self.index_impl, geometry), + streaming = 'AB', float_type = self.float_type, extras = extras ) 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> diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako index 1dcafe8..a53184e 100644 --- a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako +++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako @@ -1,32 +1,15 @@ +<%namespace name="pattern" file="${'/pattern/%s.cpp.mako' % context['streaming']}"/> <% import sympy +subexpr, assignment = model.collision(f_eq = model.equilibrium(resolve_moments = True)) %> -void bounce_back_boundary( ${float_type}* f_next, - const ${float_type}* f_prev, - std::size_t gid) -{ - ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; - const ${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 -} - +</%call> diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako index 39317e3..8adf295 100644 --- a/boltzgen/kernel/template/collect_moments.cl.mako +++ b/boltzgen/kernel/template/collect_moments.cl.mako @@ -1,36 +1,26 @@ +<%namespace name="pattern" file="${'/pattern/%s.cl.mako' % context['streaming']}"/> <% import sympy +moments_subexpr, moments_assignment = model.moments() %> -__kernel void collect_moments_gid(__global ${float_type}* f, - __global ${float_type}* m, - unsigned int gid) -{ - __global ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')}; - __global ${float_type}* preshifted_m = m + gid*${descriptor.d+1}; - -% for i in range(0,descriptor.q): - const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; -% endfor - -<% - moments_subexpr, moments_assignment = model.moments() -%> - +<%call expr="pattern.functor_ab('collect_moments', [('__global %s*' % float_type, 'm')])"> % for i, expr in enumerate(moments_subexpr): const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])}; % endfor + __global ${float_type}* preshifted_m = m + gid*${descriptor.d+1}; + % for i, expr in enumerate(moments_assignment): preshifted_m[${i}] = ${sympy.ccode(expr.rhs)}; % endfor -} +</%call> % if 'cell_list_dispatch' in extras: __kernel void collect_moments_cells(__global ${float_type}* f, - __global ${float_type}* moments, + __global ${float_type}* m, __global unsigned int* cells) { - collect_moments_gid(f, moments, cells[get_global_id(0)]); + collect_moments(f, cells[get_global_id(0)], m); } % endif diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako index 493c53d..570cf92 100644 --- a/boltzgen/kernel/template/collect_moments.cpp.mako +++ b/boltzgen/kernel/template/collect_moments.cpp.mako @@ -1,22 +1,10 @@ +<%namespace name="pattern" file="${'/pattern/%s.cpp.mako' % context['streaming']}"/> <% import sympy +moments_subexpr, moments_assignment = model.moments() %> -void collect_moments(const ${float_type}* f, - std::size_t gid, - ${float_type}& rho, - ${float_type} u[${descriptor.d}]) -{ - const ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')}; - -% for i in range(0,descriptor.q): - const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; -% endfor - -<% - moments_subexpr, moments_assignment = model.moments() -%> - +<%call expr="pattern.functor_ab('collect_moments', [('%s&' % float_type, 'rho'), (float_type, 'u[%d]' % descriptor.d)])"> % for i, expr in enumerate(moments_subexpr): const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])}; % endfor @@ -28,5 +16,4 @@ void collect_moments(const ${float_type}* f, u[${i-1}] = ${sympy.ccode(expr.rhs)}; % endif % endfor -} - +</%call> 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> diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako index ee42eb7..dec009f 100644 --- a/boltzgen/kernel/template/collide_and_stream.cpp.mako +++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako @@ -1,22 +1,10 @@ +<%namespace name="pattern" file="${'/pattern/%s.cpp.mako' % context['streaming']}"/> <% import sympy +subexpr, assignment = model.collision(f_eq = model.equilibrium(resolve_moments = True)) %> -void collide_and_stream( ${float_type}* f_next, - const ${float_type}* f_prev, - std::size_t gid) -{ - ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; - const ${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,8 +12,4 @@ void collide_and_stream( ${float_type}* f_next, % for i, expr in enumerate(assignment): const ${float_type} ${sympy.ccode(expr)} % endfor - -% for i, expr in enumerate(assignment): - preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i}; -% endfor -} +</%call> diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako index 4ee8d41..ba01712 100644 --- a/boltzgen/kernel/template/equilibrilize.cl.mako +++ b/boltzgen/kernel/template/equilibrilize.cl.mako @@ -1,21 +1,7 @@ -__kernel void equilibrilize_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')}; +<%namespace name="pattern" file="${'/pattern/%s.cl.mako' % context['streaming']}"/> +<%call expr="pattern.operator_ab('equilibrilize')"> % for i, w_i in enumerate(descriptor.w): - preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f; - preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f; + ${float_type} f_next_${i} = ${w_i.evalf()}; % endfor -} - -% if 'cell_list_dispatch' in extras: -__kernel void equilibrilize_cells(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - __global unsigned int* cells) -{ - equilibrilize_gid(f_next, f_prev, cells[get_global_id(0)]); -} -% endif +</%call> diff --git a/boltzgen/kernel/template/equilibrilize.cpp.mako b/boltzgen/kernel/template/equilibrilize.cpp.mako index 3b95a31..b7d6cde 100644 --- a/boltzgen/kernel/template/equilibrilize.cpp.mako +++ b/boltzgen/kernel/template/equilibrilize.cpp.mako @@ -1,13 +1,7 @@ -void equilibrilize(${float_type}* f_next, - ${float_type}* f_prev, - std::size_t gid) -{ - ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; - ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')}; +<%namespace name="pattern" file="${'/pattern/%s.cpp.mako' % context['streaming']}"/> +<%call expr="pattern.operator_ab('equilibrilize')"> % for i, w_i in enumerate(descriptor.w): - preshifted_f_next[${layout.pop_offset(i)}] = ${w_i.evalf()}; - preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i.evalf()}; + ${float_type} f_next_${i} = ${w_i.evalf()}; % endfor -} - +</%call> 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 diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako index ae78e9f..bea6cc1 100644 --- a/boltzgen/kernel/template/momenta_boundary.cpp.mako +++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako @@ -1,23 +1,12 @@ +<%namespace name="pattern" file="${'/pattern/%s.cpp.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)"> -void ${name}_momenta_boundary( - ${float_type}* f_next, - const ${float_type}* f_prev, - std::size_t gid, ${param}) -{ - ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; - const ${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,21 +20,17 @@ void ${name}_momenta_boundary( % 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 velocity[%d]' % (float_type, descriptor.d))"> +<%call expr="momenta_boundary('velocity', [(float_type, 'velocity[%d]' % descriptor.d)])"> ${float_type} ${sympy.ccode(moments_assignment[0])} % for i, expr in enumerate(moments_assignment[1:]): ${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)} diff --git a/boltzgen/kernel/template/pattern/AB.cl.mako b/boltzgen/kernel/template/pattern/AB.cl.mako new file mode 100644 index 0000000..c6391fa --- /dev/null +++ b/boltzgen/kernel/template/pattern/AB.cl.mako @@ -0,0 +1,69 @@ +<%def name="operator_ab(name, params = None)"> +__kernel void ${name}( + __global ${float_type}* f_next + , __global ${float_type}* f_prev + , unsigned int gid +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + __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 + + ${caller.body()} + +% for i, _ in enumerate(descriptor.c): + preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i}; +% endfor +} + +% if 'cell_list_dispatch' in extras: +__kernel void ${name}_cells( + __global ${float_type}* f_next + , __global ${float_type}* f_prev + , __global unsigned int* cells +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + ${name}( + f_next + , f_prev + , cells[get_global_id(0)] +% if params is not None: +% for param_type, param_name in params: + , ${param_name} +% endfor +% endif + ); +} +% endif +</%def> + +<%def name="functor_ab(name, params = None)"> +__kernel void ${name}( + __global ${float_type}* f + , unsigned int gid +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + __global ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')}; + +% for i in range(0,descriptor.q): + const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; +% endfor + + ${caller.body()} +} +</%def> diff --git a/boltzgen/kernel/template/pattern/AB.cpp.mako b/boltzgen/kernel/template/pattern/AB.cpp.mako new file mode 100644 index 0000000..ced8524 --- /dev/null +++ b/boltzgen/kernel/template/pattern/AB.cpp.mako @@ -0,0 +1,45 @@ +<%def name="operator_ab(name, params = None)"> +void ${name}( + ${float_type}* f_next + , const ${float_type}* f_prev + , std::size_t gid +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; + const ${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 + + ${caller.body()} + +% for i, _ in enumerate(descriptor.c): + preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i}; +% endfor +} +</%def> + +<%def name="functor_ab(name, params = None)"> +void ${name}( + const ${float_type}* f + , std::size_t gid +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + const ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')}; + +% for i in range(0,descriptor.q): + const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; +% endfor + + ${caller.body()} +} +</%def> |