diff options
Diffstat (limited to 'template/kernel.mako')
| -rw-r--r-- | template/kernel.mako | 44 |
1 files changed, 22 insertions, 22 deletions
diff --git a/template/kernel.mako b/template/kernel.mako index 417851a..41edcbf 100644 --- a/template/kernel.mako +++ b/template/kernel.mako @@ -9,13 +9,13 @@ def pop_offset(i): return i * geometry.volume %> -__kernel void equilibrilize(__global __write_only float* f_next, - __global __write_only float* f_prev) +__kernel void equilibrilize(__global __write_only ${float_type}* f_next, + __global __write_only ${float_type}* f_prev) { const unsigned int gid = ${gid()}; - __global __write_only float* preshifted_f_next = f_next + gid; - __global __write_only float* preshifted_f_prev = f_prev + gid; + __global __write_only ${float_type}* preshifted_f_next = f_next + gid; + __global __write_only ${float_type}* preshifted_f_prev = f_prev + gid; % if pop_eq_src == '': % for i, w_i in enumerate(descriptor.w): @@ -36,8 +36,8 @@ def neighbor_offset(c_i): %> -__kernel void collide_and_stream(__global __write_only float* f_next, - __global __read_only float* f_prev, +__kernel void collide_and_stream(__global __write_only ${float_type}* f_next, + __global __read_only ${float_type}* f_prev, __global __read_only int* material) { const unsigned int gid = ${gid()}; @@ -48,29 +48,29 @@ __kernel void collide_and_stream(__global __write_only float* f_next, return; } - __global __write_only float* preshifted_f_next = f_next + gid; - __global __read_only float* preshifted_f_prev = f_prev + gid; + __global __write_only ${float_type}* preshifted_f_next = f_next + gid; + __global __read_only ${float_type}* preshifted_f_prev = f_prev + gid; % for i, c_i in enumerate(descriptor.c): - const float f_curr_${i} = preshifted_f_prev[${pop_offset(i) + neighbor_offset(-c_i)}]; + 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 ${expr[0]} = ${ccode(expr[1])}; + const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor % for i, expr in enumerate(moments_assignment): - float ${ccode(expr)} + ${float_type} ${ccode(expr)} % endfor ${boundary_src} % for i, expr in enumerate(collide_subexpr): - const float ${expr[0]} = ${ccode(expr[1])}; + const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor % for i, expr in enumerate(collide_assignment): - const float ${ccode(expr)} + const ${float_type} ${ccode(expr)} % endfor % for i in range(0,descriptor.q): @@ -78,19 +78,19 @@ __kernel void collide_and_stream(__global __write_only float* f_next, % endfor } -__kernel void collect_moments(__global __read_only float* f, - __global __write_only float* moments) +__kernel void collect_moments(__global __read_only ${float_type}* f, + __global __write_only ${float_type}* moments) { const unsigned int gid = ${gid()}; - __global __read_only float* preshifted_f = f + gid; + __global __read_only ${float_type}* preshifted_f = f + gid; % for i in range(0,descriptor.q): - const float f_curr_${i} = preshifted_f[${pop_offset(i)}]; + const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}]; % endfor % for i, expr in enumerate(moments_subexpr): - const float ${expr[0]} = ${ccode(expr[1])}; + const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor % for i, expr in enumerate(moments_assignment): @@ -98,19 +98,19 @@ __kernel void collect_moments(__global __read_only float* f, % endfor } -__kernel void collect_gl_moments(__global __read_only float* f, +__kernel void collect_gl_moments(__global __read_only ${float_type}* f, __global __write_only float4* moments) { const unsigned int gid = ${gid()}; - __global __read_only float* preshifted_f = f + gid; + __global __read_only ${float_type}* preshifted_f = f + gid; % for i in range(0,descriptor.q): - const float f_curr_${i} = preshifted_f[${pop_offset(i)}]; + const ${float_type} f_curr_${i} = preshifted_f[${pop_offset(i)}]; % endfor % for i, expr in enumerate(moments_subexpr): - const float ${expr[0]} = ${ccode(expr[1])}; + const ${float_type} ${expr[0]} = ${ccode(expr[1])}; % endfor float4 data; |