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/generator.py | 7 +++--- boltzgen/kernel/target/cl.py | 39 ------------------------------- boltzgen/kernel/target/cpp.py | 33 -------------------------- boltzgen/kernel/target/layout/__init__.py | 2 ++ boltzgen/kernel/target/layout/cl.py | 39 +++++++++++++++++++++++++++++++ boltzgen/kernel/target/layout/cpp.py | 33 ++++++++++++++++++++++++++ 6 files changed, 77 insertions(+), 76 deletions(-) delete mode 100644 boltzgen/kernel/target/cl.py delete mode 100644 boltzgen/kernel/target/cpp.py create mode 100644 boltzgen/kernel/target/layout/__init__.py create mode 100644 boltzgen/kernel/target/layout/cl.py create mode 100644 boltzgen/kernel/target/layout/cpp.py diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py index 04b0a64..ebe1018 100644 --- a/boltzgen/kernel/generator.py +++ b/boltzgen/kernel/generator.py @@ -3,8 +3,7 @@ import sympy from mako.template import Template from pathlib import Path -import kernel.target.cl -import kernel.target.cpp +import kernel.target.layout class Generator: def __init__(self, descriptor, moments, collision): @@ -33,8 +32,8 @@ class Generator: extras = extras ) - def kernel(self, target, precision, layout, geometry, functions = ['collide_and_stream'], extras = []): - layout_impl = eval("kernel.target.%s.%s" % (target, layout)) + def kernel(self, target, precision, layout, geometry, functions, extras = []): + layout_impl = eval("kernel.target.layout.%s.%s" % (target, layout)) if layout_impl is None: raise Exception("Target '%s' doesn't support layout '%s'" % (target, layout)) else: diff --git a/boltzgen/kernel/target/cl.py b/boltzgen/kernel/target/cl.py deleted file mode 100644 index 9707cb3..0000000 --- a/boltzgen/kernel/target/cl.py +++ /dev/null @@ -1,39 +0,0 @@ -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)() diff --git a/boltzgen/kernel/target/cpp.py b/boltzgen/kernel/target/cpp.py deleted file mode 100644 index bb5dba0..0000000 --- a/boltzgen/kernel/target/cpp.py +++ /dev/null @@ -1,33 +0,0 @@ -class AOS: - def __init__(self, descriptor, geometry): - self.descriptor = descriptor - self.geometry = geometry - - def gid_offset(self): - return self.descriptor.q - - def pop_offset(self, i): - return i - - def neighbor_offset(self, c_i): - return self.descriptor.q * { - 2: lambda: c_i[0]*self.geometry.size_y + c_i[1], - 3: lambda: c_i[0]*self.geometry.size_y*self.geometry.size_z + c_i[1]*self.geometry.size_z + c_i[2] - }.get(self.descriptor.d)() - -class SOA: - def __init__(self, descriptor, geometry): - self.descriptor = descriptor - self.geometry = geometry - - def gid_offset(self): - return 1 - - def pop_offset(self, i): - return i * self.geometry.volume - - def neighbor_offset(self, c_i): - return { - 2: lambda: c_i[0]*self.geometry.size_y + c_i[1], - 3: lambda: c_i[0]*self.geometry.size_y*self.geometry.size_z + c_i[1]*self.geometry.size_z + c_i[2] - }.get(self.descriptor.d)() diff --git a/boltzgen/kernel/target/layout/__init__.py b/boltzgen/kernel/target/layout/__init__.py new file mode 100644 index 0000000..3887e07 --- /dev/null +++ b/boltzgen/kernel/target/layout/__init__.py @@ -0,0 +1,2 @@ +import kernel.target.layout.cl +import kernel.target.layout.cpp 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)() diff --git a/boltzgen/kernel/target/layout/cpp.py b/boltzgen/kernel/target/layout/cpp.py new file mode 100644 index 0000000..bb5dba0 --- /dev/null +++ b/boltzgen/kernel/target/layout/cpp.py @@ -0,0 +1,33 @@ +class AOS: + def __init__(self, descriptor, geometry): + self.descriptor = descriptor + self.geometry = geometry + + def gid_offset(self): + return self.descriptor.q + + def pop_offset(self, i): + return i + + def neighbor_offset(self, c_i): + return self.descriptor.q * { + 2: lambda: c_i[0]*self.geometry.size_y + c_i[1], + 3: lambda: c_i[0]*self.geometry.size_y*self.geometry.size_z + c_i[1]*self.geometry.size_z + c_i[2] + }.get(self.descriptor.d)() + +class SOA: + def __init__(self, descriptor, geometry): + self.descriptor = descriptor + self.geometry = geometry + + def gid_offset(self): + return 1 + + def pop_offset(self, i): + return i * self.geometry.volume + + def neighbor_offset(self, c_i): + return { + 2: lambda: c_i[0]*self.geometry.size_y + c_i[1], + 3: lambda: c_i[0]*self.geometry.size_y*self.geometry.size_z + c_i[1]*self.geometry.size_z + c_i[2] + }.get(self.descriptor.d)() -- cgit v1.2.3