Collapse SOA into single array

Weirdly the expected performance gains due to better coalescence of memory access is not achieved.

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()