aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/pattern/AB.cl.mako
blob: 33836ea37d12dbd1a1cc82bd8d8a5b67a1501db5 (plain)
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
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
<%def name="operator(name, params = None)">
__kernel void ${name}(
      __global ${float_type}* f_next
    , __global ${float_type}* f_prev
    , unsigned int gid
% if params is not None:
% for param_type, param_name in params:
    , ${param_type} ${param_name}
% endfor
% endif
) {
    __global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
    __global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};

% for i, c_i in enumerate(descriptor.c):
    const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
% endfor

    ${caller.body()}

% for i, _ in enumerate(descriptor.c):
    preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
% endfor
}

% if 'cell_list_dispatch' in extras:
__kernel void ${name}_cells(
      __global ${float_type}* f_next
    , __global ${float_type}* f_prev
    , __global unsigned int* cells
% if params is not None:
% for param_type, param_name in params:
    , ${param_type} ${param_name}
% endfor
% endif
) {
    ${name}(
          f_next
        , f_prev
        , cells[get_global_id(0)]
% if params is not None:
% for param_type, param_name in params:
        , ${param_name}
% endfor
% endif
    );
}
% endif
</%def>

<%def name="functor(name, params = None)">
__kernel void ${name}(
      __global ${float_type}* f
    , unsigned int gid
% if params is not None:
% for param_type, param_name in params:
    , ${param_type} ${param_name}
% endfor
% endif
) {
    __global ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')};

% for i in range(0,descriptor.q):
    const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
% endfor

    ${caller.body()}
}

% if 'cell_list_dispatch' in extras:
__kernel void ${name}_cells(
      __global ${float_type}* f
    , __global unsigned int* cells
% if params is not None:
% for param_type, param_name in params:
    , ${param_type} ${param_name}
% endfor
% endif
) {
    ${name}(
          f
        , cells[get_global_id(0)]
% if params is not None:
% for param_type, param_name in params:
        , ${param_name}
% endfor
% endif
    );
}
% endif
</%def>