bulk commit

[LbmBenchmarkKernelsPublic.git] / src / BenchKernelD3Q19ListAaPvCommon.c

diff --git a/src/BenchKernelD3Q19ListAaPvCommon.c b/src/BenchKernelD3Q19ListAaPvCommon.c
index 9cdf0e13ab057c297da8aa0053b5fe27d8445f91..7b479cfa62aea674796eb98a458d2b642af7a9b6 100644 (file)
--- a/src/BenchKernelD3Q19ListAaPvCommon.c
+++ b/src/BenchKernelD3Q19ListAaPvCommon.c
@@ -28,6 +28,7 @@
 
 #include "Memory.h"
 #include "Vtk.h"
+#include "Padding.h"
 
 #include <math.h>
 
@@ -225,15 +226,19 @@ static void ParameterUsage()
 {
        printf("Kernel parameters:\n");
        printf("  [-blk <n>] [-blk-[xyz] <n>]\n");
+#ifdef DATA_LAYOUT_SOA
+       printf("  [-pad auto|modulus_1+offset_1(,modulus_n+offset_n)*]\n");
+#endif
 
        return;
 }
 
-static void ParseParameters(Parameters * params, int * blk)
+static void ParseParameters(Parameters * params, int * blk, PadInfo ** padInfo)
 {
        Assert(blk != NULL);
 
        blk[0] = 0; blk[1] = 0; blk[2] = 0;
+       *padInfo = NULL;
 
        #define ARG_IS(param)                   (!strcmp(params->KernelArgs[i], param))
        #define NEXT_ARG_PRESENT() \
@@ -251,8 +256,8 @@ static void ParseParameters(Parameters * params, int * blk)
 
                        int tmp = strtol(params->KernelArgs[++i], NULL, 0);
 
-                       if (tmp <= 0) {
-                               printf("ERROR: blocking parameter must be > 0.\n");
+                       if (tmp < 0) {
+                               printf("ERROR: blocking parameter must be >= 0.\n");
                                exit(1);
                        }
 
@@ -263,8 +268,8 @@ static void ParseParameters(Parameters * params, int * blk)
 
                        int tmp = strtol(params->KernelArgs[++i], NULL, 0);
 
-                       if (tmp <= 0) {
-                               printf("ERROR: blocking parameter must be > 0.\n");
+                       if (tmp < 0) {
+                               printf("ERROR: blocking parameter must be >= 0.\n");
                                exit(1);
                        }
 
@@ -275,8 +280,8 @@ static void ParseParameters(Parameters * params, int * blk)
 
                        int tmp = strtol(params->KernelArgs[++i], NULL, 0);
 
-                       if (tmp <= 0) {
-                               printf("ERROR: blocking parameter must be > 0.\n");
+                       if (tmp < 0) {
+                               printf("ERROR: blocking parameter must be >= 0.\n");
                                exit(1);
                        }
 
@@ -287,13 +292,20 @@ static void ParseParameters(Parameters * params, int * blk)
 
                        int tmp = strtol(params->KernelArgs[++i], NULL, 0);
 
-                       if (tmp <= 0) {
-                               printf("ERROR: blocking parameter must be > 0.\n");
+                       if (tmp < 0) {
+                               printf("ERROR: blocking parameter must be >= 0.\n");
                                exit(1);
                        }
 
                        blk[2] = tmp;
                }
+#ifdef DATA_LAYOUT_SOA
+               else if (ARG_IS("-pad") || ARG_IS("--pad")) {
+                       NEXT_ARG_PRESENT();
+
+                       *padInfo = PadInfoFromStr(params->KernelArgs[++i]);
+               }
+#endif
                else if (ARG_IS("-h") || ARG_IS("-help") || ARG_IS("--help")) {
                        ParameterUsage();
                        exit(1);
@@ -502,6 +514,11 @@ void FNAME(D3Q19ListAaPvInit)(LatticeDesc * ld, KernelData ** kernelData, Parame
        kdlr->nConsecThreadIndices = 0;
 #endif
 
+       int blk[3] = { 0 };
+       PadInfo * padInfo = NULL;
+
+       ParseParameters(params, blk, &padInfo);
+
        // Ajust the dimensions according to padding, if used.
        kd->Dims[0] = kd->GlobalDims[0] = ld->Dims[0];
        kd->Dims[1] = kd->GlobalDims[1] = ld->Dims[1];
@@ -514,17 +531,21 @@ void FNAME(D3Q19ListAaPvInit)(LatticeDesc * ld, KernelData ** kernelData, Parame
        int lZ = lDims[2];
 
        int nTotalCells = lX * lY * lZ;
-       int nCells = ld->nFluid; // TODO: + padding
+       int nCells = ld->nFluid;
        int nFluid = ld->nFluid;
 
+#ifdef DATA_LAYOUT_SOA
+       {
+               nCells = PadCellsAndReport(nCells, sizeof(PdfT), &padInfo);
+               PadInfoFree(padInfo); padInfo = NULL;
+       }
+#endif
+
        kdl->nCells = nCells;
        kdl->nFluid = nFluid;
 
        PdfT * pdfs[2];
 
-       int blk[3] = { 0 };
-
-       ParseParameters(params, blk);
 
        if (blk[0] == 0) blk[0] = lX;
        if (blk[1] == 0) blk[1] = lY;
@@ -607,18 +628,18 @@ void FNAME(D3Q19ListAaPvInit)(LatticeDesc * ld, KernelData ** kernelData, Parame
 
        // Blocking is implemented via setup of the adjacency list. The kernel later will
        // walk through the lattice blocked automatically.
-       for (int bZ = 0; bZ < lZ; bZ += blk[2]) {
-       for (int bY = 0; bY < lY; bY += blk[1]) {
        for (int bX = 0; bX < lX; bX += blk[0]) {
+       for (int bY = 0; bY < lY; bY += blk[1]) {
+       for (int bZ = 0; bZ < lZ; bZ += blk[2]) {
 
                int eX = MIN(bX + blk[0], lX);
                int eY = MIN(bY + blk[1], lY);
                int eZ = MIN(bZ + blk[2], lZ);
 
 
-               for (int z = bZ; z < eZ; ++z) {
-               for (int y = bY; y < eY; ++y) {
                for (int x = bX; x < eX; ++x) {
+               for (int y = bY; y < eY; ++y) {
+               for (int z = bZ; z < eZ; ++z) {
 
                        latticeIndex = L_INDEX_4(lDims, x, y, z);
 
@@ -660,7 +681,15 @@ void FNAME(D3Q19ListAaPvInit)(LatticeDesc * ld, KernelData ** kernelData, Parame
 
        // Loop over all fluid nodes and compute the indices to the neighboring
        // PDFs for configured data layout (AoS/SoA).
-       // TODO: Parallelized loop to ensure correct NUMA placement.
+       #ifdef _OPENMP
+               #pragma omp parallel for
+       #endif
+       for (int index = 0; index < nFluid; ++index) {
+               for (int d = 0; d < N_D3Q19_IDX; ++d) {
+                       adjList[index * N_D3Q19_IDX + d] = -1;
+               }
+       }
+
        // #ifdef _OPENMP  --> add line continuation
        //      #pragma omp parallel for default(none)
        //              shared(nFluid, nCells, coords, D3Q19_INV, D3Q19_X, D3Q19_Y, D3Q19_Z,