From a99edaaa0e9a95354f68614cd2f4ab801179b946 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Sun, 27 Oct 2019 19:50:21 +0100 Subject: Accept cell id as parameter in OpenCL functions It is more flexible to place OpenCL thread ID dependent dispatching in a separate function. --- boltzgen/kernel/template/collect_moments.cl.mako | 5 ++--- boltzgen/kernel/template/equilibrilize.cl.mako | 5 ++--- 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; -- cgit v1.2.3