merge with kernels from MH's master thesis

[LbmBenchmarkKernelsPublic.git] / src / BenchKernelD3Q19ListPullSplitNt.c

diff --git a/src/BenchKernelD3Q19ListPullSplitNt.c b/src/BenchKernelD3Q19ListPullSplitNt.c
index 0132dc9058f26d7f18f153e8a0c11f37d38efa74..05d1d0a5e59a223c3b41669c67055ab32ded894a 100644 (file)
--- a/src/BenchKernelD3Q19ListPullSplitNt.c
+++ b/src/BenchKernelD3Q19ListPullSplitNt.c
@@ -55,8 +55,8 @@ void FNAME(KernelPullSplitNt1S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        Assert(kernelData != NULL);
        Assert(cd != NULL);
 
-       Assert(cd->Omega > 0.0);
-       Assert(cd->Omega < 2.0);
+       Assert(cd->Omega > F(0.0));
+       Assert(cd->Omega < F(2.0));
 
        KernelData        * kd   = (KernelData *)kernelData;
        KernelDataList    * kdl  = KDL(kernelData);
@@ -65,16 +65,16 @@ void FNAME(KernelPullSplitNt1S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        PdfT omega = cd->Omega;
        const PdfT omegaEven = omega;
 
-       PdfT magicParam = 1.0 / 12.0;
-       const PdfT omegaOdd = 1.0 / (0.5 + magicParam / (1.0 / omega - 0.5));
+       PdfT magicParam = F(1.0) / F(12.0);
+       const PdfT omegaOdd = F(1.0) / (F(0.5) + magicParam / (F(1.0) / omega - F(0.5)));
 
 
-       const PdfT w_0 = 1.0 /  3.0;
-       const PdfT w_1 = 1.0 / 18.0;
-       const PdfT w_2 = 1.0 / 36.0;
+       const PdfT w_0 = F(1.0) / F( 3.0);
+       const PdfT w_1 = F(1.0) / F(18.0);
+       const PdfT w_2 = F(1.0) / F(36.0);
 
-       const PdfT w_1_x3 = w_1 * 3.0;  const PdfT w_1_nine_half = w_1 * 9.0 / 2.0;
-       const PdfT w_2_x3 = w_2 * 3.0;  const PdfT w_2_nine_half = w_2 * 9.0 / 2.0;
+       const PdfT w_1_x3 = w_1 * F(3.0);       const PdfT w_1_nine_half = w_1 * F(9.0) / F(2.0);
+       const PdfT w_2_x3 = w_2 * F(3.0);       const PdfT w_2_nine_half = w_2 * F(9.0) / F(2.0);
 
        const VPDFT vw_1_x3 = VSET(w_1_x3);
        const VPDFT vw_2_x3 = VSET(w_2_x3);
@@ -85,7 +85,7 @@ void FNAME(KernelPullSplitNt1S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        const VPDFT vomegaEven = VSET(omegaEven);
        const VPDFT vomegaOdd  = VSET(omegaOdd);
 
-       const VPDFT voneHalf = VSET(0.5);
+       const VPDFT voneHalf = VSET(F(0.5));
 
        // uint32_t nConsecNodes = kdlr->nConsecNodes;
        // uint32_t * consecNodes = kdlr->ConsecNodes;
@@ -119,8 +119,10 @@ void FNAME(KernelPullSplitNt1S)(LatticeDesc * ld, KernelData * kernelData, CaseD
                KernelStatistics(kd, ld, cd, 0);
        #endif
 
+       X_KERNEL_START(kernelData);
+
+       X_LIKWID_START("list-pull-split-nt-1s");
 
-                       X_LIKWID_START("list-pull-split-nt-1s");
        #ifdef _OPENMP
                #pragma omp parallel default(none) \
                        shared(nFluid, nCells, kd, kdl, adjList, src, dst, \
@@ -244,6 +246,8 @@ void FNAME(KernelPullSplitNt1S)(LatticeDesc * ld, KernelData * kernelData, CaseD
 
        X_LIKWID_STOP("list-pull-split-nt-1s");
 
+       X_KERNEL_END(kernelData);
+
 #ifdef VTK_OUTPUT
        if (cd->VtkOutput) {
                kd->PdfsActive = src;
@@ -266,8 +270,8 @@ void FNAME(KernelPullSplitNt2S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        Assert(kernelData != NULL);
        Assert(cd != NULL);
 
-       Assert(cd->Omega > 0.0);
-       Assert(cd->Omega < 2.0);
+       Assert(cd->Omega > F(0.0));
+       Assert(cd->Omega < F(2.0));
 
        KernelData        * kd   = (KernelData *)kernelData;
        KernelDataList    * kdl  = KDL(kernelData);
@@ -276,16 +280,15 @@ void FNAME(KernelPullSplitNt2S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        PdfT omega = cd->Omega;
        const PdfT omegaEven = omega;
 
-       PdfT magicParam = 1.0 / 12.0;
-       const PdfT omegaOdd = 1.0 / (0.5 + magicParam / (1.0 / omega - 0.5));
-
+       PdfT magicParam = F(1.0) / F(12.0);
+       const PdfT omegaOdd = F(1.0) / (F(0.5) + magicParam / (F(1.0) / omega - F(0.5)));
 
-       const PdfT w_0 = 1.0 /  3.0;
-       const PdfT w_1 = 1.0 / 18.0;
-       const PdfT w_2 = 1.0 / 36.0;
+       const PdfT w_0 = F(1.0) / F( 3.0);
+       const PdfT w_1 = F(1.0) / F(18.0);
+       const PdfT w_2 = F(1.0) / F(36.0);
 
-       const PdfT w_1_x3 = w_1 * 3.0;  const PdfT w_1_nine_half = w_1 * 9.0 / 2.0;
-       const PdfT w_2_x3 = w_2 * 3.0;  const PdfT w_2_nine_half = w_2 * 9.0 / 2.0;
+       const PdfT w_1_x3 = w_1 * F(3.0);       const PdfT w_1_nine_half = w_1 * F(9.0) / F(2.0);
+       const PdfT w_2_x3 = w_2 * F(3.0);       const PdfT w_2_nine_half = w_2 * F(9.0) / F(2.0);
 
        const VPDFT vw_1_x3 = VSET(w_1_x3);
        const VPDFT vw_2_x3 = VSET(w_2_x3);
@@ -296,7 +299,7 @@ void FNAME(KernelPullSplitNt2S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        const VPDFT vomegaEven = VSET(omegaEven);
        const VPDFT vomegaOdd  = VSET(omegaOdd);
 
-       const VPDFT voneHalf = VSET(0.5);
+       const VPDFT voneHalf = VSET(F(0.5));
 
        // uint32_t nConsecNodes = kdlr->nConsecNodes;
        // uint32_t * consecNodes = kdlr->ConsecNodes;
@@ -331,7 +334,9 @@ void FNAME(KernelPullSplitNt2S)(LatticeDesc * ld, KernelData * kernelData, CaseD
        #endif
 
 
-                       X_LIKWID_START("list-pull-split-nt-2s");
+       X_KERNEL_START(kernelData);
+
+       X_LIKWID_START("list-pull-split-nt-2s");
 
 
        #ifdef _OPENMP
@@ -452,7 +457,9 @@ void FNAME(KernelPullSplitNt2S)(LatticeDesc * ld, KernelData * kernelData, CaseD
                MemFree((void **)&tmpArray);
        }
 
-                       X_LIKWID_STOP("list-pull-split-nt-2s");
+       X_LIKWID_STOP("list-pull-split-nt-2s");
+
+       X_KERNEL_END(kernelData);
 
 #ifdef VTK_OUTPUT
        if (cd->VtkOutput) {