3 files changed, 74 insertions, 0 deletions
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)()