diff options
| -rw-r--r-- | README.md | 1 | ||||
| -rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cl.mako | 15 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collect_moments.cl.mako | 15 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collide_and_stream.cl.mako | 15 | ||||
| -rw-r--r-- | boltzgen/kernel/template/equilibrilize.cl.mako | 15 | ||||
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cl.mako | 20 |
6 files changed, 68 insertions, 13 deletions
@@ -12,6 +12,7 @@ At the moment this is a more structured and cleaned up version of the OpenCL ker * equilibrilization and moment collection utility functions * optimization via common subexpression elimination * array-of-structures and structure-of-arrays memory layouts +* static resolution of memory offsets * AB streaming pattern * C++ and OpenCL targets * simple CLI frontend 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 diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako index ef520da..0ab42d1 100644 --- a/boltzgen/kernel/template/collect_moments.cl.mako +++ b/boltzgen/kernel/template/collect_moments.cl.mako @@ -1,6 +1,6 @@ -__kernel void collect_moments(__global ${float_type}* f, - __global ${float_type}* moments, - unsigned int gid) +__kernel void collect_moments_gid(__global ${float_type}* f, + __global ${float_type}* moments, + unsigned int gid) { __global ${float_type}* preshifted_f = f + gid; @@ -16,3 +16,12 @@ __kernel void collect_moments(__global ${float_type}* f, moments[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)}; % endfor } + +% if 'cell_list_dispatch' in extras: +__kernel void collect_moments_cells(__global ${float_type}* f, + __global ${float_type}* moments, + __global unsigned int* cells) +{ + collect_moments_gid(f, moments, cells[get_global_id(0)]); +} +% endif diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako index 28cfa57..a8fe532 100644 --- a/boltzgen/kernel/template/collide_and_stream.cl.mako +++ b/boltzgen/kernel/template/collide_and_stream.cl.mako @@ -1,6 +1,6 @@ -__kernel void collide_and_stream(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - unsigned int gid) +__kernel void collide_and_stream_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 collide_and_stream(__global ${float_type}* f_next, preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i}; % endfor } + +% if 'cell_list_dispatch' in extras: +__kernel void collide_and_stream_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells) +{ + collide_and_stream_gid(f_next, f_prev, cells[get_global_id(0)]); +} +% endif diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako index 0759dd5..4b9b984 100644 --- a/boltzgen/kernel/template/equilibrilize.cl.mako +++ b/boltzgen/kernel/template/equilibrilize.cl.mako @@ -1,6 +1,6 @@ -__kernel void equilibrilize(__global ${float_type}* f_next, - __global ${float_type}* f_prev, - unsigned int gid) +__kernel void equilibrilize_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; @@ -10,3 +10,12 @@ __kernel void equilibrilize(__global ${float_type}* f_next, preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f; % endfor } + +% if 'cell_list_dispatch' in extras: +__kernel void equilibrilize_cells(__global ${float_type}* f_next, + __global ${float_type}* f_prev, + __global unsigned int* cells) +{ + equilibrilize_gid(f_next, f_prev, cells[get_global_id(0)]); +} +% endif 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 |