diff options
-rw-r--r-- | D2Q9.py (renamed from codegen_lbm.py) | 60 | ||||
-rw-r--r-- | implosion.py | 242 | ||||
-rw-r--r-- | notebook/inspect_opencl_layout.ipynb (renamed from inspect_opencl_layout.ipynb) | 0 | ||||
-rw-r--r-- | notebook/lbm_codegen.ipynb (renamed from lbm_codegen.ipynb) | 0 | ||||
-rw-r--r-- | notebook/lbm_codegen_d3q19.ipynb (renamed from lbm_codegen_d3q19.ipynb) | 0 | ||||
-rw-r--r-- | symbolic/D2Q9.py (renamed from lbm_d2q9.py) | 0 |
6 files changed, 38 insertions, 264 deletions
@@ -2,18 +2,13 @@ import pyopencl as cl mf = cl.mem_flags import numpy -import time - -import matplotlib -import matplotlib.pyplot as plt -matplotlib.use('AGG') import sympy -import lbm_d2q9 as D2Q9 +import symbolic.D2Q9 as D2Q9 from mako.template import Template -class D2Q9_BGK_Lattice: +class Lattice: def idx(self, x, y): return y * self.nX + x; @@ -79,54 +74,5 @@ class D2Q9_BGK_Lattice: self.program.collect_moments(self.queue, (self.nX,self.nY), (32,1), self.cl_pop_b, self.cl_moments) else: self.program.collect_moments(self.queue, (self.nX,self.nY), (32,1), self.cl_pop_a, self.cl_moments) - cl.enqueue_copy(LBM.queue, moments, LBM.cl_moments).wait(); + cl.enqueue_copy(self.queue, moments, self.cl_moments).wait(); return moments - - -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))) - - density = numpy.ndarray(shape=(lattice.nY-2, lattice.nX-2)) - for y in range(1,lattice.nY-1): - for x in range(1,lattice.nX-1): - density[y-1,x-1] = m[0,lattice.idx(x,y)] - - plt.figure(figsize=(10, 10)) - plt.imshow(density, origin='lower', vmin=0.2, vmax=2.0, cmap=plt.get_cmap('seismic')) - plt.savefig("result/density_" + str(i) + ".png", bbox_inches='tight', pad_inches=0) - -def box(nX, nY, x, y): - if x == 1 or y == 1 or x == nX-2 or y == nY-2: - return 2 - else: - return 1 - -nUpdates = 1000 -nStat = 100 - -moments = [] - -print("Initializing simulation...\n") - -LBM = D2Q9_BGK_Lattice(nX = 1024, nY = 1024, tau = 0.8, geometry = box) - -print("Starting simulation using %d cells...\n" % LBM.nCells) - -lastStat = time.time() - -for i in range(1,nUpdates+1): - LBM.evolve() - - if i % nStat == 0: - LBM.sync() - print("i = %4d; %3.0f MLUPS" % (i, MLUPS(LBM.nCells, nStat, time.time() - lastStat))) - moments.append(LBM.get_moments()) - lastStat = time.time() - -print("\nConcluded simulation.\n") - -generate_moment_plots(LBM, moments) diff --git a/implosion.py b/implosion.py index c70f21a..75e9640 100644 --- a/implosion.py +++ b/implosion.py @@ -1,228 +1,56 @@ -import pyopencl as cl -mf = cl.mem_flags - -from string import Template - import numpy -import matplotlib.pyplot as plt - import time -kernel = """ -float constant w[9] = { - 1./36., 1./9., 1./36., - 1./9. , 4./9., 1./9. , - 1./36 , 1./9., 1./36. -}; - -unsigned int indexOfDirection(int i, int j) { - return (i+1) + 3*(1-j); -} - -unsigned int indexOfCell(int x, int y) -{ - return y * $nX + x; -} - -unsigned int idx(int x, int y, int i, int j) { - return indexOfDirection(i,j)*$nCells + indexOfCell(x,y); -} - -__global float f_i(__global __read_only float* f, int x, int y, int i, int j) { - return f[idx(x,y,i,j)]; -} - -float comp(int i, int j, float2 v) { - return i*v.x + j*v.y; -} - -float sq(float x) { - return x*x; -} - -float f_eq(float w, float d, float2 v, int i, int j, float dotv) { - return w * d * (1.f + 3.f*comp(i,j,v) + 4.5f*sq(comp(i,j,v)) - 1.5f*dotv); -} - -__kernel void collide_and_stream(__global __write_only float* f_a, - __global __read_only float* f_b, - __global __write_only float* moments, - __global __read_only int* material) -{ - const unsigned int gid = indexOfCell(get_global_id(0), get_global_id(1)); - - const uint2 cell = (uint2)(get_global_id(0), get_global_id(1)); - - const int m = material[gid]; - - if ( m == 0 ) { - return; - } - - float f0 = f_i(f_b, cell.x+1, cell.y-1, -1, 1); - float f1 = f_i(f_b, cell.x , cell.y-1, 0, 1); - float f2 = f_i(f_b, cell.x-1, cell.y-1, 1, 1); - float f3 = f_i(f_b, cell.x+1, cell.y , -1, 0); - float f4 = f_i(f_b, cell.x , cell.y , 0, 0); - float f5 = f_i(f_b, cell.x-1, cell.y , 1, 0); - float f6 = f_i(f_b, cell.x+1, cell.y+1, -1,-1); - float f7 = f_i(f_b, cell.x , cell.y+1, 0,-1); - float f8 = f_i(f_b, cell.x-1, cell.y+1, 1,-1); - - const float d = f0 + f1 + f2 + f3 + f4 + f5 + f6 + f7 + f8; - - float2 v = (float2)( - (f5 - f3 + f2 - f6 + f8 - f0) / d, - (f1 - f7 + f2 - f6 - f8 + f0) / d - ); - - if ( m == 2 ) { - v = (float2)(0.0f, 0.0f); - } - - const float dotv = dot(v,v); - - f0 += $omega * (f_eq(w[0], d,v,-1, 1, dotv) - f0); - f1 += $omega * (f_eq(w[1], d,v, 0, 1, dotv) - f1); - f2 += $omega * (f_eq(w[2], d,v, 1, 1, dotv) - f2); - f3 += $omega * (f_eq(w[3], d,v,-1, 0, dotv) - f3); - f4 += $omega * (f_eq(w[4], d,v, 0, 0, dotv) - f4); - f5 += $omega * (f_eq(w[5], d,v, 1, 0, dotv) - f5); - f6 += $omega * (f_eq(w[6], d,v,-1,-1, dotv) - f6); - f7 += $omega * (f_eq(w[7], d,v, 0,-1, dotv) - f7); - f8 += $omega * (f_eq(w[8], d,v, 1,-1, dotv) - f8); - - f_a[0*$nCells + gid] = f0; - f_a[1*$nCells + gid] = f1; - f_a[2*$nCells + gid] = f2; - f_a[3*$nCells + gid] = f3; - f_a[4*$nCells + gid] = f4; - f_a[5*$nCells + gid] = f5; - f_a[6*$nCells + gid] = f6; - f_a[7*$nCells + gid] = f7; - f_a[8*$nCells + gid] = f8; - - moments[1*gid] = d; - moments[2*gid] = v.x; - moments[3*gid] = v.y; -}""" - -class D2Q9_BGK_Lattice: - def idx(self, x, y): - return y * self.nX + x; - - def __init__(self, nX, nY): - self.nX = nX - self.nY = nY - self.nCells = nX * nY - self.tick = True - - self.platform = cl.get_platforms()[0] - self.context = cl.Context(properties=[(cl.context_properties.PLATFORM, self.platform)]) - self.queue = cl.CommandQueue(self.context) - - self.np_pop_a = numpy.ndarray(shape=(9, self.nCells), dtype=numpy.float32) - self.np_pop_b = numpy.ndarray(shape=(9, self.nCells), dtype=numpy.float32) - - self.np_moments = numpy.ndarray(shape=(3, self.nCells), dtype=numpy.float32) - self.np_material = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.int32) - - self.setup_geometry() - - self.equilibrilize() - self.setup_anomaly() - - self.cl_pop_a = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a) - self.cl_pop_b = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b) - - self.cl_material = cl.Buffer(self.context, mf.READ_ONLY | mf.USE_HOST_PTR, hostbuf=self.np_material) - self.cl_moments = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_moments) - - self.build_kernel() - - def setup_geometry(self): - self.np_material[:] = 0 - for x in range(1,self.nX-1): - for y in range(1,self.nY-1): - if x == 1 or y == 1 or x == self.nX-2 or y == self.nY-2: - self.np_material[self.idx(x,y)] = 2 - else: - self.np_material[self.idx(x,y)] = 1 - - def equilibrilize(self): - self.np_pop_a[(0,2,6,8),:] = 1./36. - self.np_pop_a[(1,3,5,7),:] = 1./9. - self.np_pop_a[4,:] = 4./9. - - self.np_pop_b[(0,2,6,8),:] = 1./36. - self.np_pop_b[(1,3,5,7),:] = 1./9. - self.np_pop_b[4,:] = 4./9. - - def setup_anomaly(self): - bubbles = [ [ self.nX//4, self.nY//4], - [ self.nX//4,self.nY-self.nY//4], - [self.nX-self.nX//4, self.nY//4], - [self.nX-self.nX//4,self.nY-self.nY//4] ] - - for x in range(0,self.nX-1): - for y in range(0,self.nY-1): - for [a,b] in bubbles: - if numpy.sqrt((x-a)*(x-a)+(y-b)*(y-b)) < self.nX//10: - self.np_pop_a[:,self.idx(x,y)] = 1./24. - self.np_pop_b[:,self.idx(x,y)] = 1./24. - - def build_kernel(self): - self.program = cl.Program(self.context, Template(kernel).substitute({ - 'nX' : self.nX, - 'nY' : self.nY, - 'nCells': self.nCells, - 'omega': 1.0/0.8 - })).build() #'-cl-single-precision-constant -cl-fast-relaxed-math') - - def evolve(self): - if self.tick: - self.tick = False - self.program.collide_and_stream(self.queue, (self.nX,self.nY), (64,1), self.cl_pop_a, self.cl_pop_b, self.cl_moments, self.cl_material) - else: - self.tick = True - self.program.collide_and_stream(self.queue, (self.nX,self.nY), (64,1), self.cl_pop_b, self.cl_pop_a, self.cl_moments, self.cl_material) - - def sync(self): - self.queue.finish() - - def show(self, i): - cl.enqueue_copy(LBM.queue, LBM.np_moments, LBM.cl_moments).wait(); - - density = numpy.ndarray(shape=(self.nX-2, self.nY-2)) - for y in range(1,self.nY-1): - for x in range(1,self.nX-1): - density[x-1,y-1] = self.np_moments[0,self.idx(x,y)] - - plt.imshow(density, vmin=0.2, vmax=2.0, cmap=plt.get_cmap("seismic")) - plt.savefig("result/density_" + str(i) + ".png") +import matplotlib +import matplotlib.pyplot as plt +matplotlib.use('AGG') +from D2Q9 import Lattice 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))) + + density = numpy.ndarray(shape=(lattice.nY-2, lattice.nX-2)) + for y in range(1,lattice.nY-1): + for x in range(1,lattice.nX-1): + density[y-1,x-1] = m[0,lattice.idx(x,y)] + + plt.figure(figsize=(10, 10)) + plt.imshow(density, origin='lower', vmin=0.2, vmax=2.0, cmap=plt.get_cmap('seismic')) + plt.savefig("result/density_" + str(i) + ".png", bbox_inches='tight', pad_inches=0) + +def box(nX, nY, x, y): + if x == 1 or y == 1 or x == nX-2 or y == nY-2: + return 2 + else: + return 1 + nUpdates = 1000 -nStat = 100 +nStat = 100 + +moments = [] print("Initializing simulation...\n") -LBM = D2Q9_BGK_Lattice(1024, 1024) +lattice = Lattice(nX = 1024, nY = 1024, tau = 0.8, geometry = box) -print("Starting simulation using %d cells...\n" % LBM.nCells) +print("Starting simulation using %d cells...\n" % lattice.nCells) lastStat = time.time() for i in range(1,nUpdates+1): + lattice.evolve() + if i % nStat == 0: - LBM.sync() - #LBM.show(i) - print("i = %4d; %3.0f MLUPS" % (i, MLUPS(LBM.nCells, nStat, time.time() - lastStat))) + lattice.sync() + print("i = %4d; %3.0f MLUPS" % (i, MLUPS(lattice.nCells, nStat, time.time() - lastStat))) + moments.append(lattice.get_moments()) lastStat = time.time() - LBM.evolve() +print("\nConcluded simulation.\n") -LBM.show(nUpdates) +generate_moment_plots(lattice, moments) diff --git a/inspect_opencl_layout.ipynb b/notebook/inspect_opencl_layout.ipynb index 521f93a..521f93a 100644 --- a/inspect_opencl_layout.ipynb +++ b/notebook/inspect_opencl_layout.ipynb diff --git a/lbm_codegen.ipynb b/notebook/lbm_codegen.ipynb index e1593c0..e1593c0 100644 --- a/lbm_codegen.ipynb +++ b/notebook/lbm_codegen.ipynb diff --git a/lbm_codegen_d3q19.ipynb b/notebook/lbm_codegen_d3q19.ipynb index face58f..face58f 100644 --- a/lbm_codegen_d3q19.ipynb +++ b/notebook/lbm_codegen_d3q19.ipynb diff --git a/lbm_d2q9.py b/symbolic/D2Q9.py index a3c2503..a3c2503 100644 --- a/lbm_d2q9.py +++ b/symbolic/D2Q9.py |