diff options
Diffstat (limited to 'boltzgen/kernel')
| -rw-r--r-- | boltzgen/kernel/generator.py | 14 | ||||
| -rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cl.mako | 8 | ||||
| -rw-r--r-- | boltzgen/kernel/template/bounce_back_boundary.cpp.mako | 8 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collect_moments.cl.mako | 8 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collect_moments.cpp.mako | 10 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collide_and_stream.cl.mako | 8 | ||||
| -rw-r--r-- | boltzgen/kernel/template/collide_and_stream.cpp.mako | 9 | ||||
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cl.mako | 12 | ||||
| -rw-r--r-- | boltzgen/kernel/template/momenta_boundary.cpp.mako | 12 |
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> |