aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/equilibrilize.cl.mako
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-11-04 23:38:36 +0100
committerAdrian Kummerlaender2019-11-04 23:38:36 +0100
commit5828235f806c3e87a5b1eed34ef69ef317a110bd (patch)
treef4fd988f2af0e0ca2cd0e1b277f26fcf8b0f3b3a /boltzgen/kernel/template/equilibrilize.cl.mako
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.
Diffstat (limited to 'boltzgen/kernel/template/equilibrilize.cl.mako')
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako22
1 files changed, 4 insertions, 18 deletions
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>