Separate functions into separate template files

Selection of the desired templates is possible via a new `functions` parameter.

10 files changed, 174 insertions, 192 deletions

diff --git a/boltzgen/kernel/template/basic.cl.mako b/boltzgen/kernel/template/basic.cl.mako
deleted file mode 100644
index b64a480..0000000
--- a/boltzgen/kernel/template/basic.cl.mako
+++ /dev/null
@@ -1,84 +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
-
-__kernel void equilibrilize(__global ${float_type}* f_next,
-                            __global ${float_type}* f_prev)
-{
-    const unsigned int gid = ${layout.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[${layout.pop_offset(i)}] = ${w_i}.f;
-    preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f;
-% endfor
-}
-
-__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 = ${layout.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[${layout.pop_offset(i) + layout.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(collision_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
-    const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i in range(0,descriptor.q):
-    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-__kernel void collect_moments(__global ${float_type}* f,
-                              __global ${float_type}* moments)
-{
-    const unsigned int gid = ${layout.gid()};
-
-    __global ${float_type}* preshifted_f = f + gid;
-
-% for i in range(0,descriptor.q):
-    const ${float_type} f_curr_${i} = preshifted_f[${layout.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[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
-% endfor
-}
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
new file mode 100644
index 0000000..b07b759
--- /dev/null
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -0,0 +1,19 @@
+__kernel void collect_moments(__global ${float_type}* f,
+                              __global ${float_type}* moments)
+{
+    const unsigned int gid = ${layout.gid()};
+
+    __global ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+    const ${float_type} f_curr_${i} = preshifted_f[${layout.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[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
+% endfor
+}
diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako
new file mode 100644
index 0000000..8c37db2
--- /dev/null
+++ b/boltzgen/kernel/template/collect_moments.cpp.mako
@@ -0,0 +1,24 @@
+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*${layout.gid_offset()};
+
+% for i in range(0,descriptor.q):
+    const ${float_type} f_curr_${i} = preshifted_f[${layout.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
+}
+
diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
new file mode 100644
index 0000000..28cfa57
--- /dev/null
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -0,0 +1,31 @@
+__kernel void collide_and_stream(__global ${float_type}* f_next,
+                                 __global ${float_type}* f_prev,
+                                 unsigned int gid)
+{
+    __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[${layout.pop_offset(i) + layout.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(collision_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i in range(0,descriptor.q):
+    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako
new file mode 100644
index 0000000..71c62b7
--- /dev/null
+++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako
@@ -0,0 +1,32 @@
+void collide_and_stream(      ${float_type}* f_next,
+                        const ${float_type}* f_prev,
+                        std::size_t gid)
+{
+          ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+    const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% 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
+
+% 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(collision_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
+
diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako
new file mode 100644
index 0000000..aa2246c
--- /dev/null
+++ b/boltzgen/kernel/template/equilibrilize.cl.mako
@@ -0,0 +1,13 @@
+__kernel void equilibrilize(__global ${float_type}* f_next,
+                            __global ${float_type}* f_prev)
+{
+    const unsigned int gid = ${layout.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[${layout.pop_offset(i)}] = ${w_i}.f;
+    preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f;
+% endfor
+}
diff --git a/boltzgen/kernel/template/equilibrilize.cpp.mako b/boltzgen/kernel/template/equilibrilize.cpp.mako
new file mode 100644
index 0000000..9f082a1
--- /dev/null
+++ b/boltzgen/kernel/template/equilibrilize.cpp.mako
@@ -0,0 +1,13 @@
+void equilibrilize(${float_type}* f_next,
+                   ${float_type}* f_prev,
+                   std::size_t gid)
+{
+    ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+    ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% 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()};
+% endfor
+}
+
diff --git a/boltzgen/kernel/template/basic.cpp.mako b/boltzgen/kernel/template/example.cpp.mako
index 118ef8c..d979bfd 100644
--- a/boltzgen/kernel/template/basic.cpp.mako
+++ b/boltzgen/kernel/template/example.cpp.mako
@@ -1,108 +1,3 @@
-void equilibrilize(${float_type}* f_next,
-                   ${float_type}* f_prev,
-                   std::size_t gid)
-{
-    ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
-    ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% 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()};
-% endfor
-}
-
-void collide_and_stream(      ${float_type}* f_next,
-                        const ${float_type}* f_prev,
-                        std::size_t gid)
-{
-          ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
-    const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% 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
-
-% 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(collision_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
-    const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
-    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-void velocity_momenta_boundary(      ${float_type}* f_next,
-                               const ${float_type}* f_prev,
-                               std::size_t gid,
-                               ${float_type} velocity[${descriptor.d}])
-{
-          ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
-    const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% 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
-
-% 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(moments_assignment[1:]):
-    ${expr.lhs} = velocity[${i}];
-% endfor
-
-% for i, expr in enumerate(collision_subexpr):
-    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
-    const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
-    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% 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*${layout.gid_offset()};
-
-% for i in range(0,descriptor.q):
-    const ${float_type} f_curr_${i} = preshifted_f[${layout.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
-}
-
 % if 'moments_vtk' in extras:
 void collect_moments_to_vtk(const std::string& path, ${float_type}* f) {
     std::ofstream fout;
@@ -186,7 +81,6 @@ void collect_moments_to_vtk(const std::string& path, ${float_type}* f) {
 }
 % endif
 
-% if 'test_ldc' in extras:
 void test_ldc(std::size_t nStep)
 {
     auto f_a = std::make_unique<${float_type}[]>(${geometry.volume*descriptor.q});
@@ -276,5 +170,3 @@ void test_ldc(std::size_t nStep)
     collect_moments_to_vtk("test.vtk", f_next);
 % endif
 }
-% endif
-
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
new file mode 100644
index 0000000..d7e0dd2
--- /dev/null
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -0,0 +1,35 @@
+void velocity_momenta_boundary(      ${float_type}* f_next,
+                               const ${float_type}* f_prev,
+                               std::size_t gid,
+                               ${float_type} velocity[${descriptor.d}])
+{
+          ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+    const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% 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
+
+% 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(moments_assignment[1:]):
+    ${expr.lhs} = velocity[${i}];
+% endfor
+
+% for i, expr in enumerate(collision_subexpr):
+    const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
diff --git a/boltzgen/kernel/template/preamble.cl.mako b/boltzgen/kernel/template/preamble.cl.mako
new file mode 100644
index 0000000..c47254d
--- /dev/null
+++ b/boltzgen/kernel/template/preamble.cl.mako
@@ -0,0 +1,7 @@
+% 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