1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
|
<%
import numpy
directions = list(filter(lambda c: numpy.sum(numpy.abs(c)) == 1, descriptor.c))
def name(direction):
names = [ ('west','','east'), ('south','','north'), ('down','','up') ]
return ''.join([ names[i][v+1] for i, v in enumerate(direction) ])
def neighbors(direction):
j = list(direction).index(next(x for x in direction if x != 0))
return [ (i,c_i) for i, c_i in enumerate(descriptor.c) if c_i[j] == direction[j] ]
%>
% for direction in directions:
__global__ void communicate_${name(direction)}(
${float_type}** f_a
, ${float_type}** f_b
, std::size_t* cells_a
, std::size_t* cells_b
, std::size_t cell_count
) {
const std::size_t index = blockIdx.x * blockDim.x + threadIdx.x;
if (!(index < cell_count)) {
return;
}
const std::size_t gid_a = cells_a[index];
const std::size_t gid_b = cells_b[index];
% for i, c_i in neighbors(direction):
${float_type}* preshifted_f_a_${i} = f_a[${i}] + ${layout.cell_preshift('gid_a')};
${float_type}* preshifted_f_b_${i} = f_b[${i}] + ${layout.cell_preshift('gid_b')};
% endfor
% for i, c_i in neighbors(direction):
*preshifted_f_b_${i} = *preshifted_f_a_${i};
% endfor
% for i, c_i in neighbors(-direction):
${float_type}* preshifted_f_a_${i} = f_a[${i}] + ${layout.cell_preshift('gid_a')};
${float_type}* preshifted_f_b_${i} = f_b[${i}] + ${layout.cell_preshift('gid_b')};
% endfor
% for i, c_i in neighbors(-direction):
*preshifted_f_a_${i} = *preshifted_f_b_${i};
% endfor
}
% endfor
|
