diff options
Diffstat (limited to 'template/opengl.mako')
-rw-r--r-- | template/opengl.mako | 100 |
1 files changed, 100 insertions, 0 deletions
diff --git a/template/opengl.mako b/template/opengl.mako new file mode 100644 index 0000000..6b09f69 --- /dev/null +++ b/template/opengl.mako @@ -0,0 +1,100 @@ +<% +def gid(): + return { + 2: 'get_global_id(1)*%d + get_global_id(0)' % memory.size_x, + 3: 'get_global_id(2)*%d + get_global_id(1)*%d + get_global_id(0)' % (memory.size_x*memory.size_y, memory.size_x) + }.get(descriptor.d) + +def pop_offset(i): + return i * memory.volume + +def moments_cell(): + return { + 2: '(int2)(get_global_id(0), get_global_id(1))', + 3: '(int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0)' + }.get(descriptor.d) +%> + +__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 collect_gl_moments_to_texture(__global __read_only ${float_type}* f, + __global __read_only int* material, +% if descriptor.d == 2: + __write_only image2d_t moments) +% elif descriptor.d == 3: + __write_only image3d_t moments) +% endif +{ + 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] != 5) { +% 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 = 1.0; + } + + write_imagef(moments, ${moments_cell()}, data); +} |