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/momenta_boundary.cl.mako')
-rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cl.mako | 20 |
1 files changed, 19 insertions, 1 deletions
diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako index 50044d2..e4a8ff3 100644 --- a/boltzgen/kernel/template/momenta_boundary.cl.mako +++ b/boltzgen/kernel/template/momenta_boundary.cl.mako @@ -1,5 +1,5 @@ <%def name="momenta_boundary(name, param)"> -__kernel void ${name}_momenta_boundary( +__kernel void ${name}_momenta_boundary_gid( __global ${float_type}* f_next, __global ${float_type}* f_prev, unsigned int gid, ${param}) @@ -44,3 +44,21 @@ __kernel void ${name}_momenta_boundary( ${float_type} ${ccode(expr)} % endfor </%call> + +% if 'cell_list_dispatch' in extras: +__kernel void velocity_momenta_boundary_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells, + ${float_type}${descriptor.d} velocity) +{ + velocity_momenta_boundary_gid(f_next, f_prev, cells[get_global_id(0)], velocity); +} + +__kernel void density_momenta_boundary_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells, + ${float_type} density) +{ + density_momenta_boundary_gid(f_next, f_prev, cells[get_global_id(0)], density); +} +% endif |