aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--ldc_2d_gl_interop.py98
-rw-r--r--simulation.py40
-rw-r--r--template/kernel.mako25
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;
+}