aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/pattern/AB.cl.mako
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-11-05 23:33:47 +0100
committerAdrian Kummerlaender2019-11-05 23:34:14 +0100
commit62e3d5708470415b9ea2f0a737acaf4e2d00bb21 (patch)
tree69f4bb0156cade8dfaa4ec0b71db00283cd8ebb7 /boltzgen/kernel/template/pattern/AB.cl.mako
parent019e8d57c6266ce0b26d8eacab984f303442a184 (diff)
downloadboltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar.gz
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar.bz2
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar.lz
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar.xz
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.tar.zst
boltzgen-62e3d5708470415b9ea2f0a737acaf4e2d00bb21.zip
Implement AA pattern for OpenCL target
Works well but function naming is getting kind of clunky, e.g. "velocity_momenta_boundary_tick_cells" This could be hidden to a degree by proving branching wrappers for the odd and even time step implementations. However this would not vectorize when targeting Intel via OpenCL.
Diffstat (limited to 'boltzgen/kernel/template/pattern/AB.cl.mako')
-rw-r--r--boltzgen/kernel/template/pattern/AB.cl.mako22
1 files changed, 22 insertions, 0 deletions
diff --git a/boltzgen/kernel/template/pattern/AB.cl.mako b/boltzgen/kernel/template/pattern/AB.cl.mako
index d2d8b8a..33836ea 100644
--- a/boltzgen/kernel/template/pattern/AB.cl.mako
+++ b/boltzgen/kernel/template/pattern/AB.cl.mako
@@ -66,4 +66,26 @@ __kernel void ${name}(
${caller.body()}
}
+
+% if 'cell_list_dispatch' in extras:
+__kernel void ${name}_cells(
+ __global ${float_type}* f
+ , __global unsigned int* cells
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_type} ${param_name}
+% endfor
+% endif
+) {
+ ${name}(
+ f
+ , cells[get_global_id(0)]
+% if params is not None:
+% for param_type, param_name in params:
+ , ${param_name}
+% endfor
+% endif
+ );
+}
+% endif
</%def>