aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 22:22:24 +0100
committerAdrian Kummerlaender2019-10-27 22:22:24 +0100
commit536b25e2c2b742c17d023d2b3386fed4dc60a339 (patch)
tree7424a6431d8076257125a70adcd45bf417aa38c1
parent1b9ac6e7aee3cf63495a65c2d7dbf79a0be23d7d (diff)
downloadboltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.gz
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.bz2
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.tar.xz
boltzgen-536b25e2c2b742c17d023d2b3386fed4dc60a339.zip
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.
-rw-r--r--README.md1
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako15
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako15
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako15
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako15
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cl.mako20
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