diff options
author | Adrian Kummerlaender | 2019-10-31 13:13:00 +0100 |
---|---|---|
committer | Adrian Kummerlaender | 2019-10-31 13:13:00 +0100 |
commit | e2b00f4ec963060be98939c7b0d12d6c00e50a02 (patch) | |
tree | 9b8729d491a2d9cb2b5fafe8284bc807ad21665a /boltzgen/kernel/template | |
parent | c82b38122cc3ab7717cb0ba9ec530b4658bd03e4 (diff) | |
download | boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar.gz boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar.bz2 boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar.lz boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar.xz boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.tar.zst boltzgen-e2b00f4ec963060be98939c7b0d12d6c00e50a02.zip |
Call symbolic generator inside code templates
This paves the way for dropping in other LBM collision models.
As a side benefit the default momenta calulcation is now fully inlined where possible.
Diffstat (limited to 'boltzgen/kernel/template')
8 files changed, 39 insertions, 37 deletions
diff --git a/boltzgen/kernel/template/bounce_back_boundary.cl.mako b/boltzgen/kernel/template/bounce_back_boundary.cl.mako index cb59cbd..8bb4f6c 100644 --- a/boltzgen/kernel/template/bounce_back_boundary.cl.mako +++ b/boltzgen/kernel/template/bounce_back_boundary.cl.mako @@ -9,19 +9,15 @@ __kernel void bounce_back_boundary_gid(__global ${float_type}* f_next, 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 +<% + subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True)) +%> -% for i, expr in enumerate(collision_subexpr): +% for i, expr in enumerate(subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor -% for i, expr in enumerate(collision_assignment): +% for i, expr in enumerate(assignment): const ${float_type} ${ccode(expr)} % endfor diff --git a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako index c7abd2a..91dcfa0 100644 --- a/boltzgen/kernel/template/bounce_back_boundary.cpp.mako +++ b/boltzgen/kernel/template/bounce_back_boundary.cpp.mako @@ -9,19 +9,15 @@ void bounce_back_boundary( ${float_type}* f_next, 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 +<% + subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True)) +%> -% for i, expr in enumerate(collision_subexpr): +% for i, expr in enumerate(subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor -% for i, expr in enumerate(collision_assignment): +% for i, expr in enumerate(assignment): const ${float_type} ${ccode(expr)} % endfor diff --git a/boltzgen/kernel/template/collect_moments.cl.mako b/boltzgen/kernel/template/collect_moments.cl.mako index ddf5102..587f58a 100644 --- a/boltzgen/kernel/template/collect_moments.cl.mako +++ b/boltzgen/kernel/template/collect_moments.cl.mako @@ -9,6 +9,10 @@ __kernel void collect_moments_gid(__global ${float_type}* f, const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; % endfor +<% + moments_subexpr, moments_assignment = model.moments() +%> + % for i, expr in enumerate(moments_subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor diff --git a/boltzgen/kernel/template/collect_moments.cpp.mako b/boltzgen/kernel/template/collect_moments.cpp.mako index d72ff5d..1605084 100644 --- a/boltzgen/kernel/template/collect_moments.cpp.mako +++ b/boltzgen/kernel/template/collect_moments.cpp.mako @@ -9,6 +9,10 @@ void collect_moments(const ${float_type}* f, const ${float_type} f_curr_${i} = preshifted_f[${layout.pop_offset(i)}]; % endfor +<% + moments_subexpr, moments_assignment = model.moments() +%> + % for i, expr in enumerate(moments_subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor diff --git a/boltzgen/kernel/template/collide_and_stream.cl.mako b/boltzgen/kernel/template/collide_and_stream.cl.mako index c586884..303b21b 100644 --- a/boltzgen/kernel/template/collide_and_stream.cl.mako +++ b/boltzgen/kernel/template/collide_and_stream.cl.mako @@ -9,19 +9,15 @@ __kernel void collide_and_stream_gid(__global ${float_type}* f_next, 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 +<% + subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True)) +%> -% for i, expr in enumerate(collision_subexpr): +% for i, expr in enumerate(subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor -% for i, expr in enumerate(collision_assignment): +% for i, expr in enumerate(assignment): const ${float_type} ${ccode(expr)} % endfor diff --git a/boltzgen/kernel/template/collide_and_stream.cpp.mako b/boltzgen/kernel/template/collide_and_stream.cpp.mako index 20fc529..47dd6dd 100644 --- a/boltzgen/kernel/template/collide_and_stream.cpp.mako +++ b/boltzgen/kernel/template/collide_and_stream.cpp.mako @@ -9,23 +9,19 @@ void collide_and_stream( ${float_type}* f_next, 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 +<% + subexpr, assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = True)) +%> -% for i, expr in enumerate(collision_subexpr): +% for i, expr in enumerate(subexpr): const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor -% for i, expr in enumerate(collision_assignment): +% for i, expr in enumerate(assignment): const ${float_type} ${ccode(expr)} % endfor -% for i, expr in enumerate(collision_assignment): +% 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 f5edb61..4e29601 100644 --- a/boltzgen/kernel/template/momenta_boundary.cl.mako +++ b/boltzgen/kernel/template/momenta_boundary.cl.mako @@ -1,3 +1,8 @@ +<% +moments_subexpr, moments_assignment = model.moments() +collision_subexpr, collision_assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = False)) +%> + <%def name="momenta_boundary(name, param)"> __kernel void ${name}_momenta_boundary_gid( __global ${float_type}* f_next, diff --git a/boltzgen/kernel/template/momenta_boundary.cpp.mako b/boltzgen/kernel/template/momenta_boundary.cpp.mako index 3f51087..6d05288 100644 --- a/boltzgen/kernel/template/momenta_boundary.cpp.mako +++ b/boltzgen/kernel/template/momenta_boundary.cpp.mako @@ -1,3 +1,8 @@ +<% +moments_subexpr, moments_assignment = model.moments() +collision_subexpr, collision_assignment = model.bgk(f_eq = model.equilibrium(resolve_moments = False)) +%> + <%def name="momenta_boundary(name, param)"> void ${name}_momenta_boundary( ${float_type}* f_next, |