Prototype OpenGL interoperation

3 files changed, 158 insertions, 5 deletions

diff --git a/ldc_2d_gl_interop.py b/ldc_2d_gl_interop.py
new file mode 100644
index 0000000..6ff10c7
--- /dev/null
+++ b/ldc_2d_gl_interop.py
@@ -0,0 +1,98 @@
+import numpy
+import time
+
+from simulation         import Lattice, Geometry
+from symbolic.generator import LBM
+
+import symbolic.D2Q9 as D2Q9
+
+from OpenGL.GL   import *
+from OpenGL.GLUT import *
+
+screen_x = 1920
+screen_y = 1200
+
+def cavity(geometry, x, y):
+    if x == 1 or y == 1 or x == geometry.size_x-2:
+        return 2
+    elif y == geometry.size_y-2:
+        return 3
+    else:
+        return 1
+
+boundary = """
+    if ( m == 2 ) {
+        u_0 = 0.0;
+        u_1 = 0.0;
+    }
+    if ( m == 3 ) {
+        u_0 = 0.1;
+        u_1 = 0.0;
+    }
+"""
+
+def glut_window(fullscreen = False):
+    glutInit(sys.argv)
+    glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_DEPTH)
+
+    if fullscreen:
+        window = glutEnterGameMode()
+    else:
+        glutInitWindowSize(screen_x, screen_y)
+        glutInitWindowPosition(0, 0)
+        window = glutCreateWindow("LBM")
+
+    glMatrixMode(GL_PROJECTION)
+    glLoadIdentity()
+
+    glOrtho(
+        0, screen_x,
+        0, screen_y,
+        0.1, 100.0
+    )
+
+    return window
+
+lbm = LBM(D2Q9)
+
+window = glut_window(fullscreen = False)
+
+lattice = Lattice(
+    descriptor   = D2Q9,
+    geometry     = Geometry(screen_x//4, screen_y//4),
+    moments      = lbm.moments(optimize = False),
+    collide      = lbm.bgk(f_eq = lbm.equilibrium(), tau = 0.52),
+    boundary_src = boundary,
+    opengl       = True
+)
+
+lattice.setup_geometry(cavity)
+
+def on_display():
+    for i in range(0,100):
+        lattice.evolve()
+
+    lattice.sync_gl_moments()
+
+    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT)
+    glMatrixMode(GL_MODELVIEW)
+    glLoadIdentity()
+
+    glTranslatef(0., 0., -1.)
+
+    lattice.gl_moments.bind()
+    glEnableClientState(GL_VERTEX_ARRAY)
+    glVertexPointer(4, GL_FLOAT, 0, lattice.gl_moments)
+    glDrawArrays(GL_POINTS, 0, lattice.geometry.volume)
+    glDisableClientState(GL_VERTEX_ARRAY)
+
+    glutSwapBuffers()
+
+def on_timer(t):
+    glutTimerFunc(t, on_timer, t)
+    glutPostRedisplay()
+
+glutDisplayFunc(on_display)
+glutTimerFunc(10, on_timer, 10)
+
+glutMainLoop()
diff --git a/simulation.py b/simulation.py
index 672003a..cbc9840 100644
--- a/simulation.py
+++ b/simulation.py
@@ -7,6 +7,10 @@ import sympy
 from mako.template import Template
 from pathlib import Path
 
+from pyopencl.tools import get_gl_sharing_context_properties
+import OpenGL.GL as gl
+from OpenGL.arrays import vbo
+
 class Geometry:
     def __init__(self, size_x, size_y, size_z = 1):
         self.size_x = size_x
@@ -30,9 +34,8 @@ class Geometry:
         else:
             return (self.size_x-2, self.size_y-2, self.size_z-2)
 
-
 class Lattice:
-    def __init__(self, descriptor, geometry, moments, collide, pop_eq_src = '', boundary_src = ''):
+    def __init__(self, descriptor, geometry, moments, collide, pop_eq_src = '', boundary_src = '', opengl = False):
         self.descriptor = descriptor
         self.geometry   = geometry
 
@@ -43,7 +46,12 @@ class Lattice:
         self.boundary_src = boundary_src
 
         self.platform = cl.get_platforms()[0]
-        self.context  = cl.Context(properties=[(cl.context_properties.PLATFORM, self.platform)])
+        if opengl:
+            self.context = cl.Context(
+                properties=[(cl.context_properties.PLATFORM, self.platform)] + get_gl_sharing_context_properties())
+        else:
+            self.context = cl.Context(
+                properties=[(cl.context_properties.PLATFORM, self.platform)])
         self.queue = cl.CommandQueue(self.context)
 
         self.np_material = numpy.ndarray(shape=(self.geometry.volume, 1), dtype=numpy.int32)
@@ -56,14 +64,21 @@ class Lattice:
         self.cl_pop_a = cl.Buffer(self.context, mf.READ_WRITE, size=self.pop_size)
         self.cl_pop_b = cl.Buffer(self.context, mf.READ_WRITE, size=self.pop_size)
 
-        self.cl_moments  = cl.Buffer(self.context, mf.WRITE_ONLY, size=self.moments_size)
+        if opengl:
+            self.np_moments = numpy.ndarray(shape=(self.geometry.volume, 4), dtype=numpy.float32)
+            self.gl_moments = vbo.VBO(data=self.np_moments, usage=gl.GL_DYNAMIC_DRAW, target=gl.GL_ARRAY_BUFFER)
+            self.gl_moments.bind()
+            self.cl_gl_moments  = cl.GLBuffer(self.context, mf.READ_WRITE, int(self.gl_moments))
+        else:
+            self.cl_moments  = cl.Buffer(self.context, mf.WRITE_ONLY, size=self.moments_size)
+
         self.cl_material = cl.Buffer(self.context, mf.READ_ONLY | mf.USE_HOST_PTR, hostbuf=self.np_material)
 
         self.build_kernel()
 
         self.layout = {
             (2, 9): (32,1),
-            (3,19): (32,4,4),
+            (3,19): (32,1,1),
             (3,27): (32,1,1)
         }.get((descriptor.d, descriptor.q), None)
 
@@ -117,11 +132,26 @@ class Lattice:
 
     def get_moments(self):
         moments = numpy.ndarray(shape=(self.descriptor.d+1, self.geometry.volume), dtype=numpy.float32)
+
         if self.tick:
             self.program.collect_moments(
                 self.queue, self.geometry.size(), self.layout, self.cl_pop_b, self.cl_moments)
         else:
             self.program.collect_moments(
                 self.queue, self.geometry.size(), self.layout, self.cl_pop_a, self.cl_moments)
+
         cl.enqueue_copy(self.queue, moments, self.cl_moments).wait();
         return moments
+
+    def sync_gl_moments(self):
+        cl.enqueue_acquire_gl_objects(self.queue, [self.cl_gl_moments])
+
+        if self.tick:
+            self.program.collect_gl_moments(
+                self.queue, self.geometry.size(), self.layout, self.cl_pop_b, self.cl_gl_moments)
+        else:
+            self.program.collect_gl_moments(
+                self.queue, self.geometry.size(), self.layout, self.cl_pop_a, self.cl_gl_moments)
+
+        #cl.enqueue_release_gl_objects(self.queue, [self.cl_gl_moments])
+        self.sync()
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;
+}