blob: 2a98b946fd11af55ca6a3d323a5f069bf3535eef (
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)
%>
__global__ void init_sss_control_structure(${float_type}* f, ${float_type}** control) {
% for i, c_i in enumerate(descriptor.c):
control[${i}] = f + ${padding + layout.pop_offset(i, 2*padding)};
% endfor
}
__global__ void update_sss_control_structure(${float_type}** f) {
% for i, c_i in enumerate(descriptor.c):
${float_type}* f_old_${i} = f[${i}];
% endfor
% for i, c_i in enumerate(descriptor.c):
f[${i}] = f_old_${descriptor.c.index(-c_i)} + ${layout.neighbor_offset(-c_i)};
% endfor
}
|