aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/collect_moments.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/collect_moments.cl.mako
parent1b9ac6e7aee3cf63495a65c2d7dbf79a0be23d7d (diff)
downloadboltzgen-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/collect_moments.cl.mako')
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako15
1 files changed, 12 insertions, 3 deletions
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
index ef520da..0ab42d1 100644
--- a/boltzgen/kernel/template/collect_moments.cl.mako
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -1,6 +1,6 @@
-__kernel void collect_moments(__global ${float_type}* f,
- __global ${float_type}* moments,
- unsigned int gid)
+__kernel void collect_moments_gid(__global ${float_type}* f,
+ __global ${float_type}* moments,
+ unsigned int gid)
{
__global ${float_type}* preshifted_f = f + gid;
@@ -16,3 +16,12 @@ __kernel void collect_moments(__global ${float_type}* f,
moments[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
% endfor
}
+
+% if 'cell_list_dispatch' in extras:
+__kernel void collect_moments_cells(__global ${float_type}* f,
+ __global ${float_type}* moments,
+ __global unsigned int* cells)
+{
+ collect_moments_gid(f, moments, cells[get_global_id(0)]);
+}
+% endif