aboutsummaryrefslogtreecommitdiff
diff options
context:
space:
mode:
-rw-r--r--implosion.py275
1 files changed, 78 insertions, 197 deletions
diff --git a/implosion.py b/implosion.py
index d39a2ff..75b2fee 100644
--- a/implosion.py
+++ b/implosion.py
@@ -16,19 +16,28 @@ float constant w[9] = {
1./36 , 1./9., 1./36.
};
-uint2 cellAtGid(unsigned int gid)
-{
- const int y = gid / $nX;
- return (uint2)(gid - $nX*y, y);
+unsigned int indexOfDirection(int i, int j) {
+ return (i+1) + 3*(1-j);
}
-unsigned int gidOfCell(int x, int y)
+unsigned int indexOfCell(int x, int y)
{
return y * $nX + x;
}
-unsigned int indexOfDirection(int i, int j) {
- return 3*(i+1) + (j+1);
+unsigned int idx(int x, int y, int i, int j) {
+ return indexOfDirection(i,j)*$nX*$nY + indexOfCell(x,y);
+}
+
+uint2 cellAtIndex(unsigned int gid)
+{
+ const int y = gid / $nX;
+ return (uint2)(gid - $nX*y, y);
+}
+
+
+__global float* f_i(__global 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) {
@@ -39,37 +48,37 @@ float sq(float x) {
return x*x;
}
-float equilibrium(float d, float2 v, int i, int j) {
- return w[indexOfDirection(i,j)] * d * (1 + 3*comp(i,j,v) + 4.5*sq(comp(i,j,v)) - 1.5*dot(v,v));
+float density(__global const float* f, unsigned int gid) {
+ return f[0*$nX*$nY + gid]
+ + f[1*$nX*$nY + gid]
+ + f[2*$nX*$nY + gid]
+ + f[3*$nX*$nY + gid]
+ + f[4*$nX*$nY + gid]
+ + f[5*$nX*$nY + gid]
+ + f[6*$nX*$nY + gid]
+ + f[7*$nX*$nY + gid]
+ + f[8*$nX*$nY + gid];
+}
+
+float2 velocity(__global const float* f, float d, unsigned int gid)
+{
+ return (float2)(
+ (f[5*$nX*$nY+gid] - f[3*$nX*$nY+gid] + f[2*$nX*$nY+gid] - f[6*$nX*$nY+gid] + f[8*$nX*$nY+gid] - f[0*$nX*$nY+gid]) / d,
+ (f[1*$nX*$nY+gid] - f[7*$nX*$nY+gid] + f[2*$nX*$nY+gid] - f[6*$nX*$nY+gid] - f[8*$nX*$nY+gid] + f[0*$nX*$nY+gid]) / d
+ );
}
-float bgk(__global const float* pop, uint ngid, int i, int j, float d, float2 v) {
- return pop[ngid] + $tau * (equilibrium(d,v,i,j) - pop[ngid]);
+float f_eq(float d, float2 v, int i, int j) {
+ return w[indexOfDirection(i,j)] * d * (1 + 3*comp(i,j,v) + 4.5*sq(comp(i,j,v)) - 1.5*dot(v,v));
}
-__kernel void collide_and_stream(__global float* pop_a_0,
- __global float* pop_a_1,
- __global float* pop_a_2,
- __global float* pop_a_3,
- __global float* pop_a_4,
- __global float* pop_a_5,
- __global float* pop_a_6,
- __global float* pop_a_7,
- __global float* pop_a_8,
- __global const float* pop_b_0,
- __global const float* pop_b_1,
- __global const float* pop_b_2,
- __global const float* pop_b_3,
- __global const float* pop_b_4,
- __global const float* pop_b_5,
- __global const float* pop_b_6,
- __global const float* pop_b_7,
- __global const float* pop_b_8,
- __global float* moments,
+__kernel void collide_and_stream(__global float* f_a,
+ __global const float* f_b,
+ __global float* moments,
__global const int* material)
{
const unsigned int gid = get_global_id(0);
- const uint2 cell = cellAtGid(gid);
+ const uint2 cell = cellAtIndex(gid);
const int m = material[gid];
@@ -77,42 +86,19 @@ __kernel void collide_and_stream(__global float* pop_a_0,
return;
}
- const float d = pop_b_0[gid] + pop_b_1[gid] + pop_b_2[gid] + pop_b_3[gid] + pop_b_4[gid] + pop_b_5[gid] + pop_b_6[gid] + pop_b_7[gid] + pop_b_8[gid];
-
- const float2 v = (float2)(
- (pop_b_5[gid] - pop_b_3[gid] + pop_b_2[gid] - pop_b_6[gid] + pop_b_8[gid] - pop_b_0[gid]) / d,
- (pop_b_1[gid] - pop_b_7[gid] + pop_b_2[gid] - pop_b_6[gid] - pop_b_8[gid] + pop_b_0[gid]) / d
- );
+ const float d = density(f_b, gid);
+ const float2 v = velocity(f_b, d, gid);
- if ( m == 1 ) {
- pop_a_0[gid] = bgk(pop_b_0, gidOfCell(cell.x+1, cell.y-1), -1, 1, d, v);
- pop_a_1[gid] = bgk(pop_b_1, gidOfCell(cell.x , cell.y-1), 0, 1, d, v);
- pop_a_2[gid] = bgk(pop_b_2, gidOfCell(cell.x-1, cell.y-1), 1, 1, d, v);
-
- pop_a_3[gid] = bgk(pop_b_3, gidOfCell(cell.x+1, cell.y ), -1, 0, d, v);
- pop_a_4[gid] = bgk(pop_b_4, gidOfCell(cell.x , cell.y ), 0, 0, d, v);
- pop_a_5[gid] = bgk(pop_b_5, gidOfCell(cell.x-1, cell.y ), 1, 0, d, v);
-
- pop_a_6[gid] = bgk(pop_b_6, gidOfCell(cell.x+1, cell.y+1), -1,-1, d, v);
- pop_a_7[gid] = bgk(pop_b_7, gidOfCell(cell.x , cell.y+1), 0,-1, d, v);
- pop_a_8[gid] = bgk(pop_b_8, gidOfCell(cell.x-1, cell.y+1), 1,-1, d, v);
- } else {
- pop_a_8[gid] = bgk(pop_b_0, gidOfCell(cell.x+1, cell.y-1), -1, 1, d, v);
- pop_a_7[gid] = bgk(pop_b_1, gidOfCell(cell.x , cell.y-1), 0, 1, d, v);
- pop_a_6[gid] = bgk(pop_b_2, gidOfCell(cell.x-1, cell.y-1), 1, 1, d, v);
-
- pop_a_5[gid] = bgk(pop_b_3, gidOfCell(cell.x+1, cell.y ), -1, 0, d, v);
- pop_a_4[gid] = bgk(pop_b_4, gidOfCell(cell.x , cell.y ), 0, 0, d, v);
- pop_a_3[gid] = bgk(pop_b_5, gidOfCell(cell.x-1, cell.y ), 1, 0, d, v);
-
- pop_a_2[gid] = bgk(pop_b_6, gidOfCell(cell.x+1, cell.y+1), -1,-1, d, v);
- pop_a_1[gid] = bgk(pop_b_7, gidOfCell(cell.x , cell.y+1), 0,-1, d, v);
- pop_a_0[gid] = bgk(pop_b_8, gidOfCell(cell.x-1, cell.y+1), 1,-1, d, v);
+ for ( int i = -1; i <= 1; ++i ) {
+ for ( int j = 1; j >= -1; --j ) {
+ *f_i(f_a, cell.x, cell.y, m*i, m*j) = *f_i(f_b, cell.x-i, cell.y-j, i, j)
+ + $tau * (f_eq(d,v,i,j) - *f_i(f_b, cell.x-i, cell.y-j, i, j));
+ }
}
- moments[gid*3+0] = d;
- moments[gid*3+1] = v.x;
- moments[gid*3+2] = v.y;
+ moments[gid] = d;
+ moments[2*gid] = v.x;
+ moments[3*gid] = v.y;
}"""
class D2Q9_BGK_Lattice:
@@ -129,27 +115,10 @@ class D2Q9_BGK_Lattice:
self.context = cl.Context(properties=[(cl.context_properties.PLATFORM, self.platform)])
self.queue = cl.CommandQueue(self.context)
- self.np_pop_a_0 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_1 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_2 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_3 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_4 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_5 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_6 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_7 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_a_8 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
-
- self.np_pop_b_0 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_1 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_2 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_3 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_4 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_5 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_6 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_7 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
- self.np_pop_b_8 = numpy.ndarray(shape=(self.nCells, 1), dtype=numpy.float32)
-
- self.np_moments = numpy.ndarray(shape=(self.nCells, 3), dtype=numpy.float32)
+ 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()
@@ -157,28 +126,11 @@ class D2Q9_BGK_Lattice:
self.equilibrilize()
self.setup_anomaly()
- self.cl_pop_a_0 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_0)
- self.cl_pop_a_1 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_1)
- self.cl_pop_a_2 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_2)
- self.cl_pop_a_3 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_3)
- self.cl_pop_a_4 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_4)
- self.cl_pop_a_5 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_5)
- self.cl_pop_a_6 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_6)
- self.cl_pop_a_7 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_7)
- self.cl_pop_a_8 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_a_8)
-
- self.cl_pop_b_0 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_0)
- self.cl_pop_b_1 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_1)
- self.cl_pop_b_2 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_2)
- self.cl_pop_b_3 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_3)
- self.cl_pop_b_4 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_4)
- self.cl_pop_b_5 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_5)
- self.cl_pop_b_6 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_6)
- self.cl_pop_b_7 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_7)
- self.cl_pop_b_8 = cl.Buffer(self.context, mf.READ_WRITE | mf.USE_HOST_PTR, hostbuf=self.np_pop_b_8)
-
- self.cl_moments = cl.Buffer(self.context, mf.WRITE_ONLY | mf.USE_HOST_PTR, hostbuf=self.np_moments)
- self.cl_material = cl.Buffer(self.context, mf.READ_ONLY | mf.USE_HOST_PTR, hostbuf=self.np_material)
+ 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()
@@ -192,26 +144,13 @@ class D2Q9_BGK_Lattice:
self.np_material[self.idx(x,y)] = 1
def equilibrilize(self):
- self.np_pop_a_0[:] = 1./36.
- self.np_pop_a_1[:] = 1./9.
- self.np_pop_a_2[:] = 1./36.
- self.np_pop_a_3[:] = 1./9.
- self.np_pop_a_4[:] = 4./9.
- self.np_pop_a_5[:] = 1./9.
- self.np_pop_a_6[:] = 1./36
- self.np_pop_a_7[:] = 1./9.
- self.np_pop_a_8[:] = 1./36.
-
- self.np_pop_b_0[:] = 1./36.
- self.np_pop_b_1[:] = 1./9.
- self.np_pop_b_2[:] = 1./36.
- self.np_pop_b_3[:] = 1./9.
- self.np_pop_b_4[:] = 4./9.
- self.np_pop_b_5[:] = 1./9.
- self.np_pop_b_6[:] = 1./36
- self.np_pop_b_7[:] = 1./9.
- self.np_pop_b_8[:] = 1./36.
+ 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],
@@ -223,25 +162,8 @@ class D2Q9_BGK_Lattice:
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_0[self.idx(x,y)] = 1./24.
- self.np_pop_a_1[self.idx(x,y)] = 1./24.
- self.np_pop_a_2[self.idx(x,y)] = 1./24.
- self.np_pop_a_3[self.idx(x,y)] = 1./24.
- self.np_pop_a_4[self.idx(x,y)] = 1./24.
- self.np_pop_a_5[self.idx(x,y)] = 1./24.
- self.np_pop_a_6[self.idx(x,y)] = 1./24.
- self.np_pop_a_7[self.idx(x,y)] = 1./24.
- self.np_pop_a_8[self.idx(x,y)] = 1./24.
-
- self.np_pop_b_0[self.idx(x,y)] = 1./24.
- self.np_pop_b_1[self.idx(x,y)] = 1./24.
- self.np_pop_b_2[self.idx(x,y)] = 1./24.
- self.np_pop_b_3[self.idx(x,y)] = 1./24.
- self.np_pop_b_4[self.idx(x,y)] = 1./24.
- self.np_pop_b_5[self.idx(x,y)] = 1./24.
- self.np_pop_b_6[self.idx(x,y)] = 1./24.
- self.np_pop_b_7[self.idx(x,y)] = 1./24.
- self.np_pop_b_8[self.idx(x,y)] = 1./24.
+ 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({
@@ -253,76 +175,35 @@ class D2Q9_BGK_Lattice:
def evolve(self):
if self.tick:
self.tick = False
- self.program.collide_and_stream(self.queue, (self.nCells,), None,
- self.cl_pop_a_0,
- self.cl_pop_a_1,
- self.cl_pop_a_2,
- self.cl_pop_a_3,
- self.cl_pop_a_4,
- self.cl_pop_a_5,
- self.cl_pop_a_6,
- self.cl_pop_a_7,
- self.cl_pop_a_8,
- self.cl_pop_b_0,
- self.cl_pop_b_1,
- self.cl_pop_b_2,
- self.cl_pop_b_3,
- self.cl_pop_b_4,
- self.cl_pop_b_5,
- self.cl_pop_b_6,
- self.cl_pop_b_7,
- self.cl_pop_b_8,
- self.cl_moments,
- self.cl_material)
+ self.program.collide_and_stream(self.queue, (self.nCells,), None, self.cl_pop_a, self.cl_pop_b, self.cl_moments, self.cl_material)
self.queue.finish()
else:
self.tick = True
- self.program.collide_and_stream(self.queue, (self.nCells,), None,
- self.cl_pop_b_0,
- self.cl_pop_b_1,
- self.cl_pop_b_2,
- self.cl_pop_b_3,
- self.cl_pop_b_4,
- self.cl_pop_b_5,
- self.cl_pop_b_6,
- self.cl_pop_b_7,
- self.cl_pop_b_8,
- self.cl_pop_a_0,
- self.cl_pop_a_1,
- self.cl_pop_a_2,
- self.cl_pop_a_3,
- self.cl_pop_a_4,
- self.cl_pop_a_5,
- self.cl_pop_a_6,
- self.cl_pop_a_7,
- self.cl_pop_a_8,
- self.cl_moments,
- self.cl_material)
+ self.program.collide_and_stream(self.queue, (self.nCells,), None, self.cl_pop_b, self.cl_pop_a, self.cl_moments, self.cl_material)
self.queue.finish()
def show(self, i):
- cl.enqueue_copy(self.queue, self.np_moments, self.cl_moments).wait();
-
- density = numpy.ndarray(shape=(self.nX, self.nY))
+ cl.enqueue_copy(LBM.queue, LBM.np_moments, LBM.cl_moments).wait();
- for y in range(0,self.nY-1):
- for x in range(0,self.nX-1):
- density[x,y] = self.np_moments[self.idx(x,y),0]
+ 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, cmap=plt.get_cmap("seismic"))
+ plt.imshow(density, vmin=0.2, vmax=2.0, cmap=plt.get_cmap("seismic"))
plt.savefig("result/density_" + str(i) + ".png")
def MLUPS(cells, steps, time):
return ((cells*steps) / time) / 1000000
-LBM = D2Q9_BGK_Lattice(1000, 1000)
+LBM = D2Q9_BGK_Lattice(2000, 2000)
-nUpdates = 10000
+nUpdates = 100
start = timer()
-for i in range(1,nUpdates):
+for i in range(0,nUpdates):
LBM.evolve()
end = timer()