From dbd9a340a7809a770d52d10154712278431acdc3 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Tue, 29 Oct 2019 10:56:41 +0100 Subject: Unify AOS, SOA specific cell preshift between targets SOA and AOS should not be target specific, neighbor offset calculation / bijection between gid and cell coordinates should be customizable. --- boltzgen/kernel/template/collide_and_stream.cl.mako | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'boltzgen/kernel/template/collide_and_stream.cl.mako') diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako index a8fe532..c586884 100644 --- a/boltzgen/kernel/template/collide_and_stream.cl.mako +++ b/boltzgen/kernel/template/collide_and_stream.cl.mako @@ -2,8 +2,8 @@ __kernel void collide_and_stream_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; + __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; + __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')}; % for i, c_i in enumerate(descriptor.c): const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}]; -- cgit v1.2.3