aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-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