diff options
author | Adrian Kummerlaender | 2019-10-27 22:22:24 +0100 |
---|---|---|
committer | Adrian Kummerlaender | 2019-10-27 22:22:24 +0100 |
commit | 536b25e2c2b742c17d023d2b3386fed4dc60a339 (patch) | |
tree | 7424a6431d8076257125a70adcd45bf417aa38c1 /boltzgen/kernel/template/equilibrilize.cl.mako | |
parent | 1b9ac6e7aee3cf63495a65c2d7dbf79a0be23d7d (diff) | |
download | boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.gz boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.bz2 boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.lz boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.xz boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.zst boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.zip |
Optionally generate cell-list-based OpenCL dispatch functions
Requires different function naming as OpenCL 1.2 doesn't support overloads.
The OpenCL kernel code generated using this commit was successfully tested
on an actual GPU. Time to set up some automatic validation.
Diffstat (limited to 'boltzgen/kernel/template/equilibrilize.cl.mako')
-rw-r--r-- | boltzgen/kernel/template/equilibrilize.cl.mako | 15 |
1 files changed, 12 insertions, 3 deletions
diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako index 0759dd5..4b9b984 100644 --- a/boltzgen/kernel/template/equilibrilize.cl.mako +++ b/boltzgen/kernel/template/equilibrilize.cl.mako @@ -1,6 +1,6 @@ -__kernel void equilibrilize(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - unsigned int gid) +__kernel void equilibrilize_gid(__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; @@ -10,3 +10,12 @@ __kernel void equilibrilize(__global ${float_type}* f_next, preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f; % endfor } + +% if 'cell_list_dispatch' in extras: +__kernel void equilibrilize_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells) +{ + equilibrilize_gid(f_next, f_prev, cells[get_global_id(0)]); +} +% endif |