aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--boltzgen/kernel/template/basic.cl.mako (renamed from boltzgen/kernel/template/basic.opencl.mako)31
-rw-r--r--boltzgen/kernel/template/basic.cpp.mako56
2 files changed, 53 insertions, 34 deletions
diff --git a/boltzgen/kernel/template/basic.opencl.mako b/boltzgen/kernel/template/basic.cl.mako
index 3ecb16c..1b02c63 100644
--- a/boltzgen/kernel/template/basic.opencl.mako
+++ b/boltzgen/kernel/template/basic.cl.mako
@@ -1,11 +1,3 @@
-% 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
-
<%
def gid():
return {
@@ -15,8 +7,22 @@ def gid():
def pop_offset(i):
return i * geometry.volume
+
+def neighbor_offset(c_i):
+ return {
+ 2: lambda: c_i[1]*geometry.size_x + c_i[0],
+ 3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
+ }.get(descriptor.d)()
%>
+% 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)
{
@@ -31,15 +37,6 @@ __kernel void equilibrilize(__global ${float_type}* f_next,
% endfor
}
-<%
-def neighbor_offset(c_i):
- return {
- 2: lambda: c_i[1]*geometry.size_x + c_i[0],
- 3: lambda: c_i[2]*geometry.size_x*geometry.size_y + c_i[1]*geometry.size_x + c_i[0]
- }.get(descriptor.d)()
-
-%>
-
__kernel void collide_and_stream(__global ${float_type}* f_next,
__global ${float_type}* f_prev,
__global int* material,
diff --git a/boltzgen/kernel/template/basic.cpp.mako b/boltzgen/kernel/template/basic.cpp.mako
index ef3486f..d284a1c 100644
--- a/boltzgen/kernel/template/basic.cpp.mako
+++ b/boltzgen/kernel/template/basic.cpp.mako
@@ -1,22 +1,7 @@
<%
def pop_offset(i):
return i * geometry.volume
-%>
-
-void equilibrilize(${float_type}* f_next,
- ${float_type}* f_prev,
- const std::size_t gid)
-{
- ${float_type}* preshifted_f_next = f_next + gid;
- ${float_type}* preshifted_f_prev = f_prev + gid;
-% for i, w_i in enumerate(descriptor.w):
- preshifted_f_next[${pop_offset(i)}] = ${w_i.evalf()};
- preshifted_f_prev[${pop_offset(i)}] = ${w_i.evalf()};
-% endfor
-}
-
-<%
def neighbor_offset(c_i):
return {
2: lambda: c_i[1]*geometry.size_x + c_i[0],
@@ -30,10 +15,23 @@ def padding():
}.get(descriptor.d)()
%>
+void equilibrilize(${float_type}* f_next,
+ ${float_type}* f_prev,
+ std::size_t gid)
+{
+ ${float_type}* preshifted_f_next = f_next + gid;
+ ${float_type}* preshifted_f_prev = f_prev + gid;
+
+% for i, w_i in enumerate(descriptor.w):
+ preshifted_f_next[${pop_offset(i)}] = ${w_i.evalf()};
+ preshifted_f_prev[${pop_offset(i)}] = ${w_i.evalf()};
+% endfor
+}
+
void collide_and_stream( ${float_type}* f_next,
const ${float_type}* f_prev,
- const int* material,
- const std::size_t gid)
+ const int* material,
+ std::size_t gid)
{
const int m = material[gid];
@@ -64,3 +62,27 @@ void collide_and_stream( ${float_type}* f_next,
preshifted_f_next[${pop_offset(i)}] = m*f_next_${i} + (1.0-m)*${descriptor.w[i].evalf()};
% 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;
+
+% for i in range(0,descriptor.q):
+ const ${float_type} f_curr_${i} = preshifted_f[${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
+}