aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-26 14:21:17 +0200
committerAdrian Kummerlaender2019-10-26 14:21:17 +0200
commit71e0d07ec0cdf528eb5d962ce97bd6c55da52eae (patch)
treed8c7a290de7c7d7b092ac18d63b4b321a0606161
parent769bc4c97d5686ab614bb134abddef87f3f41519 (diff)
downloadboltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar.gz
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar.bz2
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar.lz
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar.xz
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.tar.zst
boltzgen-71e0d07ec0cdf528eb5d962ce97bd6c55da52eae.zip
Implement AOS layout for the OpenCL target
-rw-r--r--boltzgen/kernel/target/cl.py20
1 files changed, 20 insertions, 0 deletions
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)()