From dfcb729ff162aba45e3d83acf864fe1ff92e5a06 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Thu, 20 Jun 2019 15:43:58 +0200
Subject: Prototype OpenGL interoperation

---
 template/kernel.mako | 25 +++++++++++++++++++++++++
 1 file changed, 25 insertions(+)

(limited to 'template/kernel.mako')

diff --git a/template/kernel.mako b/template/kernel.mako
index 1790f88..d9b1c80 100644
--- a/template/kernel.mako
+++ b/template/kernel.mako
@@ -97,3 +97,28 @@ __kernel void collect_moments(__global __read_only  float* f,
     moments[${pop_offset(i)} + gid] = ${ccode(expr.rhs)};
 % endfor
 }
+
+__kernel void collect_gl_moments(__global __read_only  float* f,
+                                 __global __write_only float4* moments)
+{
+    const unsigned int gid = ${gid()};
+
+    __global __read_only float* preshifted_f = f + gid;
+
+% for i in range(0,descriptor.q):
+    const float f_curr_${i} = preshifted_f[${pop_offset(i)}];
+% endfor
+
+% for i, expr in enumerate(moments_subexpr):
+    const float ${expr[0]} = ${ccode(expr[1])};
+% endfor
+
+    float4 data;
+
+    data.x = 4.0*((float)(get_global_id(0))) + ${ccode(2000*moments_assignment[1].rhs)};
+    data.y = 4.0*((float)(get_global_id(1))) + ${ccode(2000*moments_assignment[2].rhs)};
+    data.z = 0.0;
+    data.w = 1.0;
+
+    moments[gid] = data;
+}
-- 
cgit v1.2.3