From 71e0d07ec0cdf528eb5d962ce97bd6c55da52eae Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sat, 26 Oct 2019 14:21:17 +0200
Subject: Implement AOS layout for the OpenCL target

---
 boltzgen/kernel/target/cl.py | 20 ++++++++++++++++++++
 1 file changed, 20 insertions(+)

diff --git a/boltzgen/kernel/target/cl.py b/boltzgen/kernel/target/cl.py
index a2fc819..9707cb3 100644
--- a/boltzgen/kernel/target/cl.py
+++ b/boltzgen/kernel/target/cl.py
@@ -17,3 +17,23 @@ class SOA:
             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