aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--ldc_3d_gl_interop.py185
-rw-r--r--template/kernel.mako23
-rw-r--r--trugfeuer_2d_gl_interop.py2
-rw-r--r--utility/particles.py22
4 files changed, 215 insertions, 17 deletions
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();