projects
/
LbmBenchmarkKernelsPublic.git
/ blobdiff
commit
grep
author
committer
pickaxe
?
search:
re
summary
|
shortlog
|
log
|
commit
|
commitdiff
|
tree
raw
|
inline
| side by side
add citation information
[LbmBenchmarkKernelsPublic.git]
/
src
/
BenchKernelD3Q19List.c
diff --git
a/src/BenchKernelD3Q19List.c
b/src/BenchKernelD3Q19List.c
index e01853aab7cf668a25f5fce990120c8d6e6c73c0..0cb98f8b227d2b99ecf5acb6c463dd0aff21851a 100644
(file)
--- 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(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;
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;
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
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");
X_LIKWID_START("list-os");
+ // TODO: outer openmp parallel
#ifdef _OPENMP
#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, \
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, \
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
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))
#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;
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
#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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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;
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
#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;
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) {
#ifdef VTK_OUTPUT
if (cd->VtkOutput) {
This page took
0.141403 seconds
and
5
git commands to generate.