From e2b00f4ec963060be98939c7b0d12d6c00e50a02 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Thu, 31 Oct 2019 13:13:00 +0100
Subject: Call symbolic generator inside code templates

This paves the way for dropping in other LBM collision models.
As a side benefit the default momenta calulcation is now fully inlined where possible.
---
 boltzgen/kernel/template/bounce_back_boundary.cl.mako  | 14 +++++---------
 boltzgen/kernel/template/bounce_back_boundary.cpp.mako | 14 +++++---------
 boltzgen/kernel/template/collect_moments.cl.mako       |  4 ++++
 boltzgen/kernel/template/collect_moments.cpp.mako      |  4 ++++
 boltzgen/kernel/template/collide_and_stream.cl.mako    | 14 +++++---------
 boltzgen/kernel/template/collide_and_stream.cpp.mako   | 16 ++++++----------
 boltzgen/kernel/template/momenta_boundary.cl.mako      |  5 +++++
 boltzgen/kernel/template/momenta_boundary.cpp.mako     |  5 +++++
 8 files changed, 39 insertions(+), 37 deletions(-)

(limited to 'boltzgen/kernel/template')

diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
index cb59cbd..8bb4f6c 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
@@ -9,19 +9,15 @@ __kernel void bounce_back_boundary_gid(__global ${float_type}* f_next,
     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
+<%
+    subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True))
+%>
 
-% for i, expr in enumerate(collision_subexpr):
+% for i, expr in enumerate(subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
 
-% for i, expr in enumerate(collision_assignment):
+% for i, expr in enumerate(assignment):
     const ${float_type} ${ccode(expr)}
 % endfor
 
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
index c7abd2a..91dcfa0 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
@@ -9,19 +9,15 @@ void bounce_back_boundary(      ${float_type}* f_next,
     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
+<%
+    subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True))
+%>
 
-% for i, expr in enumerate(collision_subexpr):
+% for i, expr in enumerate(subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
 
-% for i, expr in enumerate(collision_assignment):
+% for i, expr in enumerate(assignment):
     const ${float_type} ${ccode(expr)}
 % endfor
 
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
index ddf5102..587f58a 100644
--- a/boltzgen/kernel/template/collect_moments.cl.mako
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -9,6 +9,10 @@ __kernel void collect_moments_gid(__global ${float_type}* f,
     const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
 % endfor
 
+<%
+    moments_subexpr, moments_assignment = model.moments()
+%>
+
 % for i, expr in enumerate(moments_subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako
index d72ff5d..1605084 100644
--- a/boltzgen/kernel/template/collect_moments.cpp.mako
+++ b/boltzgen/kernel/template/collect_moments.cpp.mako
@@ -9,6 +9,10 @@ void collect_moments(const ${float_type}* f,
     const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
 % endfor
 
+<%
+    moments_subexpr, moments_assignment = model.moments()
+%>
+
 % for i, expr in enumerate(moments_subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
index c586884..303b21b 100644
--- a/boltzgen/kernel/template/collide_and_stream.cl.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -9,19 +9,15 @@ __kernel void collide_and_stream_gid(__global ${float_type}* f_next,
     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
+<%
+    subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True))
+%>
 
-% for i, expr in enumerate(collision_subexpr):
+% for i, expr in enumerate(subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
 
-% for i, expr in enumerate(collision_assignment):
+% for i, expr in enumerate(assignment):
     const ${float_type} ${ccode(expr)}
 % endfor
 
diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako
index 20fc529..47dd6dd 100644
--- a/boltzgen/kernel/template/collide_and_stream.cpp.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako
@@ -9,23 +9,19 @@ void collide_and_stream(      ${float_type}* f_next,
     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
+<%
+    subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True))
+%>
 
-% for i, expr in enumerate(collision_subexpr):
+% for i, expr in enumerate(subexpr):
     const ${float_type} ${expr[0]} = ${ccode(expr[1])};
 % endfor
 
-% for i, expr in enumerate(collision_assignment):
+% for i, expr in enumerate(assignment):
     const ${float_type} ${ccode(expr)}
 % endfor
 
-% for i, expr in enumerate(collision_assignment):
+% for i, expr in enumerate(assignment):
     preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
 % endfor
 }
diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
index f5edb61..4e29601 100644
--- a/boltzgen/kernel/template/momenta_boundary.cl.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -1,3 +1,8 @@
+<%
+moments_subexpr, moments_assignment = model.moments()
+collision_subexpr, collision_assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = False))
+%>
+
 <%def name="momenta_boundary(name, param)">
 __kernel void ${name}_momenta_boundary_gid(
    __global ${float_type}* f_next,
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
index 3f51087..6d05288 100644
--- a/boltzgen/kernel/template/momenta_boundary.cpp.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -1,3 +1,8 @@
+<%
+moments_subexpr, moments_assignment = model.moments()
+collision_subexpr, collision_assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = False))
+%>
+
 <%def name="momenta_boundary(name, param)">
 void ${name}_momenta_boundary(
           ${float_type}* f_next,
-- 
cgit v1.2.3