aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-27 14:05:21 +0100
committerAdrian Kummerlaender2019-10-27 14:05:21 +0100
commit18c54d79699db7554faa851c87d7113db67a8a08 (patch)
treec89243b5796bb8f1e6a8cea241f3a5e1bf934ba1
parent7fa72d8718d96727bcfd60cc3bcb1609526d3c9b (diff)
downloadboltzgen-18c54d79699db7554faa851c87d7113db67a8a08.tar
boltzgen-18c54d79699db7554faa851c87d7113db67a8a08.tar.gz
boltzgen-18c54d79699db7554faa851c87d7113db67a8a08.tar.bz2
boltzgen-18c54d79699db7554faa851c87d7113db67a8a08.tar.xz
boltzgen-18c54d79699db7554faa851c87d7113db67a8a08.zip
Separate functions into separate template files
Selection of the desired templates is possible via a new `functions` parameter.
-rwxr-xr-xboltzgen.py15
-rw-r--r--boltzgen/kernel/generator.py46
-rw-r--r--boltzgen/kernel/template/basic.cl.mako84
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako19
-rw-r--r--boltzgen/kernel/template/collect_moments.cpp.mako24
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako31
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cpp.mako32
-rw-r--r--boltzgen/kernel/template/equilibrilize.cl.mako13
-rw-r--r--boltzgen/kernel/template/equilibrilize.cpp.mako13
-rw-r--r--boltzgen/kernel/template/example.cpp.mako (renamed from boltzgen/kernel/template/basic.cpp.mako)108
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cpp.mako35
-rw-r--r--boltzgen/kernel/template/preamble.cl.mako7
12 files changed, 209 insertions, 218 deletions
diff --git a/boltzgen.py b/boltzgen.py
index 2b32661..5224942 100755
--- a/boltzgen.py
+++ b/boltzgen.py
@@ -15,7 +15,8 @@ argparser.add_argument('--geometry', required = True, help = 'Size of the block
argparser.add_argument('--tau', required = True, help = 'BGK relaxation time')
argparser.add_argument('--disable-cse', action = 'store_const', const = True, help = 'Disable common subexpression elimination')
-argparser.add_argument('--extra', action = 'append', nargs = '+', default = [], help = 'Additional generator parameters')
+argparser.add_argument('--functions', action = 'append', nargs = '+', default = [], help = 'Function templates to be generated')
+argparser.add_argument('--extras', action = 'append', nargs = '+', default = [], help = 'Additional generator parameters')
args = argparser.parse_args()
@@ -29,5 +30,15 @@ generator = Generator(
geometry = Geometry.parse(args.geometry)
-src = generator.kernel(args.language, args.precision, args.layout, geometry, sum(args.extra, []))
+functions = sum(args.functions, [])
+if len(functions) == 0:
+ functions += ['default']
+if 'default' in functions:
+ for f in ['collide_and_stream', 'equilibrilize', 'collect_moments', 'momenta_boundary']:
+ functions.insert(functions.index('default'), f)
+ functions.remove('default')
+
+extras = sum(args.extras, [])
+
+src = generator.kernel(args.language, args.precision, args.layout, geometry, functions, extras)
print(src)
diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py
index 8d0de37..04b0a64 100644
--- a/boltzgen/kernel/generator.py
+++ b/boltzgen/kernel/generator.py
@@ -7,25 +7,15 @@ import kernel.target.cl
import kernel.target.cpp
class Generator:
- def __init__(self, descriptor, moments, collision, boundary = ''):
+ def __init__(self, descriptor, moments, collision):
self.descriptor = descriptor
self.moments = moments
self.collision = collision
- self.boundary = boundary
- def kernel(self, target, precision, layout, geometry, extras = []):
- template_path = Path(__file__).parent/('template/basic.' + target + '.mako')
+ def instantiate(self, target, template, float_type, layout_impl, geometry, extras = []):
+ template_path = Path(__file__).parent/("template/%s.%s.mako" % (template, target))
if not template_path.exists():
- raise Exception("Target '%s' not supported" % target)
-
- layout_impl = eval("kernel.target.%s.%s" % (target, layout))
- if layout_impl is None:
- raise Exception("Target '%s' doesn't support layout '%s'" % (target, layout))
- else:
- layout_impl = layout_impl(self.descriptor, geometry)
-
- if geometry.dimension() != self.descriptor.d:
- raise Exception('Geometry dimension must match descriptor dimension')
+ raise Exception("Target '%s' doesn't provide '%s'" % (target, template))
return Template(filename = str(template_path)).render(
descriptor = self.descriptor,
@@ -38,16 +28,24 @@ class Generator:
collision_assignment = self.collision[1],
ccode = sympy.ccode,
- float_type = {
- 'single': 'float',
- 'double': 'double'
- }.get(precision),
-
- boundary_src = Template(self.boundary).render(
- descriptor = self.descriptor,
- geometry = geometry,
- float_type = precision
- ),
+ float_type = float_type,
extras = extras
)
+
+ def kernel(self, target, precision, layout, geometry, functions = ['collide_and_stream'], extras = []):
+ layout_impl = eval("kernel.target.%s.%s" % (target, layout))
+ if layout_impl is None:
+ raise Exception("Target '%s' doesn't support layout '%s'" % (target, layout))
+ else:
+ layout_impl = layout_impl(self.descriptor, geometry)
+
+ if geometry.dimension() != self.descriptor.d:
+ raise Exception('Geometry dimension must match descriptor dimension')
+
+ float_type = {
+ 'single': 'float',
+ 'double': 'double'
+ }.get(precision)
+
+ return "\n".join(map(lambda f: self.instantiate(target, f, float_type, layout_impl, geometry, extras), functions))
diff --git a/boltzgen/kernel/template/basic.cl.mako b/boltzgen/kernel/template/basic.cl.mako
deleted file mode 100644
index b64a480..0000000
--- a/boltzgen/kernel/template/basic.cl.mako
+++ /dev/null
@@ -1,84 +0,0 @@
-% if float_type == 'double':
-#if defined(cl_khr_fp64)
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-#elif defined(cl_amd_fp64)
-#pragma OPENCL EXTENSION cl_amd_fp64 : enable
-#endif
-% endif
-
-__kernel void equilibrilize(__global ${float_type}* f_next,
- __global ${float_type}* f_prev)
-{
- const unsigned int gid = ${layout.gid()};
-
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
-
-% for i, w_i in enumerate(descriptor.w):
- preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f;
- preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f;
-% endfor
-}
-
-__kernel void collide_and_stream(__global ${float_type}* f_next,
- __global ${float_type}* f_prev,
- __global int* material,
- unsigned int time)
-{
- const unsigned int gid = ${layout.gid()};
-
- const int m = material[gid];
-
- if ( m == 0 ) {
- return;
- }
-
- __global ${float_type}* preshifted_f_next = f_next + gid;
- __global ${float_type}* preshifted_f_prev = f_prev + gid;
-
-% for i, c_i in enumerate(descriptor.c):
- const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
- ${float_type} ${ccode(expr)}
-% endfor
-
- ${boundary_src}
-
-% for i, expr in enumerate(collision_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
- const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i in range(0,descriptor.q):
- preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-__kernel void collect_moments(__global ${float_type}* f,
- __global ${float_type}* moments)
-{
- const unsigned int gid = ${layout.gid()};
-
- __global ${float_type}* preshifted_f = f + gid;
-
-% for i in range(0,descriptor.q):
- const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
- moments[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
-% endfor
-}
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
new file mode 100644
index 0000000..b07b759
--- /dev/null
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -0,0 +1,19 @@
+__kernel void collect_moments(__global ${float_type}* f,
+ __global ${float_type}* moments)
+{
+ const unsigned int gid = ${layout.gid()};
+
+ __global ${float_type}* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+ const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+ moments[${layout.pop_offset(i)} + gid] = ${ccode(expr.rhs)};
+% endfor
+}
diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako
new file mode 100644
index 0000000..8c37db2
--- /dev/null
+++ b/boltzgen/kernel/template/collect_moments.cpp.mako
@@ -0,0 +1,24 @@
+void collect_moments(const ${float_type}* f,
+ std::size_t gid,
+ ${float_type}& rho,
+ ${float_type} u[${descriptor.d}])
+{
+ const ${float_type}* preshifted_f = f + gid*${layout.gid_offset()};
+
+% for i in range(0,descriptor.q):
+ const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+% if i == 0:
+ rho = ${ccode(expr.rhs)};
+% else:
+ u[${i-1}] = ${ccode(expr.rhs)};
+% endif
+% endfor
+}
+
diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako
new file mode 100644
index 0000000..28cfa57
--- /dev/null
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -0,0 +1,31 @@
+__kernel void collide_and_stream(__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;
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+ ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+ const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i in range(0,descriptor.q):
+ preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako
new file mode 100644
index 0000000..71c62b7
--- /dev/null
+++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako
@@ -0,0 +1,32 @@
+void collide_and_stream( ${float_type}* f_next,
+ const ${float_type}* f_prev,
+ std::size_t gid)
+{
+ ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+ const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+ ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+ const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+ preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
+
diff --git a/boltzgen/kernel/template/equilibrilize.cl.mako b/boltzgen/kernel/template/equilibrilize.cl.mako
new file mode 100644
index 0000000..aa2246c
--- /dev/null
+++ b/boltzgen/kernel/template/equilibrilize.cl.mako
@@ -0,0 +1,13 @@
+__kernel void equilibrilize(__global ${float_type}* f_next,
+ __global ${float_type}* f_prev)
+{
+ const unsigned int gid = ${layout.gid()};
+
+ __global ${float_type}* preshifted_f_next = f_next + gid;
+ __global ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, w_i in enumerate(descriptor.w):
+ preshifted_f_next[${layout.pop_offset(i)}] = ${w_i}.f;
+ preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i}.f;
+% endfor
+}
diff --git a/boltzgen/kernel/template/equilibrilize.cpp.mako b/boltzgen/kernel/template/equilibrilize.cpp.mako
new file mode 100644
index 0000000..9f082a1
--- /dev/null
+++ b/boltzgen/kernel/template/equilibrilize.cpp.mako
@@ -0,0 +1,13 @@
+void equilibrilize(${float_type}* f_next,
+ ${float_type}* f_prev,
+ std::size_t gid)
+{
+ ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+ ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% for i, w_i in enumerate(descriptor.w):
+ preshifted_f_next[${layout.pop_offset(i)}] = ${w_i.evalf()};
+ preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i.evalf()};
+% endfor
+}
+
diff --git a/boltzgen/kernel/template/basic.cpp.mako b/boltzgen/kernel/template/example.cpp.mako
index 118ef8c..d979bfd 100644
--- a/boltzgen/kernel/template/basic.cpp.mako
+++ b/boltzgen/kernel/template/example.cpp.mako
@@ -1,108 +1,3 @@
-void equilibrilize(${float_type}* f_next,
- ${float_type}* f_prev,
- std::size_t gid)
-{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% for i, w_i in enumerate(descriptor.w):
- preshifted_f_next[${layout.pop_offset(i)}] = ${w_i.evalf()};
- preshifted_f_prev[${layout.pop_offset(i)}] = ${w_i.evalf()};
-% endfor
-}
-
-void collide_and_stream( ${float_type}* f_next,
- const ${float_type}* f_prev,
- std::size_t gid)
-{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% for i, c_i in enumerate(descriptor.c):
- const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
- ${float_type} ${ccode(expr)}
-% endfor
-
-% for i, expr in enumerate(collision_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
- const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
- preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-void velocity_momenta_boundary( ${float_type}* f_next,
- const ${float_type}* f_prev,
- std::size_t gid,
- ${float_type} velocity[${descriptor.d}])
-{
- ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
- const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
-
-% for i, c_i in enumerate(descriptor.c):
- const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
- ${float_type} ${ccode(expr)}
-% endfor
-% for i, expr in enumerate(moments_assignment[1:]):
- ${expr.lhs} = velocity[${i}];
-% endfor
-
-% for i, expr in enumerate(collision_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
- const ${float_type} ${ccode(expr)}
-% endfor
-
-% for i, expr in enumerate(collision_assignment):
- preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-void collect_moments(const ${float_type}* f,
- std::size_t gid,
- ${float_type}& rho,
- ${float_type} u[${descriptor.d}])
-{
- const ${float_type}* preshifted_f = f + gid*${layout.gid_offset()};
-
-% for i in range(0,descriptor.q):
- const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}];
-% endfor
-
-% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
-% endfor
-
-% for i, expr in enumerate(moments_assignment):
-% if i == 0:
- rho = ${ccode(expr.rhs)};
-% else:
- u[${i-1}] = ${ccode(expr.rhs)};
-% endif
-% endfor
-}
-
% if 'moments_vtk' in extras:
void collect_moments_to_vtk(const std::string& path, ${float_type}* f) {
std::ofstream fout;
@@ -186,7 +81,6 @@ void collect_moments_to_vtk(const std::string& path, ${float_type}* f) {
}
% endif
-% if 'test_ldc' in extras:
void test_ldc(std::size_t nStep)
{
auto f_a = std::make_unique<${float_type}[]>(${geometry.volume*descriptor.q});
@@ -276,5 +170,3 @@ void test_ldc(std::size_t nStep)
collect_moments_to_vtk("test.vtk", f_next);
% endif
}
-% endif
-
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
new file mode 100644
index 0000000..d7e0dd2
--- /dev/null
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -0,0 +1,35 @@
+void velocity_momenta_boundary( ${float_type}* f_next,
+ const ${float_type}* f_prev,
+ std::size_t gid,
+ ${float_type} velocity[${descriptor.d}])
+{
+ ${float_type}* preshifted_f_next = f_next + gid*${layout.gid_offset()};
+ const ${float_type}* preshifted_f_prev = f_prev + gid*${layout.gid_offset()};
+
+% for i, c_i in enumerate(descriptor.c):
+ const ${float_type} f_curr_${i} = preshifted_f_prev[${layout.pop_offset(i) + layout.neighbor_offset(-c_i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(moments_assignment):
+ ${float_type} ${ccode(expr)}
+% endfor
+% for i, expr in enumerate(moments_assignment[1:]):
+ ${expr.lhs} = velocity[${i}];
+% endfor
+
+% for i, expr in enumerate(collision_subexpr):
+ const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+ const ${float_type} ${ccode(expr)}
+% endfor
+
+% for i, expr in enumerate(collision_assignment):
+ preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
+% endfor
+}
diff --git a/boltzgen/kernel/template/preamble.cl.mako b/boltzgen/kernel/template/preamble.cl.mako
new file mode 100644
index 0000000..c47254d
--- /dev/null
+++ b/boltzgen/kernel/template/preamble.cl.mako
@@ -0,0 +1,7 @@
+% if float_type == 'double':
+#if defined(cl_khr_fp64)
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+#elif defined(cl_amd_fp64)
+#pragma OPENCL EXTENSION cl_amd_fp64 : enable
+#endif
+% endif