aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 19:11:26 +0100
committerAdrian Kummerlaender2019-10-27 19:11:26 +0100
commit4315b5c62fe862dcacbe2a8958732c8a66bdb6e1 (patch)
tree786657b96b483b4f50dea20d9b3f83c6db84bbe0
parent30f262287bd43015b155b0e882e6478cfae2780c (diff)
downloadboltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar.gz
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar.bz2
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar.lz
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar.xz
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.tar.zst
boltzgen-4315b5c62fe862dcacbe2a8958732c8a66bdb6e1.zip
Use Mako defines to generate momenta boundaries
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cl.mako46
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cpp.mako31
2 files changed, 67 insertions, 10 deletions
diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
new file mode 100644
index 0000000..50044d2
--- /dev/null
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -0,0 +1,46 @@
+<%def name="momenta_boundary(name, param)">
+__kernel void ${name}_momenta_boundary(
+ __global ${float_type}* f_next,
+ __global ${float_type}* f_prev,
+ unsigned int gid, ${param})
+{
+ __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
+
+ ${caller.body()}
+
+% 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
+}
+</%def>
+
+<%call expr="momenta_boundary('velocity', '%s%d velocity' % (float_type, descriptor.d))">
+ ${float_type} ${ccode(moments_assignment[0])}
+% for i, expr in enumerate(moments_assignment[1:]):
+ ${float_type} ${expr.lhs} = velocity.${['x', 'y', 'z'][i]};
+% endfor
+</%call>
+
+<%call expr="momenta_boundary('density', '%s density' % float_type)">
+ ${float_type} ${moments_assignment[0].lhs} = density;
+% for i, expr in enumerate(moments_assignment[1:]):
+ ${float_type} ${ccode(expr)}
+% endfor
+</%call>
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
index d7e0dd2..5475241 100644
--- a/boltzgen/kernel/template/momenta_boundary.cpp.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -1,7 +1,8 @@
-void velocity_momenta_boundary( ${float_type}* f_next,
- const ${float_type}* f_prev,
- std::size_t gid,
- ${float_type} velocity[${descriptor.d}])
+<%def name="momenta_boundary(name, param)">
+void ${name}_momenta_boundary(
+ ${float_type}* f_next,
+ const ${float_type}* f_prev,
+ std::size_t gid, ${param})
{
${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
@@ -14,12 +15,7 @@ void velocity_momenta_boundary( ${float_type}* f_next,
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
+ ${caller.body()}
% for i, expr in enumerate(collision_subexpr):
const ${float_type} ${expr[0]} = ${ccode(expr[1])};
@@ -33,3 +29,18 @@ void velocity_momenta_boundary( ${float_type}* f_next,
preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
% endfor
}
+</%def>
+
+<%call expr="momenta_boundary('velocity', '%s velocity[%d]' % (float_type, descriptor.d))">
+ ${float_type} ${ccode(moments_assignment[0])}
+% for i, expr in enumerate(moments_assignment[1:]):
+ ${float_type} ${expr.lhs} = velocity[${i}];
+% endfor
+</%call>
+
+<%call expr="momenta_boundary('density', '%s density' % float_type)">
+ ${float_type} ${moments_assignment[0].lhs} = density;
+% for i, expr in enumerate(moments_assignment[1:]):
+ ${float_type} ${ccode(expr)}
+% endfor
+</%call>