aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 19:50:21 +0100
committerAdrian Kummerlaender2019-10-27 19:50:21 +0100
commita99edaaa0e9a95354f68614cd2f4ab801179b946 (patch)
treea412c83652838fc09e154e25ea93bdf5df08bbc9
parentdcd162aef328cde66082a3333740ec6f58298a4c (diff)
downloadboltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar.gz
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar.bz2
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar.lz
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar.xz
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.tar.zst
boltzgen-a99edaaa0e9a95354f68614cd2f4ab801179b946.zip
Accept cell id as parameter in OpenCL functions
It is more flexible to place OpenCL thread ID dependent dispatching in a separate function.
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako5
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako5
2 files changed, 4 insertions, 6 deletions
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
index b07b759..ef520da 100644
--- a/boltzgen/kernel/template/collect_moments.cl.mako
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -1,8 +1,7 @@
__kernel void collect_moments(__global ${float_type}* f,
- __global ${float_type}* moments)
+ __global ${float_type}* moments,
+ unsigned int gid)
{
- const unsigned int gid = ${layout.gid()};
-
__global ${float_type}* preshifted_f = f + gid;
% for i in range(0,descriptor.q):
diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako
index aa2246c..0759dd5 100644
--- a/boltzgen/kernel/template/equilibrilize.cl.mako
+++ b/boltzgen/kernel/template/equilibrilize.cl.mako
@@ -1,8 +1,7 @@
__kernel void equilibrilize(__global ${float_type}* f_next,
- __global ${float_type}* f_prev)
+ __global ${float_type}* f_prev,
+ unsigned int gid)
{
- const unsigned int gid = ${layout.gid()};
-
__global ${float_type}* preshifted_f_next = f_next + gid;
__global ${float_type}* preshifted_f_prev = f_prev + gid;