From 6938758f6c1754f0ee49d0709dd0ca376a146010 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Sun, 27 Oct 2019 16:03:07 +0100 Subject: Move layout implementations into separate folder --- boltzgen/kernel/target/layout/cl.py | 39 +++++++++++++++++++++++++++++++++++++ 1 file changed, 39 insertions(+) create mode 100644 boltzgen/kernel/target/layout/cl.py (limited to 'boltzgen/kernel/target/layout/cl.py') diff --git a/boltzgen/kernel/target/layout/cl.py b/boltzgen/kernel/target/layout/cl.py new file mode 100644 index 0000000..9707cb3 --- /dev/null +++ b/boltzgen/kernel/target/layout/cl.py @@ -0,0 +1,39 @@ +class SOA: + def __init__(self, descriptor, geometry): + self.descriptor = descriptor + self.geometry = geometry + + def gid(self): + return { + 2: 'get_global_id(1)*%d + get_global_id(0)' % self.geometry.size_x, + 3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (self.geometry.size_x*self.geometry.size_y, self.geometry.size_x) + }.get(self.descriptor.d) + + def pop_offset(self, i): + return i * self.geometry.volume + + def neighbor_offset(self, c_i): + return { + 2: lambda: c_i[1]*self.geometry.size_x + c_i[0], + 3: lambda: c_i[2]*self.geometry.size_x*self.geometry.size_y + c_i[1]*self.geometry.size_x + c_i[0] + }.get(self.descriptor.d)() + +class AOS: + def __init__(self, descriptor, geometry): + self.descriptor = descriptor + self.geometry = geometry + + def gid(self): + return { + 2: '%d*(get_global_id(1)*%d + get_global_id(0))' % (self.descriptor.q, self.geometry.size_x), + 3: '%d*(get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0))' % (self.descriptor.q, self.geometry.size_x*self.geometry.size_y, self.geometry.size_x) + }.get(self.descriptor.d) + + def pop_offset(self, i): + return i + + def neighbor_offset(self, c_i): + return self.descriptor.q * { + 2: lambda: c_i[1]*self.geometry.size_x + c_i[0], + 3: lambda: c_i[2]*self.geometry.size_x*self.geometry.size_y + c_i[1]*self.geometry.size_x + c_i[0] + }.get(self.descriptor.d)() -- cgit v1.2.3