aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 19:40:56 +0100
committerAdrian Kummerlaender2019-10-27 19:40:56 +0100
commitdcd162aef328cde66082a3333740ec6f58298a4c (patch)
treea82322b9e0ed46d885f50bc9c558d350b22023d5
parent6e99286c9a756272e972fe21212b094de95d36f6 (diff)
downloadboltzgen-dcd162aef328cde66082a3333740ec6f58298a4c.tar
boltzgen-dcd162aef328cde66082a3333740ec6f58298a4c.tar.gz
boltzgen-dcd162aef328cde66082a3333740ec6f58298a4c.tar.bz2
boltzgen-dcd162aef328cde66082a3333740ec6f58298a4c.tar.xz
boltzgen-dcd162aef328cde66082a3333740ec6f58298a4c.zip
Add bounce back boundary condition
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako31
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cpp.mako32
2 files changed, 63 insertions, 0 deletions
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
new file mode 100644
index 0000000..7a9bc2f
--- /dev/null
+++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
@@ -0,0 +1,31 @@
+__kernel void bounce_back_boundary(__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, c_i in enumerate(descriptor.c):
+ preshifted_f_next[${layout.pop_offset(i)}] = f_next_${descriptor.c.index(-c_i)};
+% endfor
+}
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
new file mode 100644
index 0000000..06dd718
--- /dev/null
+++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
@@ -0,0 +1,32 @@
+void bounce_back_boundary( ${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, c_i in enumerate(descriptor.c):
+ preshifted_f_next[${layout.pop_offset(i)}] = f_next_${descriptor.c.index(-c_i)};
+% endfor
+}
+