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