aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template
diff options
context:
space:
mode:
Diffstat (limited to 'boltzgen/kernel/template')
-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
10 files changed, 174 insertions, 192 deletions
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