diff options
author | Adrian Kummerlaender | 2019-11-05 23:33:47 +0100 |
---|---|---|
committer | Adrian Kummerlaender | 2019-11-05 23:34:14 +0100 |
commit | 62e3d5708470415b9ea2f0a737acaf4e2d00bb21 (patch) | |
tree | 69f4bb0156cade8dfaa4ec0b71db00283cd8ebb7 /boltzgen/kernel/template/pattern/AB.cl.mako | |
parent | 019e8d57c6266ce0b26d8eacab984f303442a184 (diff) | |
download | boltzgen-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.mako | 22 |
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> |