aboutsummaryrefslogtreecommitdiff
path: root/lid_driven_cavity/opencl
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-11-12 22:59:21 +0100
committerAdrian Kummerlaender2019-11-12 22:59:21 +0100
commit8c1f317010eeba28b3d4ae3ef158b1e760997ed1 (patch)
treef2aa1e5e2128bfdd108c51c02efa6c7a3be83e87 /lid_driven_cavity/opencl
parent2dba6db0c94ccb260f3e9345b25146ce75ccea9a (diff)
downloadboltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar.gz
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar.bz2
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar.lz
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar.xz
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.tar.zst
boltzgen_examples-8c1f317010eeba28b3d4ae3ef158b1e760997ed1.zip
Share Lattice implementation between plain and interop OpenCL example
Diffstat (limited to 'lid_driven_cavity/opencl')
-rw-r--r--lid_driven_cavity/opencl/AA.py28
-rw-r--r--lid_driven_cavity/opencl/AB.py34
-rw-r--r--lid_driven_cavity/opencl/ldc_2d.py42
3 files changed, 37 insertions, 67 deletions
diff --git a/lid_driven_cavity/opencl/AA.py b/lid_driven_cavity/opencl/AA.py
index 74979c5..7ea3417 100644
--- a/lid_driven_cavity/opencl/AA.py
+++ b/lid_driven_cavity/opencl/AA.py
@@ -24,7 +24,7 @@ class Memory:
return z * (self.size_x*self.size_y) + y * self.size_x + x;
class Lattice:
- def __init__(self, geometry, kernel_src, descriptor, platform = 0, precision = 'single'):
+ def __init__(self, geometry, kernel_src, descriptor, context, queue, precision = 'single'):
self.geometry = geometry
self.descriptor = descriptor
@@ -33,13 +33,10 @@ class Lattice:
'double': (numpy.float64, 'double'),
}.get(precision, None)
- self.platform = cl.get_platforms()[platform]
self.layout = None
- self.context = cl.Context(
- properties=[(cl.context_properties.PLATFORM, self.platform)])
-
- self.queue = cl.CommandQueue(self.context)
+ self.context = context
+ self.queue = queue
self.memory = Memory(descriptor, self.geometry, self.context, self.float_type[0])
self.tick = False
@@ -51,7 +48,7 @@ class Lattice:
self.build_kernel(kernel_src)
- self.program.equilibrilize_all(
+ self.program.equilibrilize_domain_tick(
self.queue, self.geometry.size(), self.layout, self.memory.cl_pop).wait()
self.tick_tasks = []
@@ -82,24 +79,9 @@ class Lattice:
def get_moments(self):
moments = numpy.ndarray(shape=(self.memory.volume*(self.descriptor.d+1),1), dtype=self.float_type[0])
- self.program.collect_moments_all(
+ self.program.collect_moments_domain_tock(
self.queue, self.geometry.inner_size(), self.layout, self.memory.cl_pop, self.memory.cl_moments)
cl.enqueue_copy(self.queue, moments, self.memory.cl_moments).wait();
return moments
-
-HelperTemplate = """
-__kernel void equilibrilize_all(__global ${float_type}* f_next)
-{
- const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
- equilibrilize_tick(f_next, gid);
-}
-
-__kernel void collect_moments_all(__global ${float_type}* f,
- __global ${float_type}* moments)
-{
- const unsigned int gid = ${index.gid('get_global_id(0)+1', 'get_global_id(1)+1')};
- collect_moments_tock(f, gid, moments);
-}
-"""
diff --git a/lid_driven_cavity/opencl/AB.py b/lid_driven_cavity/opencl/AB.py
index 7cbf1c6..797bc83 100644
--- a/lid_driven_cavity/opencl/AB.py
+++ b/lid_driven_cavity/opencl/AB.py
@@ -25,7 +25,7 @@ class Memory:
return z * (self.size_x*self.size_y) + y * self.size_x + x;
class Lattice:
- def __init__(self, geometry, kernel_src, descriptor, platform = 0, precision = 'single'):
+ def __init__(self, geometry, kernel_src, descriptor, context, queue, precision = 'single'):
self.geometry = geometry
self.descriptor = descriptor
@@ -34,13 +34,10 @@ class Lattice:
'double': (numpy.float64, 'double'),
}.get(precision, None)
- self.platform = cl.get_platforms()[platform]
self.layout = None
- self.context = cl.Context(
- properties=[(cl.context_properties.PLATFORM, self.platform)])
-
- self.queue = cl.CommandQueue(self.context)
+ self.context = context
+ self.queue = queue
self.memory = Memory(descriptor, self.geometry, self.context, self.float_type[0])
self.tick = False
@@ -52,8 +49,10 @@ class Lattice:
self.build_kernel(kernel_src)
- self.program.equilibrilize_all(
+ self.program.equilibrilize_domain(
self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_a, self.memory.cl_pop_b).wait()
+ self.program.equilibrilize_domain(
+ self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_b, self.memory.cl_pop_a).wait()
self.tasks = []
@@ -80,29 +79,12 @@ class Lattice:
moments = numpy.ndarray(shape=(self.memory.volume*(self.descriptor.d+1),1), dtype=self.float_type[0])
if self.tick:
- self.program.collect_moments_all(
+ self.program.collect_moments_domain(
self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_b, self.memory.cl_moments)
else:
- self.program.collect_moments_all(
+ self.program.collect_moments_domain(
self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_a, self.memory.cl_moments)
cl.enqueue_copy(self.queue, moments, self.memory.cl_moments).wait();
return moments
-
-HelperTemplate = """
-__kernel void equilibrilize_all(__global ${float_type}* f_next,
- __global ${float_type}* f_prev)
-{
- const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
- equilibrilize(f_next, f_prev, gid);
- equilibrilize(f_prev, f_next, gid);
-}
-
-__kernel void collect_moments_all(__global ${float_type}* f,
- __global ${float_type}* moments)
-{
- const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
- collect_moments(f, gid, moments);
-}
-"""
diff --git a/lid_driven_cavity/opencl/ldc_2d.py b/lid_driven_cavity/opencl/ldc_2d.py
index 7380aa9..1ea80aa 100644
--- a/lid_driven_cavity/opencl/ldc_2d.py
+++ b/lid_driven_cavity/opencl/ldc_2d.py
@@ -1,16 +1,18 @@
import numpy
import time
+import pyopencl as cl
+
from boltzgen import Generator, Geometry
from boltzgen.lbm.lattice import D2Q9
from boltzgen.lbm.model import BGK
from common import CellList, generate_moment_plots
-nUpdates = 100000
-nStat = 10000
+nUpdates = 20000
+nStat = 1000
-geometry = Geometry(256, 256)
+geometry = Geometry(512, 512)
print("Generating kernel using boltzgen...\n")
@@ -24,7 +26,6 @@ import AA
import AB
Lattice = eval('%s.Lattice' % streaming)
-HelperTemplate = eval('%s.HelperTemplate' % streaming)
def MLUPS(cells, steps, time):
return cells * steps / time * 1e-6
@@ -37,12 +38,17 @@ generator = Generator(
index = 'ZYX',
layout = 'SOA')
-kernel_src = generator.kernel(geometry, functions, extras)
-kernel_src += generator.custom(geometry, HelperTemplate)
+kernel_src = generator.kernel(geometry, functions, extras)
+
+print("Constructing OpenCL context...\n")
+
+cl_platform = cl.get_platforms()[0]
+cl_context = cl.Context(properties=[(cl.context_properties.PLATFORM, cl_platform)])
+cl_queue = cl.CommandQueue(cl_context)
print("Initializing simulation...\n")
-lattice = Lattice(geometry, kernel_src, D2Q9, precision = precision)
+lattice = Lattice(geometry, kernel_src, D2Q9, cl_context, cl_queue, precision = precision)
gid = lattice.memory.gid
ghost_cells = CellList(lattice.context, lattice.queue, lattice.float_type,
@@ -55,19 +61,19 @@ lid_cells = CellList(lattice.context, lattice.queue, lattice.float_type,
[ gid(x,y) for x, y in geometry.inner_cells() if y == geometry.size_y-2 ])
if streaming == 'AB':
- lattice.schedule('collide_and_stream_cells', bulk_cells)
- lattice.schedule('velocity_momenta_boundary_cells', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
- lattice.schedule('velocity_momenta_boundary_cells', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
+ lattice.schedule('collide_and_stream', bulk_cells)
+ lattice.schedule('velocity_momenta_boundary', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
+ lattice.schedule('velocity_momenta_boundary', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
elif streaming == 'AA':
- lattice.schedule_tick('collide_and_stream_cells_tick', bulk_cells)
- lattice.schedule_tick('velocity_momenta_boundary_cells_tick', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
- lattice.schedule_tick('velocity_momenta_boundary_cells_tick', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
-
- lattice.schedule_tock('equilibrilize_cells_tick', ghost_cells)
- lattice.schedule_tock('collide_and_stream_cells_tock', bulk_cells)
- lattice.schedule_tock('velocity_momenta_boundary_cells_tock', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
- lattice.schedule_tock('velocity_momenta_boundary_cells_tock', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
+ lattice.schedule_tick('collide_and_stream_tick', bulk_cells)
+ lattice.schedule_tick('velocity_momenta_boundary_tick', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
+ lattice.schedule_tick('velocity_momenta_boundary_tick', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
+
+ lattice.schedule_tock('equilibrilize_tick', ghost_cells)
+ lattice.schedule_tock('collide_and_stream_tock', bulk_cells)
+ lattice.schedule_tock('velocity_momenta_boundary_tock', wall_cells, numpy.array([0.0, 0.0], dtype=lattice.float_type[0]))
+ lattice.schedule_tock('velocity_momenta_boundary_tock', lid_cells, numpy.array([0.1, 0.0], dtype=lattice.float_type[0]))
print("Starting simulation using %d cells...\n" % lattice.geometry.volume)