Some cleanup, add `collect_moments` to C++ template

2 files changed, 53 insertions, 34 deletions

diff --git a/boltzgen/kernel/template/basic.opencl.mako b/boltzgen/kernel/template/basic.cl.mako
index 3ecb16c..1b02c63 100644
--- a/boltzgen/kernel/template/basic.opencl.mako
+++ b/boltzgen/kernel/template/basic.cl.mako
@@ -1,11 +1,3 @@
-% 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 {
@@ -15,8 +7,22 @@ def gid():
 
 def pop_offset(i):
     return i * geometry.volume
+
+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)()
 %>
 
+% 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
+
 __kernel void equilibrilize(__global ${float_type}* f_next,
                             __global ${float_type}* f_prev)
 {
@@ -31,15 +37,6 @@ __kernel void equilibrilize(__global ${float_type}* f_next,
 % 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,
diff --git a/boltzgen/kernel/template/basic.cpp.mako b/boltzgen/kernel/template/basic.cpp.mako
index ef3486f..d284a1c 100644
--- a/boltzgen/kernel/template/basic.cpp.mako
+++ b/boltzgen/kernel/template/basic.cpp.mako
@@ -1,22 +1,7 @@
 <%
 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],
@@ -30,10 +15,23 @@ def padding():
     }.get(descriptor.d)()
 %>
 
+void equilibrilize(${float_type}* f_next,
+                   ${float_type}* f_prev,
+                   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
+}
+
 void collide_and_stream(      ${float_type}* f_next,
                         const ${float_type}* f_prev,
-                        const int*           material,
-                        const std::size_t    gid)
+                        const int*  material,
+                        std::size_t gid)
 {
     const int m = material[gid];
 
@@ -64,3 +62,27 @@ void collide_and_stream(      ${float_type}* f_next,
     preshifted_f_next[${pop_offset(i)}] = m*f_next_${i} + (1.0-m)*${descriptor.w[i].evalf()};
 % endfor
 }
+
+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 + 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):
+%   if i == 0:
+    rho = ${ccode(expr.rhs)};
+%   else:
+    u[${i-1}] = ${ccode(expr.rhs)};
+%   endif
+% endfor
+}