aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/pattern/AB.cl.mako
diff options
context:
space:
mode:
Diffstat (limited to 'boltzgen/kernel/template/pattern/AB.cl.mako')
-rw-r--r--boltzgen/kernel/template/pattern/AB.cl.mako64
1 files changed, 29 insertions, 35 deletions
diff --git a/boltzgen/kernel/template/pattern/AB.cl.mako b/boltzgen/kernel/template/pattern/AB.cl.mako
index 8e6421e..8fa941b 100644
--- a/boltzgen/kernel/template/pattern/AB.cl.mako
+++ b/boltzgen/kernel/template/pattern/AB.cl.mako
@@ -2,13 +2,20 @@
__kernel void ${name}(
__global ${float_type}* f_next
, __global ${float_type}* f_prev
+% if 'cell_list_dispatch' in extras:
+ , __global unsigned int* cells
+% else:
, unsigned int 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
__global ${float_type}* preshifted_f_next = f_next + ${layout.cell_preshift('gid')};
__global ${float_type}* preshifted_f_prev = f_prev + ${layout.cell_preshift('gid')};
@@ -22,42 +29,51 @@ __kernel void ${name}(
preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
% endfor
}
+</%def>
-% if 'cell_list_dispatch' in extras:
-__kernel void ${name}_cells(
+<%def name="operator_with_domain_dispatch(name, params = None)">
+__kernel void ${name}(
__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}
+ const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)', 'get_global_id(2)')};
+ __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
-% endif
- );
}
-% endif
</%def>
<%def name="functor(name, params = None)">
__kernel void ${name}(
__global ${float_type}* f
+% if 'cell_list_dispatch' in extras:
+ , __global unsigned int* cells
+% else:
, unsigned int 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
__global ${float_type}* preshifted_f = f + ${layout.cell_preshift('gid')};
% for i in range(0,descriptor.q):
@@ -66,28 +82,6 @@ __kernel void ${name}(
${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>
<%def name="functor_with_domain_dispatch(name, params = None)">