aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/update_sss_control_structure.cl.mako
blob: ec6a5d54c89733e5dbbe38a2bd5c33a371b476d7 (plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
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
}