aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/equilibrilize.cl.mako
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 22:22:24 +0100
committerAdrian Kummerlaender2019-10-27 22:22:24 +0100
commit536b25e2c2b742c17d023d2b3386fed4dc60a339 (patch)
tree7424a6431d8076257125a70adcd45bf417aa38c1 /boltzgen/kernel/template/equilibrilize.cl.mako
parent1b9ac6e7aee3cf63495a65c2d7dbf79a0be23d7d (diff)
downloadboltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.gz
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.bz2
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.xz
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.mako15
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