aboutsummaryrefslogtreecommitdiff
path: root/boltzgen/kernel/template/basic.opencl.mako
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-23 22:03:46 +0200
committerAdrian Kummerlaender2019-10-23 22:03:46 +0200
commit2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983 (patch)
tree33dc0d73468a173378f1f6d8d7cffde2c8734c8d /boltzgen/kernel/template/basic.opencl.mako
parent54efae7153dd9a6ef1768ff3a054c082e090b4ca (diff)
downloadboltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar.gz
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar.bz2
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar.lz
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar.xz
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.tar.zst
boltzgen-2e07d5fac4d3fedcd2cc9ccbab27f50c69f65983.zip
Some cleanup, add `collect_moments` to C++ template
Diffstat (limited to 'boltzgen/kernel/template/basic.opencl.mako')
-rw-r--r--boltzgen/kernel/template/basic.opencl.mako104
1 files changed, 0 insertions, 104 deletions
diff --git a/boltzgen/kernel/template/basic.opencl.mako b/boltzgen/kernel/template/basic.opencl.mako
deleted file mode 100644
index 3ecb16c..0000000
--- a/boltzgen/kernel/template/basic.opencl.mako
+++ /dev/null
@@ -1,104 +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
-
-<%
-def gid():
- return {
- 2: 'get_global_id(1)*%d + get_global_id(0)' % geometry.size_x,
- 3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (geometry.size_x*geometry.size_y, geometry.size_x)
- }.get(descriptor.d)
-
-def pop_offset(i):
- return i * geometry.volume
-%>
-
-__kernel void equilibrilize(__global ${float_type}* f_next,
- __global ${float_type}* f_prev)
-{
- const unsigned int gid = ${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[${pop_offset(i)}] = ${w_i}.f;
- preshifted_f_prev[${pop_offset(i)}] = ${w_i}.f;
-% 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,
- unsigned int time)
-{
- const unsigned int gid = ${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[${pop_offset(i) + 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[${pop_offset(i)}] = f_next_${i};
-% endfor
-}
-
-__kernel void collect_moments(__global ${float_type}* f,
- __global ${float_type}* moments)
-{
- const unsigned int gid = ${gid()};
-
- __global ${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):
- moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
-% endfor
-}