From ab616f6f6762a895eb5d49b0048d68ff3ca01650 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sat, 20 Jun 2020 15:06:03 +0200
Subject: Fix usage of MomentsTexture in volumetric examples

---
 template/opengl.mako | 44 ++++++++++++++++++++++++++++++++++++++++----
 1 file changed, 40 insertions(+), 4 deletions(-)

(limited to 'template/opengl.mako')

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);
+}
-- 
cgit v1.2.3