From 8f88aea2c963ebe63db83480fc9a4b48451220af Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Mon, 21 Oct 2019 22:23:41 +0200
Subject: Pull in C++ template from symlbm_playground's standalone branch

---
 boltzgen/kernel/generator.py               |   4 +-
 boltzgen/kernel/template/basic.cpp.mako    |  66 ++++++++++++++++++
 boltzgen/kernel/template/basic.opencl.mako | 104 +++++++++++++++++++++++++++++
 boltzgen/kernel/template/kernel.mako       | 104 -----------------------------
 test.py                                    |   1 +
 5 files changed, 173 insertions(+), 106 deletions(-)
 create mode 100644 boltzgen/kernel/template/basic.cpp.mako
 create mode 100644 boltzgen/kernel/template/basic.opencl.mako
 delete mode 100644 boltzgen/kernel/template/kernel.mako

diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py
index bd1bb86..59654e2 100644
--- a/boltzgen/kernel/generator.py
+++ b/boltzgen/kernel/generator.py
@@ -3,8 +3,8 @@ import sympy
 from mako.template import Template
 from pathlib import Path
 
-def source(descriptor, moments, collide, boundary_src, float_type, geometry):
-    return Template(filename = str(Path(__file__).parent/'template/kernel.mako')).render(
+def source(target, descriptor, moments, collide, boundary_src, float_type, geometry):
+    return Template(filename = str(Path(__file__).parent/('template/basic.' + target + '.mako'))).render(
         descriptor = descriptor,
         geometry   = geometry,
 
diff --git a/boltzgen/kernel/template/basic.cpp.mako b/boltzgen/kernel/template/basic.cpp.mako
new file mode 100644
index 0000000..1879c35
--- /dev/null
+++ b/boltzgen/kernel/template/basic.cpp.mako
@@ -0,0 +1,66 @@
+<%
+def pop_offset(i):
+    return i * geometry.volume
+%>
+
+void equilibrilize(${float_type}*    f_next,
+                   ${float_type}*    f_prev,
+                   const std::size_t gid)
+{
+    ${float_type}* preshifted_f_next = f_next + gid;
+    ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, w_i in enumerate(descriptor.w):
+    preshifted_f_next[${pop_offset(i)}] = ${w_i.evalf()};
+    preshifted_f_prev[${pop_offset(i)}] = ${w_i.evalf()};
+% endfor
+}
+
+<%
+def neighbor_offset(c_i):
+    return {
+        2: lambda:                                          c_i[1]*geometry.size_x + c_i[0],
+        3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
+    }.get(descriptor.d)()
+
+def padding():
+    return {
+        2: lambda:                                     1*geometry.size_x + 1,
+        3: lambda: 1*geometry.size_x*geometry.size_y + 1*geometry.size_x + 1
+    }.get(descriptor.d)()
+%>
+
+void collide_and_stream(      ${float_type}* f_next,
+                        const ${float_type}* f_prev,
+                        const int*           material,
+                        const std::size_t    gid)
+{
+    const int m = material[gid];
+
+          ${float_type}* preshifted_f_next = f_next + gid;
+    const ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, c_i in enumerate(descriptor.c):
+    const ${float_type} f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+    ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collide_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collide_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collide_assignment):
+    preshifted_f_next[${pop_offset(i)}] = m*f_next_${i} + (1.0-m)*${descriptor.w[i].evalf()};
+% endfor
+}
diff --git a/boltzgen/kernel/template/basic.opencl.mako b/boltzgen/kernel/template/basic.opencl.mako
new file mode 100644
index 0000000..5ddf64c
--- /dev/null
+++ b/boltzgen/kernel/template/basic.opencl.mako
@@ -0,0 +1,104 @@
+% if float_type == 'double':
+#if defined(cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#elif defined(cl_amd_fp64)
+#pragma OPENCL EXTENSION cl_amd_fp64 : enable
+#endif
+% endif
+
+<%
+def gid():
+    return {
+        2: 'get_global_id(1)*%d + get_global_id(0)' % geometry.size_x,
+        3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (geometry.size_x*geometry.size_y, geometry.size_x)
+    }.get(descriptor.d)
+
+def pop_offset(i):
+    return i * geometry.volume
+%>
+
+__kernel void equilibrilize(__global ${float_type}* f_next,
+                            __global ${float_type}* f_prev)
+{
+    const unsigned int gid = ${gid()};
+
+    __global ${float_type}* preshifted_f_next = f_next + gid;
+    __global ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, w_i in enumerate(descriptor.w):
+    preshifted_f_next[${pop_offset(i)}] = ${w_i}.f;
+    preshifted_f_prev[${pop_offset(i)}] = ${w_i}.f;
+% endfor
+}
+
+<%
+def neighbor_offset(c_i):
+    return {
+        2: lambda:                                          c_i[1]*geometry.size_x + c_i[0],
+        3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
+    }.get(descriptor.d)()
+
+%>
+
+__kernel void collide_and_stream(__global ${float_type}* f_next,
+                                 __global ${float_type}* f_prev,
+                                 __global int* material,
+                                 unsigned int time)
+{
+    const unsigned int gid = ${gid()};
+
+    const int m = material[gid];
+
+    if ( m == 0 ) {
+        return;
+    }
+
+    __global ${float_type}* preshifted_f_next = f_next + gid;
+    __global ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, c_i in enumerate(descriptor.c):
+    const ${float_type} f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+    ${float_type} ${ccode(expr)}
+% endfor
+
+  ${boundary_src}
+
+% for i, expr in enumerate(collide_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collide_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i in range(0,descriptor.q):
+    preshifted_f_next[${pop_offset(i)}] = f_next_${i};
+% endfor
+}
+
+__kernel void collect_moments(__global ${float_type}* f,
+                              __global ${float_type}* moments)
+{
+    const unsigned int gid = ${gid()};
+
+    __global ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+    moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
+% endfor
+}
diff --git a/boltzgen/kernel/template/kernel.mako b/boltzgen/kernel/template/kernel.mako
deleted file mode 100644
index 5ddf64c..0000000
--- a/boltzgen/kernel/template/kernel.mako
+++ /dev/null
@@ -1,104 +0,0 @@
-% if float_type == 'double':
-#if defined(cl_khr_fp64)
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-#elif defined(cl_amd_fp64)
-#pragma OPENCL EXTENSION cl_amd_fp64 : enable
-#endif
-% endif
-
-<%
-def gid():
-    return {
-        2: 'get_global_id(1)*%d + get_global_id(0)' % geometry.size_x,
-        3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (geometry.size_x*geometry.size_y, geometry.size_x)
-    }.get(descriptor.d)
-
-def pop_offset(i):
-    return i * geometry.volume
-%>
-
-__kernel void equilibrilize(__global ${float_type}* f_next,
-                            __global ${float_type}* f_prev)
-{
-    const unsigned int gid = ${gid()};
-
-    __global ${float_type}* preshifted_f_next = f_next + gid;
-    __global ${float_type}* preshifted_f_prev = f_prev + gid;
-
-% for i, w_i in enumerate(descriptor.w):
-    preshifted_f_next[${pop_offset(i)}] = ${w_i}.f;
-    preshifted_f_prev[${pop_offset(i)}] = ${w_i}.f;
-% endfor
-}
-
-<%
-def neighbor_offset(c_i):
-    return {
-        2: lambda:                                          c_i[1]*geometry.size_x + c_i[0],
-        3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
-    }.get(descriptor.d)()
-
-%>
-
-__kernel void collide_and_stream(__global ${float_type}* f_next,
-                                 __global ${float_type}* f_prev,
-                                 __global int* material,
-                                 unsigned int time)
-{
-    const unsigned int gid = ${gid()};
-
-    const int m = material[gid];
-
-    if ( m == 0 ) {
-        return;
-    }
-
-    __global ${float_type}* preshifted_f_next = f_next + gid;
-    __global ${float_type}* preshifted_f_prev = f_prev + gid;
-
-% for i, c_i in enumerate(descriptor.c):
-    const ${float_type} f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
-    ${float_type} ${ccode(expr)}
-% endfor
-
-  ${boundary_src}
-
-% for i, expr in enumerate(collide_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collide_assignment):
-    const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i in range(0,descriptor.q):
-    preshifted_f_next[${pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-__kernel void collect_moments(__global ${float_type}* f,
-                              __global ${float_type}* moments)
-{
-    const unsigned int gid = ${gid()};
-
-    __global ${float_type}* preshifted_f = f + gid;
-
-% for i in range(0,descriptor.q):
-    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
-    moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
-% endfor
-}
diff --git a/test.py b/test.py
index b5831aa..0eec050 100644
--- a/test.py
+++ b/test.py
@@ -4,6 +4,7 @@ lbm = LBM(D2Q9)
 geometry = Geometry(32,32)
 
 src = source(
+    'opencl',
     D2Q9,
     lbm.moments(),
     lbm.bgk(f_eq = lbm.equilibrium(), tau = 0.6),
-- 
cgit v1.2.3