diff options
Optionally generate cell-list-based OpenCL dispatch functions
Requires different function naming as OpenCL 1.2 doesn't support overloads.
The OpenCL kernel code generated using this commit was successfully tested
on an actual GPU. Time to set up some automatic validation.
Diffstat (limited to 'boltzgen/kernel/template/bounce_back_boundary.cl.mako')
-rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cl.mako | 15 |
1 files changed, 12 insertions, 3 deletions
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako index 7a9bc2f..0762a09 100644 --- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako +++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako @@ -1,6 +1,6 @@ -__kernel void bounce_back_boundary(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - unsigned int gid) +__kernel void bounce_back_boundary_gid(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + unsigned int gid) { __global ${float_type}* preshifted_f_next = f_next + gid; __global ${float_type}* preshifted_f_prev = f_prev + gid; @@ -29,3 +29,12 @@ __kernel void bounce_back_boundary(__global ${float_type}* f_next, preshifted_f_next[${layout.pop_offset(i)}] = f_next_${descriptor.c.index(-c_i)}; % endfor } + +% if 'cell_list_dispatch' in extras: +__kernel void bounce_back_boundary_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells) +{ + bounce_back_boundary_gid(f_next, f_prev, cells[get_global_id(0)]); +} +% endif |