diff options
| -rw-r--r-- | boltzgen/kernel/template/pattern/SSS.cl.mako | 121 | ||||
| -rw-r--r-- | boltzgen/kernel/template/update_sss_control_structure.cl.mako | 21 |
2 files changed, 142 insertions, 0 deletions
diff --git a/boltzgen/kernel/template/pattern/SSS.cl.mako b/boltzgen/kernel/template/pattern/SSS.cl.mako new file mode 100644 index 0000000..ad5e854 --- /dev/null +++ b/boltzgen/kernel/template/pattern/SSS.cl.mako @@ -0,0 +1,121 @@ +<%def name="operator(name, params = None)"> +<% +if layout.__class__.__name__ != 'SOA': + raise Exception('SSS pattern only works for the AOS memory layout') +%> +__kernel void ${name}( + __global uintptr_t* control +% if 'cell_list_dispatch' in extras: + , __global unsigned* cells +% else: + , std::size_t gid +% endif +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { +% if 'cell_list_dispatch' in extras: + const unsigned int gid = cells[get_global_id(0)]; +% endif + +% for i, c_i in enumerate(descriptor.c): + __global ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')}; +% endfor + +% for i, c_i in enumerate(descriptor.c): + const ${float_type} f_curr_${i} = *preshifted_f_${i}; +% endfor + + ${caller.body()} + +% for i, c_i in enumerate(descriptor.c): + *preshifted_f_${i} = f_next_${descriptor.c.index(-c_i)}; +% endfor +} +</%def> + +<%def name="operator_with_domain_dispatch(name, params = None)"> +__kernel void ${name}( + __global uintptr_t* control +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)', 'get_global_id(2)')}; + +% for i, c_i in enumerate(descriptor.c): + __global ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')}; +% endfor + +% for i, c_i in enumerate(descriptor.c): + const ${float_type} f_curr_${i} = *preshifted_f_${i}; +% endfor + + ${caller.body()} + +% for i, c_i in enumerate(descriptor.c): + *preshifted_f_${i} = f_next_${descriptor.c.index(-c_i)}; +% endfor +} +</%def> + +<%def name="functor(name, params = None)"> +<% +if layout.__class__.__name__ != 'SOA': + raise Exception('SSS pattern only works for the AOS memory layout') +%> +__kernel void ${name}( + __global const uintptr_t* control +% if 'cell_list_dispatch' in extras: + , __global unsigned* cells +% else: + , std::size_t gid +% endif +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { +% if 'cell_list_dispatch' in extras: + const unsigned int gid = cells[get_global_id(0)]; +% endif + +% for i, c_i in enumerate(descriptor.c): + __global const ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')}; +% endfor + +% for i, c_i in enumerate(descriptor.c): + const ${float_type} f_curr_${i} = *preshifted_f_${descriptor.c.index(-c_i)}; +% endfor + + ${caller.body()} +} +</%def> + +<%def name="functor_with_domain_dispatch(name, params = None)"> +__kernel void ${name}( + __global const uintptr_t* control +% if params is not None: +% for param_type, param_name in params: + , ${param_type} ${param_name} +% endfor +% endif +) { + const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)', 'get_global_id(2)')}; + +% for i, c_i in enumerate(descriptor.c): + __global const ${float_type}* preshifted_f_${i} = ((__global ${float_type}*)control[${i}]) + ${layout.cell_preshift('gid')}; +% endfor + +% for i, c_i in enumerate(descriptor.c): + const ${float_type} f_curr_${i} = *preshifted_f_${i}; +% endfor + + ${caller.body()} +} +</%def> diff --git a/boltzgen/kernel/template/update_sss_control_structure.cl.mako b/boltzgen/kernel/template/update_sss_control_structure.cl.mako new file mode 100644 index 0000000..ec6a5d5 --- /dev/null +++ b/boltzgen/kernel/template/update_sss_control_structure.cl.mako @@ -0,0 +1,21 @@ +<% +if streaming != 'SSS': + raise Exception('"update_sss_control_structure" function only makes sense for the SSS pattern') + +padding = (max(geometry.size_x,geometry.size_y,geometry.size_z)+1)**(descriptor.d-1) +%> + +__kernel void init_sss_control_structure(__global ${float_type}* f, __global uintptr_t* control) { +% for i, c_i in enumerate(descriptor.c): + control[${i}] = (uintptr_t)(f + ${padding + layout.pop_offset(i, 2*padding)}); +% endfor +} + +__kernel void update_sss_control_structure(__global uintptr_t* control) { +% for i, c_i in enumerate(descriptor.c): + __global ${float_type}* f_old_${i} = (__global ${float_type}*)(control[${i}]); +% endfor +% for i, c_i in enumerate(descriptor.c): + control[${i}] = (uintptr_t)(f_old_${descriptor.c.index(-c_i)} + ${layout.neighbor_offset(-c_i)}); +% endfor +} |