aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-11-04 23:38:36 +0100
committerAdrian Kummerlaender2019-11-04 23:38:36 +0100
commit5828235f806c3e87a5b1eed34ef69ef317a110bd (patch)
treef4fd988f2af0e0ca2cd0e1b277f26fcf8b0f3b3a
parent05e74fb112f5b5f645b649c587d18052c7b7f9df (diff)
downloadboltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar
boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.gz
boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.bz2
boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.xz
boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.zip
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.
-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',