aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/target
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-24 21:52:45 +0200
committerAdrian Kummerlaender2019-10-24 21:52:45 +0200
commit73de5f16efc696cf0c88beec086eb9a4df9098dd (patch)
tree543f1eee6f9ae8391838635268a83177dc5c1c2e /boltzgen/kernel/target
parentb3d131b94f9417c0c4cd6733433c86ca780dde5e (diff)
downloadboltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar.gz
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar.bz2
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar.lz
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar.xz
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.tar.zst
boltzgen-73de5f16efc696cf0c88beec086eb9a4df9098dd.zip
Extract offset helper into target and layout specific classes
Diffstat (limited to 'boltzgen/kernel/target')
-rw-r--r--boltzgen/kernel/target/cl.py19
-rw-r--r--boltzgen/kernel/target/cpp.py45
2 files changed, 64 insertions, 0 deletions
diff --git a/boltzgen/kernel/target/cl.py b/boltzgen/kernel/target/cl.py
new file mode 100644
index 0000000..a2fc819
--- /dev/null
+++ b/boltzgen/kernel/target/cl.py
@@ -0,0 +1,19 @@
+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)()
diff --git a/boltzgen/kernel/target/cpp.py b/boltzgen/kernel/target/cpp.py
new file mode 100644
index 0000000..de79716
--- /dev/null
+++ b/boltzgen/kernel/target/cpp.py
@@ -0,0 +1,45 @@
+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)()
+
+ def padding(self):
+ return self.descriptor.q * {
+ 2: lambda: 1*self.geometry.size_y + 1,
+ 3: lambda: 1*self.geometry.size_y*self.geometry.size_z + 1*self.geometry.size_z + 1
+ }.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)()
+
+ def padding(self):
+ return {
+ 2: lambda: 1*self.geometry.size_y + 1,
+ 3: lambda: 1*self.geometry.size_y*self.geometry.size_z + 1*self.geometry.size_z + 1
+ }.get(self.descriptor.d)()