diff options
| author | Adrian Kummerlaender | 2019-11-04 23:38:36 +0100 |
|---|---|---|
| committer | Adrian Kummerlaender | 2019-11-04 23:38:36 +0100 |
| commit | 5828235f806c3e87a5b1eed34ef69ef317a110bd (patch) | |
| tree | f4fd988f2af0e0ca2cd0e1b277f26fcf8b0f3b3a /boltzgen/kernel/template/equilibrilize.cl.mako | |
| parent | 05e74fb112f5b5f645b649c587d18052c7b7f9df (diff) | |
| download | boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.gz boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.bz2 boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.lz boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.xz boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.tar.zst boltzgen-5828235f806c3e87a5b1eed34ef69ef317a110bd.zip | |
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.
Diffstat (limited to 'boltzgen/kernel/template/equilibrilize.cl.mako')
| -rw-r--r-- | boltzgen/kernel/template/equilibrilize.cl.mako | 22 |
1 files changed, 4 insertions, 18 deletions
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 +</%call> |