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 --- ldc_2d_gl_interop.py | 98 ++++++++++++++++++++++++++++++++++++++++++++++++++++ simulation.py | 40 ++++++++++++++++++--- template/kernel.mako | 25 ++++++++++++++ 3 files changed, 158 insertions(+), 5 deletions(-) create mode 100644 ldc_2d_gl_interop.py 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; +} -- cgit v1.2.3