diff options
Diffstat (limited to 'template')
| -rw-r--r-- | template/opengl.mako | 44 |
1 files changed, 40 insertions, 4 deletions
diff --git a/template/opengl.mako b/template/opengl.mako index 2029623..6958db8 100644 --- a/template/opengl.mako +++ b/template/opengl.mako @@ -64,12 +64,12 @@ def neighbor_offset(c_i): %> -__kernel void collect_gl_moments_to_texture(__global ${float_type}* f, - __global int* material, +__kernel void collect_gl_moments_and_materials_to_texture(__global ${float_type}* f, + __global int* material, % if descriptor.d == 2: - __write_only image2d_t moments) + __write_only image2d_t moments) % elif descriptor.d == 3: - __write_only image3d_t moments) + __write_only image3d_t moments) % endif { const unsigned int gid = ${gid()}; @@ -107,3 +107,39 @@ __kernel void collect_gl_moments_to_texture(__global ${float_type}* f, write_imagef(moments, ${moments_cell()}, data); } + +__kernel void collect_gl_moments_to_texture(__global ${float_type}* f, +% 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 ${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 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 + + write_imagef(moments, ${moments_cell()}, data); +} |