aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--boltzgen/kernel/generator.py12
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako35
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cpp.mako27
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako26
-rw-r--r--boltzgen/kernel/template/collect_moments.cpp.mako21
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako33
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cpp.mako24
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako22
-rw-r--r--boltzgen/kernel/template/equilibrilize.cpp.mako14
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cl.mako47
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cpp.mako27
-rw-r--r--boltzgen/kernel/template/pattern/AB.cl.mako69
-rw-r--r--boltzgen/kernel/template/pattern/AB.cpp.mako45
-rw-r--r--setup.py2
14 files changed, 176 insertions, 228 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>
diff --git a/setup.py b/setup.py
index af00a78..f386447 100644
--- a/setup.py
+++ b/setup.py
@@ -9,7 +9,7 @@ setup(
author = 'Adrian Kummerlaender',
packages = find_packages(),
include_package_data = True,
- package_data = {'boltzgen': ['kernel/template/*.mako']},
+ package_data = {'boltzgen': ['kernel/template/*.mako', 'kernel/template/pattern/*.mako']},
install_requires = [
'sympy >= 1.4',
'numpy >= 1.17.2',