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.

6 files changed, 68 insertions, 13 deletions

diff --git a/README.md b/README.md
index 679cf90..9a7052b 100644
--- a/README.md
+++ b/README.md
@@ -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