aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--boltzgen/kernel/generator.py14
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cl.mako8
-rw-r--r--boltzgen/kernel/template/bounce_back_boundary.cpp.mako8
-rw-r--r--boltzgen/kernel/template/collect_moments.cl.mako8
-rw-r--r--boltzgen/kernel/template/collect_moments.cpp.mako10
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cl.mako8
-rw-r--r--boltzgen/kernel/template/collide_and_stream.cpp.mako9
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cl.mako12
-rw-r--r--boltzgen/kernel/template/momenta_boundary.cpp.mako12
9 files changed, 54 insertions, 35 deletions
diff --git a/boltzgen/kernel/generator.py b/boltzgen/kernel/generator.py
index 4fd3c25..8da91ba 100644
--- a/boltzgen/kernel/generator.py
+++ b/boltzgen/kernel/generator.py
@@ -1,5 +1,3 @@
-import sympy
-
from mako.template import Template
from pathlib import Path
@@ -33,12 +31,8 @@ class Generator:
geometry = geometry,
index = self.index_impl(geometry),
layout = self.layout_impl(self.descriptor, self.index_impl, geometry),
-
- ccode = sympy.ccode,
-
float_type = self.float_type,
-
- extras = extras
+ extras = extras
)
def kernel(self, geometry, functions, extras = []):
@@ -47,15 +41,13 @@ class Generator:
return "\n".join(map(lambda f: self.instantiate(f, geometry, extras), functions))
- def custom(self, geometry, source):
+ def custom(self, geometry, source, extras = []):
return Template(text = source).render(
descriptor = self.descriptor,
model = self.model,
geometry = geometry,
index = self.index_impl(geometry),
layout = self.layout_impl(self.descriptor, self.index_impl, geometry),
-
- ccode = sympy.ccode,
-
float_type = self.float_type,
+ extras = extras
)
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
index 4766bba..e26cbe1 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
__kernel void bounce_back_boundary_gid(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
unsigned int gid)
@@ -14,11 +18,11 @@ __kernel void bounce_back_boundary_gid(__global ${float_type}* f_next,
%>
% for i, expr in enumerate(subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i, c_i in enumerate(descriptor.c):
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
index c37233e..1dcafe8 100644
--- a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
+++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
void bounce_back_boundary( ${float_type}* f_next,
const ${float_type}* f_prev,
std::size_t gid)
@@ -14,11 +18,11 @@ void bounce_back_boundary( ${float_type}* f_next,
%>
% for i, expr in enumerate(subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i, c_i in enumerate(descriptor.c):
diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako
index 587f58a..39317e3 100644
--- a/boltzgen/kernel/template/collect_moments.cl.mako
+++ b/boltzgen/kernel/template/collect_moments.cl.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
__kernel void collect_moments_gid(__global ${float_type}* f,
__global ${float_type}* m,
unsigned int gid)
@@ -14,11 +18,11 @@ __kernel void collect_moments_gid(__global ${float_type}* f,
%>
% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(moments_assignment):
- preshifted_m[${i}] = ${ccode(expr.rhs)};
+ preshifted_m[${i}] = ${sympy.ccode(expr.rhs)};
% endfor
}
diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako
index 1605084..493c53d 100644
--- a/boltzgen/kernel/template/collect_moments.cpp.mako
+++ b/boltzgen/kernel/template/collect_moments.cpp.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
void collect_moments(const ${float_type}* f,
std::size_t gid,
${float_type}& rho,
@@ -14,14 +18,14 @@ void collect_moments(const ${float_type}* f,
%>
% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(moments_assignment):
% if i == 0:
- rho = ${ccode(expr.rhs)};
+ rho = ${sympy.ccode(expr.rhs)};
% else:
- u[${i-1}] = ${ccode(expr.rhs)};
+ u[${i-1}] = ${sympy.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
index 6b55cd2..bfc9435 100644
--- a/boltzgen/kernel/template/collide_and_stream.cl.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cl.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
__kernel void collide_and_stream_gid(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
unsigned int gid)
@@ -14,11 +18,11 @@ __kernel void collide_and_stream_gid(__global ${float_type}* f_next,
%>
% for i, expr in enumerate(subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i in range(0,descriptor.q):
diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako
index 2c52b54..ee42eb7 100644
--- a/boltzgen/kernel/template/collide_and_stream.cpp.mako
+++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako
@@ -1,3 +1,7 @@
+<%
+import sympy
+%>
+
void collide_and_stream( ${float_type}* f_next,
const ${float_type}* f_prev,
std::size_t gid)
@@ -14,15 +18,14 @@ void collide_and_stream( ${float_type}* f_next,
%>
% for i, expr in enumerate(subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i, expr in enumerate(assignment):
preshifted_f_next[${layout.pop_offset(i)}] = f_next_${i};
% endfor
}
-
diff --git a/boltzgen/kernel/template/momenta_boundary.cl.mako b/boltzgen/kernel/template/momenta_boundary.cl.mako
index 7c1e3df..b0b4c9e 100644
--- a/boltzgen/kernel/template/momenta_boundary.cl.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cl.mako
@@ -1,4 +1,6 @@
<%
+import sympy
+
moments_subexpr, moments_assignment = model.moments()
collision_subexpr, collision_assignment = model.collision(f_eq = model.equilibrium(resolve_moments = False))
%>
@@ -17,17 +19,17 @@ __kernel void ${name}_momenta_boundary_gid(
% endfor
% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
${caller.body()}
% for i, expr in enumerate(collision_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(collision_assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i, expr in enumerate(collision_assignment):
@@ -37,7 +39,7 @@ __kernel void ${name}_momenta_boundary_gid(
</%def>
<%call expr="momenta_boundary('velocity', '%s%d velocity' % (float_type, descriptor.d))">
- ${float_type} ${ccode(moments_assignment[0])}
+ ${float_type} ${sympy.ccode(moments_assignment[0])}
% for i, expr in enumerate(moments_assignment[1:]):
${float_type} ${expr.lhs} = velocity.${['x', 'y', 'z'][i]};
% endfor
@@ -46,7 +48,7 @@ __kernel void ${name}_momenta_boundary_gid(
<%call expr="momenta_boundary('density', '%s density' % float_type)">
${float_type} ${moments_assignment[0].lhs} = density;
% for i, expr in enumerate(moments_assignment[1:]):
- ${float_type} ${ccode(expr)}
+ ${float_type} ${sympy.ccode(expr)}
% endfor
</%call>
diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako
index c41d07e..ae78e9f 100644
--- a/boltzgen/kernel/template/momenta_boundary.cpp.mako
+++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako
@@ -1,4 +1,6 @@
<%
+import sympy
+
moments_subexpr, moments_assignment = model.moments()
collision_subexpr, collision_assignment = model.collision(f_eq = model.equilibrium(resolve_moments = False))
%>
@@ -17,17 +19,17 @@ void ${name}_momenta_boundary(
% endfor
% for i, expr in enumerate(moments_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
${caller.body()}
% for i, expr in enumerate(collision_subexpr):
- const ${float_type} ${expr[0]} = ${ccode(expr[1])};
+ const ${float_type} ${expr[0]} = ${sympy.ccode(expr[1])};
% endfor
% for i, expr in enumerate(collision_assignment):
- const ${float_type} ${ccode(expr)}
+ const ${float_type} ${sympy.ccode(expr)}
% endfor
% for i, expr in enumerate(collision_assignment):
@@ -37,7 +39,7 @@ void ${name}_momenta_boundary(
</%def>
<%call expr="momenta_boundary('velocity', '%s velocity[%d]' % (float_type, descriptor.d))">
- ${float_type} ${ccode(moments_assignment[0])}
+ ${float_type} ${sympy.ccode(moments_assignment[0])}
% for i, expr in enumerate(moments_assignment[1:]):
${float_type} ${expr.lhs} = velocity[${i}];
% endfor
@@ -46,6 +48,6 @@ void ${name}_momenta_boundary(
<%call expr="momenta_boundary('density', '%s density' % float_type)">
${float_type} ${moments_assignment[0].lhs} = density;
% for i, expr in enumerate(moments_assignment[1:]):
- ${float_type} ${ccode(expr)}
+ ${float_type} ${sympy.ccode(expr)}
% endfor
</%call>