aboutsummaryrefslogtreecommitdiff
path: root/template/opengl.mako
diff options
context:
space:
mode:
Diffstat (limited to 'template/opengl.mako')
-rw-r--r--template/opengl.mako44
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);
+}