From 5828235f806c3e87a5b1eed34ef69ef317a110bd Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Mon, 4 Nov 2019 23:38:36 +0100
Subject: 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.
---
 boltzgen/kernel/template/pattern/AB.cl.mako  | 69 ++++++++++++++++++++++++++++
 boltzgen/kernel/template/pattern/AB.cpp.mako | 45 ++++++++++++++++++
 2 files changed, 114 insertions(+)
 create mode 100644 boltzgen/kernel/template/pattern/AB.cl.mako
 create mode 100644 boltzgen/kernel/template/pattern/AB.cpp.mako

(limited to 'boltzgen/kernel/template/pattern')

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>
-- 
cgit v1.2.3