diff options
author | Adrian Kummerlaender | 2019-10-27 16:03:07 +0100 |
---|---|---|
committer | Adrian Kummerlaender | 2019-10-27 16:10:33 +0100 |
commit | 6938758f6c1754f0ee49d0709dd0ca376a146010 (patch) | |
tree | 53e196ae894bba8610f26fd9b418b802f9797a03 /boltzgen/kernel/target/layout | |
parent | 74c60dcbe56862d73b000d569423cb298fb06686 (diff) | |
download | boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar.gz boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar.bz2 boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar.lz boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar.xz boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.tar.zst boltzgen-6938758f6c1754f0ee49d0709dd0ca376a146010.zip |
Move layout implementations into separate folder
Diffstat (limited to 'boltzgen/kernel/target/layout')
-rw-r--r-- | boltzgen/kernel/target/layout/__init__.py | 2 | ||||
-rw-r--r-- | boltzgen/kernel/target/layout/cl.py | 39 | ||||
-rw-r--r-- | boltzgen/kernel/target/layout/cpp.py | 33 |
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)() |