Separate functions into separate template files

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

12 files changed, 209 insertions, 218 deletions

diff --git a/boltzgen.py b/boltzgen.py
index 2b32661..5224942 100755
--- a/boltzgen.py
+++ b/boltzgen.py
@@ -15,7 +15,8 @@ argparser.add_argument('--geometry',  required = True, help = 'Size of the block
 argparser.add_argument('--tau',       required = True, help = 'BGK relaxation time')
 
 argparser.add_argument('--disable-cse', action = 'store_const', const = True, help = 'Disable common subexpression elimination')
-argparser.add_argument('--extra', action = 'append', nargs = '+', default = [], help = 'Additional generator parameters')
+argparser.add_argument('--functions',   action = 'append', nargs = '+', default = [], help = 'Function templates to be generated')
+argparser.add_argument('--extras',      action = 'append', nargs = '+', default = [], help = 'Additional generator parameters')
 
 args = argparser.parse_args()
 
@@ -29,5 +30,15 @@ generator = Generator(
 
 geometry = Geometry.parse(args.geometry)
 
-src = generator.kernel(args.language, args.precision, args.layout, geometry, sum(args.extra, []))
+functions = sum(args.functions, [])
+if len(functions) == 0:
+    functions += ['default']
+if 'default' in functions:
+    for f in ['collide_and_stream', 'equilibrilize', 'collect_moments', 'momenta_boundary']:
+        functions.insert(functions.index('default'), f)
+    functions.remove('default')
+
+extras = sum(args.extras, [])
+
+src = generator.kernel(args.language, args.precision, args.layout, geometry, functions, extras)
 print(src)
diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py
index 8d0de37..04b0a64 100644
--- a/boltzgen/kernel/generator.py
+++ b/boltzgen/kernel/generator.py
@@ -7,25 +7,15 @@ import kernel.target.cl
 import kernel.target.cpp
 
 class Generator:
-    def __init__(self, descriptor, moments, collision, boundary = ''):
+    def __init__(self, descriptor, moments, collision):
         self.descriptor = descriptor
         self.moments    = moments
         self.collision  = collision
-        self.boundary   = boundary
 
-    def kernel(self, target, precision, layout, geometry, extras = []):
-        template_path = Path(__file__).parent/('template/basic.' + target + '.mako')
+    def instantiate(self, target, template, float_type, layout_impl, geometry, extras = []):
+        template_path = Path(__file__).parent/("template/%s.%s.mako" % (template, target))
         if not template_path.exists():
-            raise Exception("Target '%s' not supported" % target)
-
-        layout_impl = eval("kernel.target.%s.%s" % (target, layout))
-        if layout_impl is None:
-            raise Exception("Target '%s' doesn't support layout '%s'" % (target, layout))
-        else:
-            layout_impl = layout_impl(self.descriptor, geometry)
-
-        if geometry.dimension() != self.descriptor.d:
-            raise Exception('Geometry dimension must match descriptor dimension')
+            raise Exception("Target '%s' doesn't provide '%s'" % (target, template))
 
         return Template(filename = str(template_path)).render(
             descriptor = self.descriptor,
@@ -38,16 +28,24 @@ class Generator:
             collision_assignment = self.collision[1],
             ccode = sympy.ccode,
 
-            float_type = {
-                'single': 'float',
-                'double': 'double'
-            }.get(precision),
-
-            boundary_src = Template(self.boundary).render(
-                descriptor = self.descriptor,
-                geometry   = geometry,
-                float_type = precision
-            ),
+            float_type = float_type,
 
             extras = extras
         )
+
+    def kernel(self, target, precision, layout, geometry, functions = ['collide_and_stream'], extras = []):
+        layout_impl = eval("kernel.target.%s.%s" % (target, layout))
+        if layout_impl is None:
+            raise Exception("Target '%s' doesn't support layout '%s'" % (target, layout))
+        else:
+            layout_impl = layout_impl(self.descriptor, geometry)
+
+        if geometry.dimension() != self.descriptor.d:
+            raise Exception('Geometry dimension must match descriptor dimension')
+
+        float_type = {
+            'single': 'float',
+            'double': 'double'
+        }.get(precision)
+
+        return "\n".join(map(lambda f: self.instantiate(target, f, float_type, layout_impl, geometry, extras), functions))
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