From 5828235f806c3e87a5b1eed34ef69ef317a110bd Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Mon, 4 Nov 2019 23:38:36 +0100 Subject: Extract streaming pattern into Mako definitions This should allow for plugging in e.g. a AA pattern implementation without without touching any file but `AA.$target.mako`. OpenCL and C++ target templates now look basically the same and could potentially be merged. However this would decrease flexibility should more differences appear in the future. Maintaining separate template files is an acceptable overhead to preserve flexibility. --- boltzgen/kernel/template/equilibrilize.cl.mako | 22 ++++------------------ 1 file changed, 4 insertions(+), 18 deletions(-) (limited to 'boltzgen/kernel/template/equilibrilize.cl.mako') diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako index 4ee8d41..ba01712 100644 --- a/boltzgen/kernel/template/equilibrilize.cl.mako +++ b/boltzgen/kernel/template/equilibrilize.cl.mako @@ -1,21 +1,7 @@ -__kernel void equilibrilize_gid(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - unsigned int gid) -{ - __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')}; - __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')}; +<%namespace name="pattern" file="${'/pattern/%s.cl.mako' % context['streaming']}"/> +<%call expr="pattern.operator_ab('equilibrilize')"> % for i, w_i in enumerate(descriptor.w): - preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f; - preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f; + ${float_type} f_next_${i} = ${w_i.evalf()}; % endfor -} - -% if 'cell_list_dispatch' in extras: -__kernel void equilibrilize_cells(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - __global unsigned int* cells) -{ - equilibrilize_gid(f_next, f_prev, cells[get_global_id(0)]); -} -% endif + -- cgit v1.2.3