X-Git-Url: http://git.rrze.uni-erlangen.de/gitweb/?a=blobdiff_plain;f=src%2FBenchKernelD3Q19List.c;h=0cb98f8b227d2b99ecf5acb6c463dd0aff21851a;hb=HEAD;hp=e01853aab7cf668a25f5fce990120c8d6e6c73c0;hpb=e3f82424829ebb623343ce0092238f83b4a1b8c2;p=LbmBenchmarkKernelsPublic.git diff --git a/src/BenchKernelD3Q19List.c b/src/BenchKernelD3Q19List.c index e01853a..0cb98f8 100644 --- a/src/BenchKernelD3Q19List.c +++ b/src/BenchKernelD3Q19List.c @@ -40,28 +40,27 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData 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 = (KernelDataList *)kernelData; PdfT omega = cd->Omega; PdfT omegaEven = omega; -// PdfT omegaOdd = 8.0*((2.0-omegaEven)/(8.0-omegaEven)); //"standard" trt odd relaxation parameter - PdfT magicParam = 1.0/12.0; // 1/4: best stability; 1/12: removes third-order advection error (best advection); 1/6: removes fourth-order diffusion error (best diffusion); 3/16: exact location of bounce back for poiseuille flow - PdfT omegaOdd = 1.0/( 0.5 + magicParam/(1.0/omega - 0.5) ); + PdfT magicParam = F(1.0) / F(12.0); + PdfT omegaOdd = F(1.0) /(F(0.5) + magicParam / (F(1.0) / omega - F(0.5))); - PdfT evenPart = 0.0; - PdfT oddPart = 0.0; - PdfT dir_indep_trm = 0.0; + PdfT evenPart = F(0.0); + PdfT oddPart = F(0.0); + PdfT dir_indep_trm = F(0.0); - PdfT w_0 = 1.0 / 3.0; - PdfT w_1 = 1.0 / 18.0; - PdfT w_2 = 1.0 / 36.0; + PdfT w_0 = F(1.0) / F( 3.0); + PdfT w_1 = F(1.0) / F(18.0); + PdfT w_2 = F(1.0) / F(36.0); - PdfT w_1_x3 = w_1 * 3.0; PdfT w_1_nine_half = w_1 * 9.0/2.0; PdfT w_1_indep = 0.0; - PdfT w_2_x3 = w_2 * 3.0; PdfT w_2_nine_half = w_2 * 9.0/2.0; PdfT w_2_indep = 0.0; + PdfT w_1_x3 = w_1 * F(3.0); PdfT w_1_nine_half = w_1 * F(9.0) / F(2.0); PdfT w_1_indep = F(0.0); + PdfT w_2_x3 = w_2 * F(3.0); PdfT w_2_nine_half = w_2 * F(9.0) / F(2.0); PdfT w_2_indep = F(0.0); PdfT ux, uy, uz, ui; PdfT dens; @@ -94,16 +93,15 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData KernelStatistics(kd, ld, cd, 0); #endif - // TODO: outer openmp parallel - for(int iter = 0; iter < maxIterations; ++iter) { - + X_KERNEL_START(kernelData); X_LIKWID_START("list-os"); + // TODO: outer openmp parallel #ifdef _OPENMP - #pragma omp parallel for default(none) \ + #pragma omp parallel default(none) \ shared(nFluid, nCells, kd, kdl, adjList, src, dst, w_0, w_1, w_2, omegaEven, omegaOdd, \ - w_1_x3, w_2_x3, w_1_nine_half, w_2_nine_half, cd) \ + w_1_x3, w_2_x3, w_1_nine_half, w_2_nine_half, cd, ld, tmp, maxIterations ) \ private(ux, uy, uz, ui, dens, dir_indep_trm, adjListIndex, \ pdf_C, \ pdf_N, pdf_E, pdf_S, pdf_W, \ @@ -112,7 +110,16 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData pdf_B, pdf_BN, pdf_BE, pdf_BS, pdf_BW, \ evenPart, oddPart, w_1_indep, w_2_indep) #endif - for (int index = 0; index < nFluid; ++index) { +{ + for(int iter = 0; iter < maxIterations; ++iter) { + + + + #ifdef INTEL_OPT_DIRECTIVES + #endif + #pragma omp for + #pragma novector + for (int index = 0; index < nFluid; ++index) { // LOOP list-os #define I(index, dir) P_INDEX_3((nCells), (index), (dir)) @@ -171,7 +178,7 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData pdf_T + pdf_TN + pdf_TE + pdf_TS + pdf_TW + pdf_B + pdf_BN + pdf_BE + pdf_BS + pdf_BW; - dir_indep_trm = dens - (ux * ux + uy * uy + uz * uz)*3.0/2.0; + dir_indep_trm = dens - (ux * ux + uy * uy + uz * uz) * F(3.0) / F(2.0); #ifdef PROP_MODEL_PUSH @@ -184,20 +191,20 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData w_1_indep = w_1*dir_indep_trm; ui = uy; - evenPart = omegaEven*( 0.5*(pdf_N + pdf_S) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_N - pdf_S) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_N + pdf_S) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_N - pdf_S) - ui*w_1_x3 ); dst[adjList[adjListIndex + D3Q19_N]] = pdf_N - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_S]] = pdf_S - evenPart + oddPart; ui = ux; - evenPart = omegaEven*( 0.5*(pdf_E + pdf_W) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_E - pdf_W) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_E + pdf_W) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_E - pdf_W) - ui*w_1_x3 ); dst[adjList[adjListIndex + D3Q19_E]] = pdf_E - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_W]] = pdf_W - evenPart + oddPart; ui = uz; - evenPart = omegaEven*( 0.5*(pdf_T + pdf_B) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_T - pdf_B) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_T + pdf_B) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_T - pdf_B) - ui*w_1_x3 ); dst[adjList[adjListIndex + D3Q19_T]] = pdf_T - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_B]] = pdf_B - evenPart + oddPart; @@ -205,38 +212,38 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData w_2_indep = w_2*dir_indep_trm; ui = -ux + uy; - evenPart = omegaEven*( 0.5*(pdf_NW + pdf_SE) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_NW - pdf_SE) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_NW + pdf_SE) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_NW - pdf_SE) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_NW]] = pdf_NW - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_SE]] = pdf_SE - evenPart + oddPart; ui = ux + uy; - evenPart = omegaEven*( 0.5*(pdf_NE + pdf_SW) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_NE - pdf_SW) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_NE + pdf_SW) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_NE - pdf_SW) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_NE]] = pdf_NE - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_SW]] = pdf_SW - evenPart + oddPart; ui = -ux + uz; - evenPart = omegaEven*( 0.5*(pdf_TW + pdf_BE) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TW - pdf_BE) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TW + pdf_BE) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TW - pdf_BE) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_TW]] = pdf_TW - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_BE]] = pdf_BE - evenPart + oddPart; ui = ux + uz; - evenPart = omegaEven*( 0.5*(pdf_TE + pdf_BW) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TE - pdf_BW) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TE + pdf_BW) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TE - pdf_BW) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_TE]] = pdf_TE - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_BW]] = pdf_BW - evenPart + oddPart; ui = -uy + uz; - evenPart = omegaEven*( 0.5*(pdf_TS + pdf_BN) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TS - pdf_BN) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TS + pdf_BN) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TS - pdf_BN) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_TS]] = pdf_TS - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_BN]] = pdf_BN - evenPart + oddPart; ui = uy + uz; - evenPart = omegaEven*( 0.5*(pdf_TN + pdf_BS) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TN - pdf_BS) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TN + pdf_BS) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TN - pdf_BS) - ui*w_2_x3 ); dst[adjList[adjListIndex + D3Q19_TN]] = pdf_TN - evenPart - oddPart; dst[adjList[adjListIndex + D3Q19_BS]] = pdf_BS - evenPart + oddPart; @@ -249,20 +256,20 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData w_1_indep = w_1*dir_indep_trm; ui = uy; - evenPart = omegaEven*( 0.5*(pdf_N + pdf_S) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_N - pdf_S) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_N + pdf_S) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_N - pdf_S) - ui*w_1_x3 ); dst[I(index, D3Q19_N )] = pdf_N - evenPart - oddPart; dst[I(index, D3Q19_S )] = pdf_S - evenPart + oddPart; ui = ux; - evenPart = omegaEven*( 0.5*(pdf_E + pdf_W) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_E - pdf_W) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_E + pdf_W) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_E - pdf_W) - ui*w_1_x3 ); dst[I(index, D3Q19_E )] = pdf_E - evenPart - oddPart; dst[I(index, D3Q19_W )] = pdf_W - evenPart + oddPart; ui = uz; - evenPart = omegaEven*( 0.5*(pdf_T + pdf_B) - ui*ui*w_1_nine_half - w_1_indep ); - oddPart = omegaOdd*(0.5*(pdf_T - pdf_B) - ui*w_1_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_T + pdf_B) - ui*ui*w_1_nine_half - w_1_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_T - pdf_B) - ui*w_1_x3 ); dst[I(index, D3Q19_T )] = pdf_T - evenPart - oddPart; dst[I(index, D3Q19_B )] = pdf_B - evenPart + oddPart; @@ -270,38 +277,38 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData w_2_indep = w_2*dir_indep_trm; ui = -ux + uy; - evenPart = omegaEven*( 0.5*(pdf_NW + pdf_SE) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_NW - pdf_SE) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_NW + pdf_SE) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_NW - pdf_SE) - ui*w_2_x3 ); dst[I(index, D3Q19_NW)] = pdf_NW - evenPart - oddPart; dst[I(index, D3Q19_SE)] = pdf_SE - evenPart + oddPart; ui = ux + uy; - evenPart = omegaEven*( 0.5*(pdf_NE + pdf_SW) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_NE - pdf_SW) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_NE + pdf_SW) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_NE - pdf_SW) - ui*w_2_x3 ); dst[I(index, D3Q19_NE)] = pdf_NE - evenPart - oddPart; dst[I(index, D3Q19_SW)] = pdf_SW - evenPart + oddPart; ui = -ux + uz; - evenPart = omegaEven*( 0.5*(pdf_TW + pdf_BE) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TW - pdf_BE) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TW + pdf_BE) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TW - pdf_BE) - ui*w_2_x3 ); dst[I(index, D3Q19_TW)] = pdf_TW - evenPart - oddPart; dst[I(index, D3Q19_BE)] = pdf_BE - evenPart + oddPart; ui = ux + uz; - evenPart = omegaEven*( 0.5*(pdf_TE + pdf_BW) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TE - pdf_BW) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TE + pdf_BW) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TE - pdf_BW) - ui*w_2_x3 ); dst[I(index, D3Q19_TE)] = pdf_TE - evenPart - oddPart; dst[I(index, D3Q19_BW)] = pdf_BW - evenPart + oddPart; ui = -uy + uz; - evenPart = omegaEven*( 0.5*(pdf_TS + pdf_BN) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TS - pdf_BN) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TS + pdf_BN) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TS - pdf_BN) - ui*w_2_x3 ); dst[I(index, D3Q19_TS)] = pdf_TS - evenPart - oddPart; dst[I(index, D3Q19_BN)] = pdf_BN - evenPart + oddPart; ui = uy + uz; - evenPart = omegaEven*( 0.5*(pdf_TN + pdf_BS) - ui*ui*w_2_nine_half - w_2_indep ); - oddPart = omegaOdd*(0.5*(pdf_TN - pdf_BS) - ui*w_2_x3 ); + evenPart = omegaEven*( F(0.5)*(pdf_TN + pdf_BS) - ui*ui*w_2_nine_half - w_2_indep ); + oddPart = omegaOdd*(F(0.5)*(pdf_TN - pdf_BS) - ui*w_2_x3 ); dst[I(index, D3Q19_TN)] = pdf_TN - evenPart - oddPart; dst[I(index, D3Q19_BS)] = pdf_BS - evenPart + oddPart; @@ -309,31 +316,36 @@ void FNAME(D3Q19ListKernel)(LatticeDesc * ld, KernelData * kernelData, CaseData #undef I } // loop over fluid nodes - X_LIKWID_STOP("list-os"); + #pragma omp single + { + #ifdef VERIFICATION + kd->PdfsActive = dst; + KernelAddBodyForce(kd, ld, cd); + #endif - #ifdef VERIFICATION - kd->PdfsActive = dst; - KernelAddBodyForce(kd, ld, cd); - #endif + #ifdef VTK_OUTPUT + if (cd->VtkOutput && (iter % cd->VtkModulus) == 0) { + kd->PdfsActive = dst; + VtkWrite(ld, kd, cd, iter); + } + #endif - #ifdef VTK_OUTPUT - if (cd->VtkOutput && (iter % cd->VtkModulus) == 0) { + #ifdef STATISTICS kd->PdfsActive = dst; - VtkWrite(ld, kd, cd, iter); - } - #endif + KernelStatistics(kd, ld, cd, iter); + #endif - #ifdef STATISTICS - kd->PdfsActive = dst; - KernelStatistics(kd, ld, cd, iter); - #endif + // swap grids + tmp = src; + src = dst; + dst = tmp; + } + } + } // for (int iter = 0; ... - // swap grids - tmp = src; - src = dst; - dst = tmp; + X_LIKWID_STOP("list-os"); - } // for (int iter = 0; ... + X_KERNEL_END(kernelData); #ifdef VTK_OUTPUT if (cd->VtkOutput) {