aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-06-16 13:51:35 +0200
committerAdrian Kummerlaender2019-06-16 13:51:35 +0200
commitb64054acaa9c9710c5b4140cfdb8e32e392a5df9 (patch)
treed02abc3884292ca3c2ecafb95906fb46a68457aa
parent54a45dbec3e046b3e3c333aa1fde68244d0d2a51 (diff)
downloadsymlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar.gz
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar.bz2
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar.lz
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar.xz
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.tar.zst
symlbm_playground-b64054acaa9c9710c5b4140cfdb8e32e392a5df9.zip
Declutter gid and offset calculation
-rw-r--r--symbolic/generator.py2
-rw-r--r--template/kernel.mako28
2 files changed, 16 insertions, 14 deletions
diff --git a/symbolic/generator.py b/symbolic/generator.py
index e57154c..161e2f4 100644
--- a/symbolic/generator.py
+++ b/symbolic/generator.py
@@ -1,7 +1,7 @@
from sympy import *
from sympy.codegen.ast import Assignment
-class LBM():
+class LBM:
def __init__(self, descriptor):
self.descriptor = descriptor
self.f_next = symarray('f_next', descriptor.q)
diff --git a/template/kernel.mako b/template/kernel.mako
index 79cae66..01c5b33 100644
--- a/template/kernel.mako
+++ b/template/kernel.mako
@@ -1,7 +1,15 @@
+<%
+def gid():
+ return {
+ 2: 'get_global_id(1)*%d + get_global_id(0)' % geometry.size_x,
+ 3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (geometry.size_x*geometry.size_y, geometry.size_x)
+ }.get(descriptor.d)
+%>
+
__kernel void equilibrilize(__global __write_only float* f_next,
__global __write_only float* f_prev)
{
- const unsigned int gid = get_global_id(2)*(${geometry.size_x*geometry.size_y}) + get_global_id(1)*${geometry.size_x} + get_global_id(0);
+ const unsigned int gid = ${gid()};
__global __write_only float* preshifted_f_next = f_next + gid;
__global __write_only float* preshifted_f_prev = f_prev + gid;
@@ -18,16 +26,10 @@ __kernel void equilibrilize(__global __write_only float* f_next,
<%
def neighbor_offset(c_i):
- if descriptor.d == 2:
- if c_i[1] == 0:
- return c_i[0]
- else:
- return c_i[1]*geometry.size_x + c_i[0]
- elif descriptor.d == 3:
- if c_i[1] == 0:
- return c_i[2]*geometry.size_x*geometry.size_y + c_i[0]
- else:
- return c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
+ return {
+ 2: lambda: c_i[1]*geometry.size_x + c_i[0],
+ 3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
+ }.get(descriptor.d)()
%>
@@ -35,7 +37,7 @@ __kernel void collide_and_stream(__global __write_only float* f_next,
__global __read_only float* f_prev,
__global __read_only int* material)
{
- const unsigned int gid = get_global_id(2)*(${geometry.size_x*geometry.size_y}) + get_global_id(1)*${geometry.size_x} + get_global_id(0);
+ const unsigned int gid = ${gid()};
const int m = material[gid];
@@ -76,7 +78,7 @@ __kernel void collide_and_stream(__global __write_only float* f_next,
__kernel void collect_moments(__global __read_only float* f,
__global __write_only float* moments)
{
- const unsigned int gid = get_global_id(2)*(${geometry.size_x*geometry.size_y}) + get_global_id(1)*${geometry.size_x} + get_global_id(0);
+ const unsigned int gid = ${gid()};
__global __read_only float* preshifted_f = f + gid;