diff options
Diffstat (limited to 'template/kernel.mako')
| -rw-r--r-- | template/kernel.mako | 109 |
1 files changed, 21 insertions, 88 deletions
diff --git a/template/kernel.mako b/template/kernel.mako index 529eb30..2e26124 100644 --- a/template/kernel.mako +++ b/template/kernel.mako @@ -1,3 +1,11 @@ +% 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 { @@ -9,13 +17,13 @@ def pop_offset(i): return i * memory.volume %> -__kernel void equilibrilize(__global __write_only ${float_type}* f_next, - __global __write_only ${float_type}* f_prev) +__kernel void equilibrilize(__global ${float_type}* f_next, + __global ${float_type}* f_prev) { const unsigned int gid = ${gid()}; - __global __write_only ${float_type}* preshifted_f_next = f_next + gid; - __global __write_only ${float_type}* preshifted_f_prev = f_prev + gid; + __global ${float_type}* preshifted_f_next = f_next + gid; + __global ${float_type}* preshifted_f_prev = f_prev + gid; % if pop_eq_src == '': % for i, w_i in enumerate(descriptor.w): @@ -36,9 +44,9 @@ def neighbor_offset(c_i): %> -__kernel void collide_and_stream(__global __write_only ${float_type}* f_next, - __global __read_only ${float_type}* f_prev, - __global __read_only int* material, +__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()}; @@ -49,8 +57,8 @@ __kernel void collide_and_stream(__global __write_only ${float_type}* f_next, return; } - __global __write_only ${float_type}* preshifted_f_next = f_next + gid; - __global __read_only ${float_type}* preshifted_f_prev = f_prev + 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[${pop_offset(i) + neighbor_offset(-c_i)}]; @@ -79,12 +87,12 @@ __kernel void collide_and_stream(__global __write_only ${float_type}* f_next, % endfor } -__kernel void collect_moments(__global __read_only ${float_type}* f, - __global __write_only ${float_type}* moments) +__kernel void collect_moments(__global ${float_type}* f, + __global ${float_type}* moments) { const unsigned int gid = ${gid()}; - __global __read_only ${float_type}* preshifted_f = f + 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)}]; @@ -97,79 +105,4 @@ __kernel void collect_moments(__global __read_only ${float_type}* f, % for i, expr in enumerate(moments_assignment): moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)}; % endfor -} - -__kernel void collect_gl_moments(__global __read_only ${float_type}* f, - __global __read_only int* material, - __global __write_only float4* moments) -{ - const unsigned int gid = ${gid()}; - - __global __read_only ${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 - - float4 data; - - if (material[gid] == 1) { -% if descriptor.d == 2: - data.x = ${ccode(moments_assignment[0].rhs)}; - data.y = ${ccode(moments_assignment[1].rhs)}; - data.z = ${ccode(moments_assignment[2].rhs)}; - data.w = sqrt(data.y*data.y + data.z*data.z); -% elif descriptor.d == 3: - data.x = ${ccode(moments_assignment[0].rhs)}; - data.y = ${ccode(moments_assignment[1].rhs)}; - data.z = ${ccode(moments_assignment[2].rhs)}; - data.w = ${ccode(moments_assignment[3].rhs)}; -% endif - } else { - data.x = 0.0; - data.y = 0.0; - data.z = 0.0; - data.w = -material[gid]; - } - - moments[gid] = data; -} - -__kernel void update_particles(__global __read_only float4* moments, - __global __read_only int* material, - __global __write_only float4* particles, - __global __read_only float4* init_particles, - float aging) -{ - const unsigned int pid = get_global_id(0); - - float4 particle = particles[pid]; - -% if descriptor.d == 2: - const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x); -% elif descriptor.d == 3: - const unsigned int gid = floor(particle.z)*${memory.size_x*memory.size_y} + floor(particle.y)*${memory.size_x} + floor(particle.x); -% endif - - const float4 moment = moments[gid]; - - if (material[gid] == 1 && particle.w < 1.0) { - particle.x += moment.y; - particle.y += moment.z; -% if descriptor.d == 2: - particle.w += min(particle.x, particle.y) * aging; -% elif descriptor.d == 3: - particle.z += moment.w; - particle.w += min(min(particle.x, particle.y), particle.z) * aging; -% endif - } else { - particle.xyz = init_particles[pid].xyz; - particle.w = particle.w-1.0; - } - - particles[pid] = particle; -} +}
\ No newline at end of file |