aboutsummaryrefslogtreecommitdiff
path: root/ldc_2d
diff options
context:
space:
mode:
authorAdrian Kummerlaender2019-10-29 20:46:21 +0100
committerAdrian Kummerlaender2019-10-29 20:50:49 +0100
commit6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3 (patch)
tree4a9b08b021471f4ff0009785ea5a5ed92443e066 /ldc_2d
parent39aa2832b3d8e1a3e1dd2efb0bf9629d8190ed39 (diff)
downloadboltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar.gz
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar.bz2
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar.lz
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar.xz
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.tar.zst
boltzgen_examples-6bf667eeb0c8f40e0f6a28df1ecf4a40824771c3.zip
Move example into subfolder
Diffstat (limited to 'ldc_2d')
-rw-r--r--ldc_2d/opencl/ldc_2d.py101
-rw-r--r--ldc_2d/opencl/shell.nix45
-rw-r--r--ldc_2d/opencl/simulation.py105
3 files changed, 251 insertions, 0 deletions
diff --git a/ldc_2d/opencl/ldc_2d.py b/ldc_2d/opencl/ldc_2d.py
new file mode 100644
index 0000000..7ca7252
--- /dev/null
+++ b/ldc_2d/opencl/ldc_2d.py
@@ -0,0 +1,101 @@
+import numpy
+import time
+
+import matplotlib
+matplotlib.use('AGG')
+import matplotlib.pyplot as plt
+
+from boltzgen import LBM, Generator, Geometry
+from boltzgen.lbm.model import D2Q9
+
+from simulation import Lattice, CellList
+
+def MLUPS(cells, steps, time):
+ return cells * steps / time * 1e-6
+
+def generate_moment_plots(lattice, moments):
+ for i, m in enumerate(moments):
+ print("Generating plot %d of %d." % (i+1, len(moments)))
+
+ gid = lattice.memory.gid
+ velocity = numpy.reshape(
+ [ numpy.sqrt(m[gid(x,y)*3+1]**2 + m[gid(x,y)*3+2]**2) for x, y in lattice.geometry.inner_cells() ],
+ lattice.geometry.inner_size())
+
+ plt.figure(figsize=(10, 10))
+ plt.imshow(velocity, origin='lower', cmap=plt.get_cmap('seismic'))
+ plt.savefig("result/ldc_2d_%02d.png" % i, bbox_inches='tight', pad_inches=0)
+
+nUpdates = 100000
+nStat = 10000
+
+geometry = Geometry(512, 512)
+
+print("Generating kernel using boltzgen...\n")
+
+functions = ['collide_and_stream', 'equilibrilize', 'collect_moments', 'momenta_boundary']
+extras = ['cell_list_dispatch']
+
+precision = 'single'
+
+lbm = LBM(D2Q9)
+generator = Generator(
+ descriptor = D2Q9,
+ moments = lbm.moments(),
+ collision = lbm.bgk(f_eq = lbm.equilibrium(), tau = 0.6),
+ target = 'cl',
+ precision = precision,
+ index = 'ZYX',
+ layout = 'SOA')
+
+kernel_src = generator.kernel(geometry, functions, extras)
+kernel_src += generator.custom(geometry, """
+__kernel void equilibrilize(__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_gid(f_next, f_prev, gid);
+}
+
+__kernel void collect_moments(__global ${float_type}* f,
+ __global ${float_type}* moments)
+{
+ const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
+ collect_moments_gid(f, moments, gid);
+}
+""")
+
+print("Initializing simulation...\n")
+
+lattice = Lattice(geometry, kernel_src, D2Q9, precision = precision)
+gid = lattice.memory.gid
+
+bulk_cells = CellList(lattice.context, lattice.queue, lattice.float_type,
+ [ gid(x,y) for x, y in geometry.inner_cells() if x > 1 and x < geometry.size_x-2 and y > 1 and y < geometry.size_y-2 ])
+wall_cells = CellList(lattice.context, lattice.queue, lattice.float_type,
+ [ gid(x,y) for x, y in geometry.inner_cells() if x == 1 or y == 1 or x == geometry.size_x-2 ])
+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 ])
+
+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]))
+
+print("Starting simulation using %d cells...\n" % lattice.geometry.volume)
+
+moments = []
+
+lastStat = time.time()
+
+for i in range(1,nUpdates+1):
+ lattice.evolve()
+
+ if i % nStat == 0:
+ lattice.sync()
+ print("i = %4d; %3.0f MLUPS" % (i, MLUPS(lattice.geometry.volume, nStat, time.time() - lastStat)))
+ moments.append(lattice.get_moments())
+ lastStat = time.time()
+
+print("\nConcluded simulation.\n")
+
+generate_moment_plots(lattice, moments)
diff --git a/ldc_2d/opencl/shell.nix b/ldc_2d/opencl/shell.nix
new file mode 100644
index 0000000..462394e
--- /dev/null
+++ b/ldc_2d/opencl/shell.nix
@@ -0,0 +1,45 @@
+{ pkgs ? import <nixpkgs> { }, ... }:
+
+pkgs.stdenvNoCC.mkDerivation rec {
+ name = "pycl-env";
+ env = pkgs.buildEnv { name = name; paths = buildInputs; };
+
+ buildInputs = let
+ boltzgen = pkgs.python3.pkgs.buildPythonPackage rec {
+ pname = "boltzgen";
+ version = "0.1";
+
+ src = pkgs.fetchFromGitHub {
+ owner = "KnairdA";
+ repo = "boltzgen";
+ rev = "v0.1.1";
+ sha256 = "03fv7krhgc43gfjill8wb4aafr8xi69i2yh9zr68knnmrkrb8vpi";
+ };
+
+ propagatedBuildInputs = with pkgs.python37Packages; [
+ sympy
+ numpy
+ Mako
+ ];
+
+ doCheck = false;
+ };
+
+ local-python = pkgs.python3.withPackages (python-packages: with python-packages; [
+ boltzgen
+ numpy
+ pyopencl setuptools
+ matplotlib
+ ]);
+
+ in [
+ local-python
+ pkgs.opencl-info
+ ];
+
+ shellHook = ''
+ export NIX_SHELL_NAME="${name}"
+ export PYOPENCL_COMPILER_OUTPUT=1
+ export PYTHONPATH="$PWD:$PYTHONPATH"
+ '';
+}
diff --git a/ldc_2d/opencl/simulation.py b/ldc_2d/opencl/simulation.py
new file mode 100644
index 0000000..7625609
--- /dev/null
+++ b/ldc_2d/opencl/simulation.py
@@ -0,0 +1,105 @@
+import pyopencl as cl
+mf = cl.mem_flags
+
+import numpy
+
+class Memory:
+ def __init__(self, descriptor, geometry, context, float_type):
+ self.context = context
+ self.float_type = float_type
+
+ self.size_x = geometry.size_x
+ self.size_y = geometry.size_y
+ self.size_z = geometry.size_z
+ self.volume = self.size_x * self.size_y * self.size_z
+
+ self.pop_size = descriptor.q * self.volume * self.float_type(0).nbytes
+ self.moments_size = 3 * self.volume * self.float_type(0).nbytes
+
+ 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)
+
+ def gid(self, x, y, z = 0):
+ return z * (self.size_x*self.size_y) + y * self.size_x + x;
+
+class CellList:
+ def __init__(self, context, queue, float_type, cells):
+ self.cl_cells = cl.Buffer(context, mf.READ_ONLY, size=len(cells) * numpy.uint32(0).nbytes)
+ self.np_cells = numpy.ndarray(shape=(len(cells), 1), dtype=numpy.uint32)
+ self.np_cells[:,0] = cells[:]
+
+ cl.enqueue_copy(queue, self.cl_cells, self.np_cells).wait();
+
+ def get(self):
+ return self.cl_cells
+
+ def size(self):
+ return (len(self.np_cells), 1, 1)
+
+class Lattice:
+ def __init__(self, geometry, kernel_src, descriptor, platform = 0, precision = 'single'):
+ self.geometry = geometry
+ self.descriptor = descriptor
+
+ self.float_type = {
+ 'single': (numpy.float32, 'float'),
+ '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.memory = Memory(descriptor, self.geometry, self.context, self.float_type[0])
+ self.tick = False
+
+ self.compiler_args = {
+ 'single': '-cl-single-precision-constant -cl-fast-relaxed-math',
+ 'double': '-cl-fast-relaxed-math'
+ }.get(precision, None)
+
+ self.build_kernel(kernel_src)
+
+ self.program.equilibrilize(
+ self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_a, self.memory.cl_pop_b).wait()
+
+ self.tasks = []
+
+ def build_kernel(self, src):
+ self.program = cl.Program(self.context, src).build(self.compiler_args)
+
+ def schedule(self, f, cells, *params):
+ self.tasks += [ (eval("self.program.%s" % f), cells, params) ]
+
+ def evolve(self):
+ if self.tick:
+ self.tick = False
+ for f, cells, params in self.tasks:
+ f(self.queue, cells.size(), self.layout, self.memory.cl_pop_a, self.memory.cl_pop_b, cells.get(), *params)
+ else:
+ self.tick = True
+ for f, cells, params in self.tasks:
+ f(self.queue, cells.size(), self.layout, self.memory.cl_pop_b, self.memory.cl_pop_a, cells.get(), *params)
+
+ def sync(self):
+ self.queue.finish()
+
+ def get_moments(self):
+ moments = numpy.ndarray(shape=(self.memory.volume*(self.descriptor.d+1),1), dtype=self.float_type[0])
+
+ if self.tick:
+ self.program.collect_moments(
+ self.queue, self.geometry.size(), self.layout, self.memory.cl_pop_b, self.memory.cl_moments)
+ else:
+ self.program.collect_moments(
+ 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