From 73de5f16efc696cf0c88beec086eb9a4df9098dd Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Thu, 24 Oct 2019 21:52:45 +0200
Subject: Extract offset helper into target and layout specific classes

---
 boltzgen/kernel/template/basic.cl.mako | 35 +++++++++-------------------------
 1 file changed, 9 insertions(+), 26 deletions(-)

(limited to 'boltzgen/kernel/template/basic.cl.mako')

diff --git a/boltzgen/kernel/template/basic.cl.mako b/boltzgen/kernel/template/basic.cl.mako
index 1b02c63..b64a480 100644
--- a/boltzgen/kernel/template/basic.cl.mako
+++ b/boltzgen/kernel/template/basic.cl.mako
@@ -1,20 +1,3 @@
-<%
-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)
-
-def pop_offset(i):
-    return i * geometry.volume
-
-def neighbor_offset(c_i):
-    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)()
-%>
-
 % if float_type == 'double':
 #if defined(cl_khr_fp64)
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
@@ -26,14 +9,14 @@ def neighbor_offset(c_i):
 __kernel void equilibrilize(__global ${float_type}* f_next,
                             __global ${float_type}* f_prev)
 {
-    const unsigned int gid = ${gid()};
+    const unsigned int gid = ${layout.gid()};
 
     __global ${float_type}* preshifted_f_next = f_next + gid;
     __global ${float_type}* preshifted_f_prev = f_prev + gid;
 
 % for i, w_i in enumerate(descriptor.w):
-    preshifted_f_next[${pop_offset(i)}] = ${w_i}.f;
-    preshifted_f_prev[${pop_offset(i)}] = ${w_i}.f;
+    preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f;
+    preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f;
 % endfor
 }
 
@@ -42,7 +25,7 @@ __kernel void collide_and_stream(__global ${float_type}* f_next,
                                  __global int* material,
                                  unsigned int time)
 {
-    const unsigned int gid = ${gid()};
+    const unsigned int gid = ${layout.gid()};
 
     const int m = material[gid];
 
@@ -54,7 +37,7 @@ __kernel void collide_and_stream(__global ${float_type}* f_next,
     __global ${float_type}* preshifted_f_prev = f_prev + gid;
 
 % for i, c_i in enumerate(descriptor.c):
-    const ${float_type} f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}];
+    const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
 % endfor
 
 % for i, expr in enumerate(moments_subexpr):
@@ -76,19 +59,19 @@ __kernel void collide_and_stream(__global ${float_type}* f_next,
 % endfor
 
 % for i in range(0,descriptor.q):
-    preshifted_f_next[${pop_offset(i)}] = f_next_${i};
+    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
 % endfor
 }
 
 __kernel void collect_moments(__global ${float_type}* f,
                               __global ${float_type}* moments)
 {
-    const unsigned int gid = ${gid()};
+    const unsigned int gid = ${layout.gid()};
 
     __global ${float_type}* preshifted_f = f + gid;
 
 % for i in range(0,descriptor.q):
-    const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}];
+    const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
 % endfor
 
 % for i, expr in enumerate(moments_subexpr):
@@ -96,6 +79,6 @@ __kernel void collect_moments(__global ${float_type}* f,
 % endfor
 
 % for i, expr in enumerate(moments_assignment):
-    moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
+    moments[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
 % endfor
 }
-- 
cgit v1.2.3