From bf728d8a33b57b2b775b9a1c3bc8f2d84388acd6 Mon Sep 17 00:00:00 2001 From: Adrian Kummerlaender Date: Fri, 13 Sep 2019 23:01:06 +0200 Subject: Add 3d lid driven cavity OpenGL visualization --- ldc_3d_gl_interop.py | 185 +++++++++++++++++++++++++++++++++++++++++++++ template/kernel.mako | 23 ++++-- trugfeuer_2d_gl_interop.py | 2 +- utility/particles.py | 22 ++++-- 4 files changed, 215 insertions(+), 17 deletions(-) create mode 100644 ldc_3d_gl_interop.py diff --git a/ldc_3d_gl_interop.py b/ldc_3d_gl_interop.py new file mode 100644 index 0000000..483ea1a --- /dev/null +++ b/ldc_3d_gl_interop.py @@ -0,0 +1,185 @@ +import numpy +from string import Template + +from simulation import Lattice, Geometry +from utility.particles import Particles +from symbolic.generator import LBM + +import symbolic.D3Q19 as D3Q19 + +from OpenGL.GL import * +from OpenGL.GLUT import * + +from OpenGL.GL import shaders + +from pyrr import matrix44 + +lattice_x = 64 +lattice_y = 32 +lattice_z = 32 + +updates_per_frame = 20 +particle_count = 50000 + +lid_speed = 0.001 +relaxation_time = 0.515 + +def get_cavity_material_map(geometry): + return [ + (lambda x, y, z: x > 0 and x < geometry.size_x-1 and + y > 0 and y < geometry.size_y-1 and + z > 0 and z < geometry.size_z-1, 1), # bulk fluid + (lambda x, y, z: x == 1 or y == 1 or z == 1 or x == geometry.size_x-2 or y == geometry.size_y-2, 2), # walls + (lambda x, y, z: z == geometry.size_z-2, 3), # lid + (lambda x, y, z: x == 0 or x == geometry.size_x-1 or + y == 0 or y == geometry.size_y-1 or + z == 0 or z == geometry.size_z-1, 0) # ghost cells + ] + +boundary = """ + if ( m == 2 ) { + u_0 = 0.0; + u_1 = 0.0; + u_2 = 0.0; + } + if ( m == 3 ) { + u_0 = 0.1; + u_1 = 0.0; + u_2 = 0.0; + } +""" + +def get_projection(width, height): + world_width = lattice_x + world_height = world_width / width * height + + projection = matrix44.create_perspective_projection(45.0, width/height, 0.1, 1000.0) + look = matrix44.create_look_at( + eye = [lattice_x/2, -2*lattice_y, lattice_z/2], + target = [lattice_x/2, lattice_y/2, lattice_z/2], + up = [0, 0, 1]) + rotate = matrix44.create_from_axis_rotation(axis=[0,1,0], theta=0.2) + + point_size = 1 + + return numpy.matmul(look, projection), point_size + +def glut_window(fullscreen = False): + glutInit(sys.argv) + glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_DEPTH) + + if fullscreen: + window = glutEnterGameMode() + else: + glutInitWindowSize(800, 500) + glutInitWindowPosition(0, 0) + window = glutCreateWindow("LBM") + + return window + +lbm = LBM(D3Q19) + +window = glut_window(fullscreen = False) + +vertex_shader = shaders.compileShader(Template(""" +#version 430 + +layout (location=0) in vec4 particles; + +out vec3 color; + +uniform mat4 projection; + +vec3 fire(float x) { + return mix( + vec3(1.0, 1.0, 0.0), + vec3(1.0, 0.0, 0.0), + x + ); +} + +void main() { + gl_Position = projection * vec4( + particles[0], + particles[1], + particles[2], + 1. + ); + + color = fire(1.0-particles[3]); +}""").substitute({}), GL_VERTEX_SHADER) + +fragment_shader = shaders.compileShader(""" +#version 430 + +in vec3 color; + +void main(){ + gl_FragColor = vec4(color.xyz, 0.0); +}""", GL_FRAGMENT_SHADER) + +shader_program = shaders.compileProgram(vertex_shader, fragment_shader) +projection_id = shaders.glGetUniformLocation(shader_program, 'projection') + +lattice = Lattice( + descriptor = D3Q19, + geometry = Geometry(lattice_x, lattice_y, lattice_z), + moments = lbm.moments(optimize = True), + collide = lbm.bgk(f_eq = lbm.equilibrium(), tau = relaxation_time), + boundary_src = boundary, + opengl = True +) + +lattice.apply_material_map( + get_cavity_material_map(lattice.geometry)) +lattice.sync_material() + +particles = Particles( + lattice.context, + lattice.queue, + lattice.memory.float_type, + numpy.mgrid[ + 6*lattice.geometry.size_x//8:7*lattice.geometry.size_x//8:10j, + lattice.geometry.size_y//8:7*lattice.geometry.size_y//8:particle_count/100j, + 6*lattice.geometry.size_z//8:7*lattice.geometry.size_z//8:10j, + ].reshape(3,-1).T) + +def on_display(): + for i in range(0,updates_per_frame): + lattice.evolve() + + lattice.collect_gl_moments() + + for i in range(0,updates_per_frame): + lattice.update_gl_particles(particles, aging = True) + + lattice.sync() + + glClear(GL_COLOR_BUFFER_BIT) + glEnableClientState(GL_VERTEX_ARRAY) + + particles.gl_particles.bind() + + shaders.glUseProgram(shader_program) + glUniformMatrix4fv(projection_id, 1, False, numpy.ascontiguousarray(projection)) + glVertexPointer(4, GL_FLOAT, 0, particles.gl_particles) + glPointSize(point_size) + glEnable(GL_POINT_SMOOTH) + glDrawArrays(GL_POINTS, 0, particles.count) + + glutSwapBuffers() + +def on_reshape(width, height): + global projection, point_size + glViewport(0,0,width,height) + projection, point_size = get_projection(width, height) + +def on_timer(t): + glutTimerFunc(t, on_timer, t) + glutPostRedisplay() + +glutDisplayFunc(on_display) +glutReshapeFunc(on_reshape) +glutTimerFunc(10, on_timer, 10) + +glutMainLoop() diff --git a/template/kernel.mako b/template/kernel.mako index 89000a9..529eb30 100644 --- a/template/kernel.mako +++ b/template/kernel.mako @@ -142,26 +142,33 @@ __kernel void collect_gl_moments(__global __read_only ${float_type}* f, __kernel void update_particles(__global __read_only float4* moments, __global __read_only int* material, __global __write_only float4* particles, - __global __read_only float2* init_particles, + __global __read_only float4* init_particles, float aging) { const unsigned int pid = get_global_id(0); float4 particle = particles[pid]; +% if descriptor.d == 2: const unsigned int gid = floor(particle.y)*${memory.size_x} + floor(particle.x); +% elif descriptor.d == 3: + const unsigned int gid = floor(particle.z)*${memory.size_x*memory.size_y} + floor(particle.y)*${memory.size_x} + floor(particle.x); +% endif - float4 moment = moments[gid]; + const float4 moment = moments[gid]; - if (material[gid] == 1 && particle.z < 1.0) { + if (material[gid] == 1 && particle.w < 1.0) { particle.x += moment.y; particle.y += moment.z; - particle.z += min(particle.x, particle.y) * aging; +% if descriptor.d == 2: + particle.w += min(particle.x, particle.y) * aging; +% elif descriptor.d == 3: + particle.z += moment.w; + particle.w += min(min(particle.x, particle.y), particle.z) * aging; +% endif } else { - float2 init_particle = init_particles[pid]; - particle.x = init_particle.x; - particle.y = init_particle.y; - particle.z = particle.z-1.0; + particle.xyz = init_particles[pid].xyz; + particle.w = particle.w-1.0; } particles[pid] = particle; diff --git a/trugfeuer_2d_gl_interop.py b/trugfeuer_2d_gl_interop.py index cfdcc17..6d66efa 100644 --- a/trugfeuer_2d_gl_interop.py +++ b/trugfeuer_2d_gl_interop.py @@ -151,7 +151,7 @@ void main() { 1. ); - color = fire(1.0-particles[2]); + color = fire(1.0-particles[3]); }""").substitute({ 'size_x': lattice_x, 'size_y': lattice_y diff --git a/utility/particles.py b/utility/particles.py index 43f32a3..a75ed34 100644 --- a/utility/particles.py +++ b/utility/particles.py @@ -12,16 +12,22 @@ class Particles: self.count = len(grid) self.np_particles = numpy.ndarray(shape=(self.count, 4), dtype=float_type) - - self.np_particles[:,0:2] = grid - self.np_particles[:,2:4] = self.np_particles[:,0:2] - - self.np_init_particles = numpy.ndarray(shape=(self.count, 2), dtype=float_type) - self.np_init_particles[:,0:2] = grid - self.cl_init_particles = cl.Buffer(context, mf.READ_ONLY, size=self.count * 2*numpy.float32(0).nbytes) - cl.enqueue_copy(queue, self.cl_init_particles, self.np_init_particles).wait(); + self.np_init_particles = numpy.ndarray(shape=(self.count, 4), dtype=float_type) + + if len(grid[0,:]) == 2: + self.np_particles[:,0:2] = grid + self.np_particles[:,2:4] = 0 + self.np_init_particles[:,0:2] = grid + self.np_init_particles[:,2:4] = 0 + elif len(grid[0,:]) == 3: + self.np_particles[:,0:3] = grid + self.np_particles[:,3] = 0 + self.np_init_particles[:,0:3] = grid + self.np_init_particles[:,3] = 0 self.gl_particles = vbo.VBO(data=self.np_particles, usage=gl.GL_DYNAMIC_DRAW, target=gl.GL_ARRAY_BUFFER) self.gl_particles.bind() self.cl_gl_particles = cl.GLBuffer(self.context, mf.READ_WRITE, int(self.gl_particles)) + self.cl_init_particles = cl.Buffer(context, mf.READ_ONLY, size=self.count * 4*numpy.float32(0).nbytes) + cl.enqueue_copy(queue, self.cl_init_particles, self.np_init_particles).wait(); -- cgit v1.2.3