From b50e017cd8a9f7b67b8462e2788eebf073f14f98 Mon Sep 17 00:00:00 2001
From: "LEGOLAS\\lenz" <lenz@irmb.tu-bs.de>
Date: Wed, 24 Apr 2019 09:44:56 +0200
Subject: [PATCH] implements overloaded operators for FlowStateData und
simplifies the interface interpolations
---
src/GksGpu/CellUpdate/CellUpdate.cu | 2 +-
src/GksGpu/FlowStateData/FlowStateData.cuh | 124 ++++++
src/GksGpu/Interface/CoarseToFineKernel.cu | 444 +++++++++++++--------
src/GksGpu/Interface/FineToCoarseKernel.cu | 66 ++-
targets/apps/GKS/PoolFire/PoolFire.cpp | 2 +-
5 files changed, 453 insertions(+), 185 deletions(-)
diff --git a/src/GksGpu/CellUpdate/CellUpdate.cu b/src/GksGpu/CellUpdate/CellUpdate.cu
index 8efa45c2b..9ffa93936 100644
--- a/src/GksGpu/CellUpdate/CellUpdate.cu
+++ b/src/GksGpu/CellUpdate/CellUpdate.cu
@@ -264,7 +264,7 @@ __host__ __device__ inline void cellUpdateFunction(DataBaseStruct dataBase, Para
dataBase.data[RHO_S_1(cellIndex, dataBase.numberOfCells)] = Z1 * updatedConserved.rho;
dataBase.data[RHO_S_2(cellIndex, dataBase.numberOfCells)] = Z2 * updatedConserved.rho;
- dataBase.data[RHO_E(cellIndex, dataBase.numberOfCells)] = updatedConserved.rhoE + releasedHeat;
+ dataBase.data[RHO_E(cellIndex, dataBase.numberOfCells)] = updatedConserved.rhoE + releasedHeat;
}
}
}
diff --git a/src/GksGpu/FlowStateData/FlowStateData.cuh b/src/GksGpu/FlowStateData/FlowStateData.cuh
index fcfdeb541..907a3ca31 100644
--- a/src/GksGpu/FlowStateData/FlowStateData.cuh
+++ b/src/GksGpu/FlowStateData/FlowStateData.cuh
@@ -125,5 +125,129 @@ struct ConservedVariables
//////////////////////////////////////////////////////////////////////////
};
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
+__host__ __device__ inline PrimitiveVariables operator+ ( const PrimitiveVariables& left, const PrimitiveVariables& right )
+{
+ PrimitiveVariables result;
+
+ result.rho = left.rho + right.rho ;
+ result.U = left.U + right.U ;
+ result.V = left.V + right.V ;
+ result.W = left.W + right.W ;
+ result.lambda = left.lambda + right.lambda;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.S_1 = left.S_1 + right.S_1 ;
+ result.S_2 = left.S_2 + right.S_2 ;
+#endif
+
+ return result;
+}
+
+__host__ __device__ inline ConservedVariables operator+ ( const ConservedVariables& left, const ConservedVariables& right )
+{
+ ConservedVariables result;
+
+ result.rho = left.rho + right.rho ;
+ result.rhoU = left.rhoU + right.rhoU ;
+ result.rhoV = left.rhoV + right.rhoV ;
+ result.rhoW = left.rhoW + right.rhoW ;
+ result.rhoE = left.rhoE + right.rhoE ;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.rhoS_1 = left.rhoS_1 + right.rhoS_1;
+ result.rhoS_2 = left.rhoS_2 + right.rhoS_2;
+#endif
+
+ return result;
+}
+
+//////////////////////////////////////////////////////////////////////////
+
+__host__ __device__ inline PrimitiveVariables operator- ( const PrimitiveVariables& left, const PrimitiveVariables& right )
+{
+ PrimitiveVariables result;
+
+ result.rho = left.rho - right.rho ;
+ result.U = left.U - right.U ;
+ result.V = left.V - right.V ;
+ result.W = left.W - right.W ;
+ result.lambda = left.lambda - right.lambda;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.S_1 = left.S_1 - right.S_1 ;
+ result.S_2 = left.S_2 - right.S_2 ;
+#endif
+
+ return result;
+}
+
+__host__ __device__ inline ConservedVariables operator- ( const ConservedVariables& left, const ConservedVariables& right )
+{
+ ConservedVariables result;
+
+ result.rho = left.rho - right.rho ;
+ result.rhoU = left.rhoU - right.rhoU ;
+ result.rhoV = left.rhoV - right.rhoV ;
+ result.rhoW = left.rhoW - right.rhoW ;
+ result.rhoE = left.rhoE - right.rhoE ;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.rhoS_1 = left.rhoS_1 - right.rhoS_1;
+ result.rhoS_2 = left.rhoS_2 - right.rhoS_2;
+#endif
+
+ return result;
+}
+
+//////////////////////////////////////////////////////////////////////////
+
+__host__ __device__ inline PrimitiveVariables operator* ( const real left, const PrimitiveVariables& right )
+{
+ PrimitiveVariables result;
+
+ result.rho = left * right.rho ;
+ result.U = left * right.U ;
+ result.V = left * right.V ;
+ result.W = left * right.W ;
+ result.lambda = left * right.lambda;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.S_1 = left * right.S_1 ;
+ result.S_2 = left * right.S_2 ;
+#endif
+
+ return result;
+}
+
+__host__ __device__ inline ConservedVariables operator* ( const real left, const ConservedVariables& right )
+{
+ ConservedVariables result;
+
+ result.rho = left * right.rho ;
+ result.rhoU = left * right.rhoU ;
+ result.rhoV = left * right.rhoV ;
+ result.rhoW = left * right.rhoW ;
+ result.rhoE = left * right.rhoE ;
+
+#ifdef USE_PASSIVE_SCALAR
+ result.rhoS_1 = left * right.rhoS_1;
+ result.rhoS_2 = left * right.rhoS_2;
+#endif
+
+ return result;
+}
+
#endif
diff --git a/src/GksGpu/Interface/CoarseToFineKernel.cu b/src/GksGpu/Interface/CoarseToFineKernel.cu
index c43f0370a..5f7d2d36b 100644
--- a/src/GksGpu/Interface/CoarseToFineKernel.cu
+++ b/src/GksGpu/Interface/CoarseToFineKernel.cu
@@ -9,6 +9,9 @@
#include "DataBase/DataBaseStruct.h"
#include "FlowStateData/FlowStateData.cuh"
+#include "FlowStateData/FlowStateDataConversion.cuh"
+
+#include "FlowStateData/AccessDeviceData.cuh"
#include "Definitions/PassiveScalar.h"
#include "Definitions/MemoryAccessPattern.h"
@@ -19,7 +22,8 @@
__global__ void coarseToFineKernel ( DataBaseStruct dataBase, uint startIndex, uint numberOfEntities );
-__host__ __device__ inline void coarseToFineFunction( DataBaseStruct dataBase, uint startIndex, uint index );
+__host__ __device__ inline void coarseToFineFunction ( DataBaseStruct dataBase, uint startIndex, uint index );
+__host__ __device__ inline void coarseToFineFunctionPrimitiveInterpolation( DataBaseStruct dataBase, uint startIndex, uint index );
//////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////
@@ -51,8 +55,213 @@ __global__ void coarseToFineKernel( DataBaseStruct dataBase, uint startIndex, ui
if( index >= numberOfEntities ) return;
coarseToFineFunction( dataBase, startIndex, index );
+ //coarseToFineFunctionPrimitiveInterpolation( dataBase, startIndex, index );
}
+//__host__ __device__ inline void coarseToFineFunction( DataBaseStruct dataBase, uint startIndex, uint index )
+//{
+// index += startIndex;
+//
+// uint cellIndex = dataBase.coarseToFine[ COARSE_TO_FINE( index, 0, dataBase.numberOfFineGhostCells ) ];
+//
+// uint cellToCell [6];
+//
+// cellToCell[0] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 0, dataBase.numberOfCells ) ];
+// cellToCell[1] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 1, dataBase.numberOfCells ) ];
+// cellToCell[2] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 2, dataBase.numberOfCells ) ];
+// cellToCell[3] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 3, dataBase.numberOfCells ) ];
+// cellToCell[4] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 4, dataBase.numberOfCells ) ];
+// cellToCell[5] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 5, dataBase.numberOfCells ) ];
+//
+// ConservedVariables childCons [8];
+//
+// {
+// real data [7];
+//
+// data[0] = dataBase.data[ RHO__(cellToCell[0], dataBase.numberOfCells) ];
+// data[1] = dataBase.data[ RHO__(cellToCell[1], dataBase.numberOfCells) ];
+// data[2] = dataBase.data[ RHO__(cellToCell[2], dataBase.numberOfCells) ];
+// data[3] = dataBase.data[ RHO__(cellToCell[3], dataBase.numberOfCells) ];
+// data[4] = dataBase.data[ RHO__(cellToCell[4], dataBase.numberOfCells) ];
+// data[5] = dataBase.data[ RHO__(cellToCell[5], dataBase.numberOfCells) ];
+// data[6] = dataBase.data[ RHO__(cellIndex , dataBase.numberOfCells) ];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rho = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
+// childCons[1].rho = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
+// childCons[2].rho = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
+// childCons[3].rho = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
+// childCons[4].rho = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
+// childCons[5].rho = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
+// childCons[6].rho = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
+// childCons[7].rho = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+// }
+//
+// {
+// real data [7];
+//
+// data[0] = dataBase.data[ RHO_U(cellToCell[0], dataBase.numberOfCells) ];
+// data[1] = dataBase.data[ RHO_U(cellToCell[1], dataBase.numberOfCells) ];
+// data[2] = dataBase.data[ RHO_U(cellToCell[2], dataBase.numberOfCells) ];
+// data[3] = dataBase.data[ RHO_U(cellToCell[3], dataBase.numberOfCells) ];
+// data[4] = dataBase.data[ RHO_U(cellToCell[4], dataBase.numberOfCells) ];
+// data[5] = dataBase.data[ RHO_U(cellToCell[5], dataBase.numberOfCells) ];
+// data[6] = dataBase.data[ RHO_U(cellIndex , dataBase.numberOfCells) ];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoU = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
+// childCons[1].rhoU = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
+// childCons[2].rhoU = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
+// childCons[3].rhoU = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
+// childCons[4].rhoU = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
+// childCons[5].rhoU = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
+// childCons[6].rhoU = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
+// childCons[7].rhoU = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+// }
+//
+// {
+// real data [7];
+//
+// data[0] = dataBase.data[ RHO_V(cellToCell[0], dataBase.numberOfCells) ];
+// data[1] = dataBase.data[ RHO_V(cellToCell[1], dataBase.numberOfCells) ];
+// data[2] = dataBase.data[ RHO_V(cellToCell[2], dataBase.numberOfCells) ];
+// data[3] = dataBase.data[ RHO_V(cellToCell[3], dataBase.numberOfCells) ];
+// data[4] = dataBase.data[ RHO_V(cellToCell[4], dataBase.numberOfCells) ];
+// data[5] = dataBase.data[ RHO_V(cellToCell[5], dataBase.numberOfCells) ];
+// data[6] = dataBase.data[ RHO_V(cellIndex , dataBase.numberOfCells) ];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoV = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
+// childCons[1].rhoV = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
+// childCons[2].rhoV = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
+// childCons[3].rhoV = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
+// childCons[4].rhoV = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
+// childCons[5].rhoV = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
+// childCons[6].rhoV = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
+// childCons[7].rhoV = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+// }
+//
+// {
+// real data [7];
+//
+// data[0] = dataBase.data[ RHO_W(cellToCell[0], dataBase.numberOfCells) ];
+// data[1] = dataBase.data[ RHO_W(cellToCell[1], dataBase.numberOfCells) ];
+// data[2] = dataBase.data[ RHO_W(cellToCell[2], dataBase.numberOfCells) ];
+// data[3] = dataBase.data[ RHO_W(cellToCell[3], dataBase.numberOfCells) ];
+// data[4] = dataBase.data[ RHO_W(cellToCell[4], dataBase.numberOfCells) ];
+// data[5] = dataBase.data[ RHO_W(cellToCell[5], dataBase.numberOfCells) ];
+// data[6] = dataBase.data[ RHO_W(cellIndex , dataBase.numberOfCells) ];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoW = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
+// childCons[1].rhoW = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
+// childCons[2].rhoW = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
+// childCons[3].rhoW = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
+// childCons[4].rhoW = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
+// childCons[5].rhoW = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
+// childCons[6].rhoW = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
+// childCons[7].rhoW = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+// }
+//
+// {
+// real data [7];
+//
+// data[0] = dataBase.data[ RHO_E(cellToCell[0], dataBase.numberOfCells) ];
+// data[1] = dataBase.data[ RHO_E(cellToCell[1], dataBase.numberOfCells) ];
+// data[2] = dataBase.data[ RHO_E(cellToCell[2], dataBase.numberOfCells) ];
+// data[3] = dataBase.data[ RHO_E(cellToCell[3], dataBase.numberOfCells) ];
+// data[4] = dataBase.data[ RHO_E(cellToCell[4], dataBase.numberOfCells) ];
+// data[5] = dataBase.data[ RHO_E(cellToCell[5], dataBase.numberOfCells) ];
+// data[6] = dataBase.data[ RHO_E(cellIndex , dataBase.numberOfCells) ];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoE = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
+// childCons[1].rhoE = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
+// childCons[2].rhoE = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
+// childCons[3].rhoE = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
+// childCons[4].rhoE = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
+// childCons[5].rhoE = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
+// childCons[6].rhoE = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
+// childCons[7].rhoE = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+// }
+//
+// #ifdef USE_PASSIVE_SCALAR
+// {
+// {
+// real data[7];
+//
+// data[0] = dataBase.data[RHO_S_1(cellToCell[0], dataBase.numberOfCells)];
+// data[1] = dataBase.data[RHO_S_1(cellToCell[1], dataBase.numberOfCells)];
+// data[2] = dataBase.data[RHO_S_1(cellToCell[2], dataBase.numberOfCells)];
+// data[3] = dataBase.data[RHO_S_1(cellToCell[3], dataBase.numberOfCells)];
+// data[4] = dataBase.data[RHO_S_1(cellToCell[4], dataBase.numberOfCells)];
+// data[5] = dataBase.data[RHO_S_1(cellToCell[5], dataBase.numberOfCells)];
+// data[6] = dataBase.data[RHO_S_1(cellIndex, dataBase.numberOfCells)];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoS_1 = data[6] + c1o8 * (+data[0] + data[2] + data[4] - data[1] - data[3] - data[5]); // PX PY PZ
+// childCons[1].rhoS_1 = data[6] + c1o8 * (+data[0] + data[2] - data[4] - data[1] - data[3] + data[5]); // PX PY MZ
+// childCons[2].rhoS_1 = data[6] + c1o8 * (+data[0] - data[2] + data[4] - data[1] + data[3] - data[5]); // PX MY PZ
+// childCons[3].rhoS_1 = data[6] + c1o8 * (+data[0] - data[2] - data[4] - data[1] + data[3] + data[5]); // PX MY MZ
+// childCons[4].rhoS_1 = data[6] + c1o8 * (-data[0] + data[2] + data[4] + data[1] - data[3] - data[5]); // MX PY PZ
+// childCons[5].rhoS_1 = data[6] + c1o8 * (-data[0] + data[2] - data[4] + data[1] - data[3] + data[5]); // MX PY MZ
+// childCons[6].rhoS_1 = data[6] + c1o8 * (-data[0] - data[2] + data[4] + data[1] + data[3] - data[5]); // MX MY PZ
+// childCons[7].rhoS_1 = data[6] + c1o8 * (-data[0] - data[2] - data[4] + data[1] + data[3] + data[5]); // MX MY MZ
+// }
+//
+// {
+// real data[7];
+//
+// data[0] = dataBase.data[RHO_S_2(cellToCell[0], dataBase.numberOfCells)];
+// data[1] = dataBase.data[RHO_S_2(cellToCell[1], dataBase.numberOfCells)];
+// data[2] = dataBase.data[RHO_S_2(cellToCell[2], dataBase.numberOfCells)];
+// data[3] = dataBase.data[RHO_S_2(cellToCell[3], dataBase.numberOfCells)];
+// data[4] = dataBase.data[RHO_S_2(cellToCell[4], dataBase.numberOfCells)];
+// data[5] = dataBase.data[RHO_S_2(cellToCell[5], dataBase.numberOfCells)];
+// data[6] = dataBase.data[RHO_S_2(cellIndex, dataBase.numberOfCells)];
+//
+// // PX PY PZ MX MY MZ
+// childCons[0].rhoS_2 = data[6] + c1o8 * (+data[0] + data[2] + data[4] - data[1] - data[3] - data[5]); // PX PY PZ
+// childCons[1].rhoS_2 = data[6] + c1o8 * (+data[0] + data[2] - data[4] - data[1] - data[3] + data[5]); // PX PY MZ
+// childCons[2].rhoS_2 = data[6] + c1o8 * (+data[0] - data[2] + data[4] - data[1] + data[3] - data[5]); // PX MY PZ
+// childCons[3].rhoS_2 = data[6] + c1o8 * (+data[0] - data[2] - data[4] - data[1] + data[3] + data[5]); // PX MY MZ
+// childCons[4].rhoS_2 = data[6] + c1o8 * (-data[0] + data[2] + data[4] + data[1] - data[3] - data[5]); // MX PY PZ
+// childCons[5].rhoS_2 = data[6] + c1o8 * (-data[0] + data[2] - data[4] + data[1] - data[3] + data[5]); // MX PY MZ
+// childCons[6].rhoS_2 = data[6] + c1o8 * (-data[0] - data[2] + data[4] + data[1] + data[3] - data[5]); // MX MY PZ
+// childCons[7].rhoS_2 = data[6] + c1o8 * (-data[0] - data[2] - data[4] + data[1] + data[3] + data[5]); // MX MY MZ
+// }
+// }
+// #endif // USE_PASSIVE_SCALAR
+//
+//#pragma unroll
+// for( uint childIndex = 0; childIndex < 8; childIndex++ ){
+//
+// uint childCellIndex = dataBase.coarseToFine[ COARSE_TO_FINE( index, ( 1 + childIndex ), dataBase.numberOfFineGhostCells ) ];
+//
+// dataBase.data[ RHO__(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rho ;
+// dataBase.data[ RHO_U(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoU;
+// dataBase.data[ RHO_V(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoV;
+// dataBase.data[ RHO_W(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoW;
+// dataBase.data[ RHO_E(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoE;
+// #ifdef USE_PASSIVE_SCALAR
+// dataBase.data[ RHO_S_1(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoS_1;
+// dataBase.data[ RHO_S_2(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoS_2;
+// #endif // USE_PASSIVE_SCALAR
+// }
+//}
+
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
+
__host__ __device__ inline void coarseToFineFunction( DataBaseStruct dataBase, uint startIndex, uint index )
{
index += startIndex;
@@ -68,179 +277,92 @@ __host__ __device__ inline void coarseToFineFunction( DataBaseStruct dataBase, u
cellToCell[4] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 4, dataBase.numberOfCells ) ];
cellToCell[5] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 5, dataBase.numberOfCells ) ];
+ ConservedVariables cons[7];
+
+ readCellData(cellToCell[0], dataBase, cons[0]);
+ readCellData(cellToCell[1], dataBase, cons[1]);
+ readCellData(cellToCell[2], dataBase, cons[2]);
+ readCellData(cellToCell[3], dataBase, cons[3]);
+ readCellData(cellToCell[4], dataBase, cons[4]);
+ readCellData(cellToCell[5], dataBase, cons[5]);
+ readCellData(cellIndex, dataBase, cons[6]);
+
ConservedVariables childCons [8];
+ ConservedVariables zeroCons;
+
+ // PX PY PZ MX MY MZ
+ childCons[0] = cons[6] + c1o8 * ( zeroCons + cons[0] + cons[2] + cons[4] - cons[1] - cons[3] - cons[5] ); // PX PY PZ
+ childCons[1] = cons[6] + c1o8 * ( zeroCons + cons[0] + cons[2] - cons[4] - cons[1] - cons[3] + cons[5] ); // PX PY MZ
+ childCons[2] = cons[6] + c1o8 * ( zeroCons + cons[0] - cons[2] + cons[4] - cons[1] + cons[3] - cons[5] ); // PX MY PZ
+ childCons[3] = cons[6] + c1o8 * ( zeroCons + cons[0] - cons[2] - cons[4] - cons[1] + cons[3] + cons[5] ); // PX MY MZ
+ childCons[4] = cons[6] + c1o8 * ( zeroCons - cons[0] + cons[2] + cons[4] + cons[1] - cons[3] - cons[5] ); // MX PY PZ
+ childCons[5] = cons[6] + c1o8 * ( zeroCons - cons[0] + cons[2] - cons[4] + cons[1] - cons[3] + cons[5] ); // MX PY MZ
+ childCons[6] = cons[6] + c1o8 * ( zeroCons - cons[0] - cons[2] + cons[4] + cons[1] + cons[3] - cons[5] ); // MX MY PZ
+ childCons[7] = cons[6] + c1o8 * ( zeroCons - cons[0] - cons[2] - cons[4] + cons[1] + cons[3] + cons[5] ); // MX MY MZ
- {
- real data [7];
-
- data[0] = dataBase.data[ RHO__(cellToCell[0], dataBase.numberOfCells) ];
- data[1] = dataBase.data[ RHO__(cellToCell[1], dataBase.numberOfCells) ];
- data[2] = dataBase.data[ RHO__(cellToCell[2], dataBase.numberOfCells) ];
- data[3] = dataBase.data[ RHO__(cellToCell[3], dataBase.numberOfCells) ];
- data[4] = dataBase.data[ RHO__(cellToCell[4], dataBase.numberOfCells) ];
- data[5] = dataBase.data[ RHO__(cellToCell[5], dataBase.numberOfCells) ];
- data[6] = dataBase.data[ RHO__(cellIndex , dataBase.numberOfCells) ];
-
- // PX PY PZ MX MY MZ
- childCons[0].rho = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
- childCons[1].rho = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
- childCons[2].rho = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
- childCons[3].rho = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
- childCons[4].rho = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
- childCons[5].rho = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
- childCons[6].rho = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
- childCons[7].rho = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
- }
+#pragma unroll
+ for( uint childIndex = 0; childIndex < 8; childIndex++ ){
- {
- real data [7];
-
- data[0] = dataBase.data[ RHO_U(cellToCell[0], dataBase.numberOfCells) ];
- data[1] = dataBase.data[ RHO_U(cellToCell[1], dataBase.numberOfCells) ];
- data[2] = dataBase.data[ RHO_U(cellToCell[2], dataBase.numberOfCells) ];
- data[3] = dataBase.data[ RHO_U(cellToCell[3], dataBase.numberOfCells) ];
- data[4] = dataBase.data[ RHO_U(cellToCell[4], dataBase.numberOfCells) ];
- data[5] = dataBase.data[ RHO_U(cellToCell[5], dataBase.numberOfCells) ];
- data[6] = dataBase.data[ RHO_U(cellIndex , dataBase.numberOfCells) ];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoU = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
- childCons[1].rhoU = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
- childCons[2].rhoU = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
- childCons[3].rhoU = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
- childCons[4].rhoU = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
- childCons[5].rhoU = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
- childCons[6].rhoU = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
- childCons[7].rhoU = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
- }
+ uint childCellIndex = dataBase.coarseToFine[ COARSE_TO_FINE( index, ( 1 + childIndex ), dataBase.numberOfFineGhostCells ) ];
- {
- real data [7];
-
- data[0] = dataBase.data[ RHO_V(cellToCell[0], dataBase.numberOfCells) ];
- data[1] = dataBase.data[ RHO_V(cellToCell[1], dataBase.numberOfCells) ];
- data[2] = dataBase.data[ RHO_V(cellToCell[2], dataBase.numberOfCells) ];
- data[3] = dataBase.data[ RHO_V(cellToCell[3], dataBase.numberOfCells) ];
- data[4] = dataBase.data[ RHO_V(cellToCell[4], dataBase.numberOfCells) ];
- data[5] = dataBase.data[ RHO_V(cellToCell[5], dataBase.numberOfCells) ];
- data[6] = dataBase.data[ RHO_V(cellIndex , dataBase.numberOfCells) ];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoV = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
- childCons[1].rhoV = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
- childCons[2].rhoV = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
- childCons[3].rhoV = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
- childCons[4].rhoV = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
- childCons[5].rhoV = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
- childCons[6].rhoV = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
- childCons[7].rhoV = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
+ writeCellData(childCellIndex, dataBase, childCons[childIndex]);
}
+}
- {
- real data [7];
-
- data[0] = dataBase.data[ RHO_W(cellToCell[0], dataBase.numberOfCells) ];
- data[1] = dataBase.data[ RHO_W(cellToCell[1], dataBase.numberOfCells) ];
- data[2] = dataBase.data[ RHO_W(cellToCell[2], dataBase.numberOfCells) ];
- data[3] = dataBase.data[ RHO_W(cellToCell[3], dataBase.numberOfCells) ];
- data[4] = dataBase.data[ RHO_W(cellToCell[4], dataBase.numberOfCells) ];
- data[5] = dataBase.data[ RHO_W(cellToCell[5], dataBase.numberOfCells) ];
- data[6] = dataBase.data[ RHO_W(cellIndex , dataBase.numberOfCells) ];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoW = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
- childCons[1].rhoW = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
- childCons[2].rhoW = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
- childCons[3].rhoW = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
- childCons[4].rhoW = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
- childCons[5].rhoW = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
- childCons[6].rhoW = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
- childCons[7].rhoW = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
- }
+__host__ __device__ inline void coarseToFineFunctionPrimitiveInterpolation( DataBaseStruct dataBase, uint startIndex, uint index )
+{
+ index += startIndex;
- {
- real data [7];
-
- data[0] = dataBase.data[ RHO_E(cellToCell[0], dataBase.numberOfCells) ];
- data[1] = dataBase.data[ RHO_E(cellToCell[1], dataBase.numberOfCells) ];
- data[2] = dataBase.data[ RHO_E(cellToCell[2], dataBase.numberOfCells) ];
- data[3] = dataBase.data[ RHO_E(cellToCell[3], dataBase.numberOfCells) ];
- data[4] = dataBase.data[ RHO_E(cellToCell[4], dataBase.numberOfCells) ];
- data[5] = dataBase.data[ RHO_E(cellToCell[5], dataBase.numberOfCells) ];
- data[6] = dataBase.data[ RHO_E(cellIndex , dataBase.numberOfCells) ];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoE = data[6] + c1o8 * ( + data[0] + data[2] + data[4] - data[1] - data[3] - data[5] ); // PX PY PZ
- childCons[1].rhoE = data[6] + c1o8 * ( + data[0] + data[2] - data[4] - data[1] - data[3] + data[5] ); // PX PY MZ
- childCons[2].rhoE = data[6] + c1o8 * ( + data[0] - data[2] + data[4] - data[1] + data[3] - data[5] ); // PX MY PZ
- childCons[3].rhoE = data[6] + c1o8 * ( + data[0] - data[2] - data[4] - data[1] + data[3] + data[5] ); // PX MY MZ
- childCons[4].rhoE = data[6] + c1o8 * ( - data[0] + data[2] + data[4] + data[1] - data[3] - data[5] ); // MX PY PZ
- childCons[5].rhoE = data[6] + c1o8 * ( - data[0] + data[2] - data[4] + data[1] - data[3] + data[5] ); // MX PY MZ
- childCons[6].rhoE = data[6] + c1o8 * ( - data[0] - data[2] + data[4] + data[1] + data[3] - data[5] ); // MX MY PZ
- childCons[7].rhoE = data[6] + c1o8 * ( - data[0] - data[2] - data[4] + data[1] + data[3] + data[5] ); // MX MY MZ
- }
+ uint cellIndex = dataBase.coarseToFine[ COARSE_TO_FINE( index, 0, dataBase.numberOfFineGhostCells ) ];
- #ifdef USE_PASSIVE_SCALAR
- {
- {
- real data[7];
-
- data[0] = dataBase.data[RHO_S_1(cellToCell[0], dataBase.numberOfCells)];
- data[1] = dataBase.data[RHO_S_1(cellToCell[1], dataBase.numberOfCells)];
- data[2] = dataBase.data[RHO_S_1(cellToCell[2], dataBase.numberOfCells)];
- data[3] = dataBase.data[RHO_S_1(cellToCell[3], dataBase.numberOfCells)];
- data[4] = dataBase.data[RHO_S_1(cellToCell[4], dataBase.numberOfCells)];
- data[5] = dataBase.data[RHO_S_1(cellToCell[5], dataBase.numberOfCells)];
- data[6] = dataBase.data[RHO_S_1(cellIndex, dataBase.numberOfCells)];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoS_1 = data[6] + c1o8 * (+data[0] + data[2] + data[4] - data[1] - data[3] - data[5]); // PX PY PZ
- childCons[1].rhoS_1 = data[6] + c1o8 * (+data[0] + data[2] - data[4] - data[1] - data[3] + data[5]); // PX PY MZ
- childCons[2].rhoS_1 = data[6] + c1o8 * (+data[0] - data[2] + data[4] - data[1] + data[3] - data[5]); // PX MY PZ
- childCons[3].rhoS_1 = data[6] + c1o8 * (+data[0] - data[2] - data[4] - data[1] + data[3] + data[5]); // PX MY MZ
- childCons[4].rhoS_1 = data[6] + c1o8 * (-data[0] + data[2] + data[4] + data[1] - data[3] - data[5]); // MX PY PZ
- childCons[5].rhoS_1 = data[6] + c1o8 * (-data[0] + data[2] - data[4] + data[1] - data[3] + data[5]); // MX PY MZ
- childCons[6].rhoS_1 = data[6] + c1o8 * (-data[0] - data[2] + data[4] + data[1] + data[3] - data[5]); // MX MY PZ
- childCons[7].rhoS_1 = data[6] + c1o8 * (-data[0] - data[2] - data[4] + data[1] + data[3] + data[5]); // MX MY MZ
- }
-
- {
- real data[7];
-
- data[0] = dataBase.data[RHO_S_2(cellToCell[0], dataBase.numberOfCells)];
- data[1] = dataBase.data[RHO_S_2(cellToCell[1], dataBase.numberOfCells)];
- data[2] = dataBase.data[RHO_S_2(cellToCell[2], dataBase.numberOfCells)];
- data[3] = dataBase.data[RHO_S_2(cellToCell[3], dataBase.numberOfCells)];
- data[4] = dataBase.data[RHO_S_2(cellToCell[4], dataBase.numberOfCells)];
- data[5] = dataBase.data[RHO_S_2(cellToCell[5], dataBase.numberOfCells)];
- data[6] = dataBase.data[RHO_S_2(cellIndex, dataBase.numberOfCells)];
-
- // PX PY PZ MX MY MZ
- childCons[0].rhoS_2 = data[6] + c1o8 * (+data[0] + data[2] + data[4] - data[1] - data[3] - data[5]); // PX PY PZ
- childCons[1].rhoS_2 = data[6] + c1o8 * (+data[0] + data[2] - data[4] - data[1] - data[3] + data[5]); // PX PY MZ
- childCons[2].rhoS_2 = data[6] + c1o8 * (+data[0] - data[2] + data[4] - data[1] + data[3] - data[5]); // PX MY PZ
- childCons[3].rhoS_2 = data[6] + c1o8 * (+data[0] - data[2] - data[4] - data[1] + data[3] + data[5]); // PX MY MZ
- childCons[4].rhoS_2 = data[6] + c1o8 * (-data[0] + data[2] + data[4] + data[1] - data[3] - data[5]); // MX PY PZ
- childCons[5].rhoS_2 = data[6] + c1o8 * (-data[0] + data[2] - data[4] + data[1] - data[3] + data[5]); // MX PY MZ
- childCons[6].rhoS_2 = data[6] + c1o8 * (-data[0] - data[2] + data[4] + data[1] + data[3] - data[5]); // MX MY PZ
- childCons[7].rhoS_2 = data[6] + c1o8 * (-data[0] - data[2] - data[4] + data[1] + data[3] + data[5]); // MX MY MZ
- }
- }
- #endif // USE_PASSIVE_SCALAR
+ uint cellToCell [6];
+
+ cellToCell[0] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 0, dataBase.numberOfCells ) ];
+ cellToCell[1] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 1, dataBase.numberOfCells ) ];
+ cellToCell[2] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 2, dataBase.numberOfCells ) ];
+ cellToCell[3] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 3, dataBase.numberOfCells ) ];
+ cellToCell[4] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 4, dataBase.numberOfCells ) ];
+ cellToCell[5] = dataBase.cellToCell[ CELL_TO_CELL( cellIndex, 5, dataBase.numberOfCells ) ];
+
+ PrimitiveVariables prim [7];
+ ConservedVariables cons[7];
+
+ readCellData(cellToCell[0], dataBase, cons[0]);
+ readCellData(cellToCell[1], dataBase, cons[1]);
+ readCellData(cellToCell[2], dataBase, cons[2]);
+ readCellData(cellToCell[3], dataBase, cons[3]);
+ readCellData(cellToCell[4], dataBase, cons[4]);
+ readCellData(cellToCell[5], dataBase, cons[5]);
+ readCellData(cellIndex, dataBase, cons[6]);
+
+ prim[0] = toPrimitiveVariables(cons[0], two);
+ prim[1] = toPrimitiveVariables(cons[1], two);
+ prim[2] = toPrimitiveVariables(cons[2], two);
+ prim[3] = toPrimitiveVariables(cons[3], two);
+ prim[4] = toPrimitiveVariables(cons[4], two);
+ prim[5] = toPrimitiveVariables(cons[5], two);
+ prim[6] = toPrimitiveVariables(cons[6], two);
+
+ PrimitiveVariables childPrim [8];
+ PrimitiveVariables zeroPrim;
+
+ // PX PY PZ MX MY MZ
+ childPrim[0] = prim[6] + c1o8 * ( zeroPrim + prim[0] + prim[2] + prim[4] - prim[1] - prim[3] - prim[5] ); // PX PY PZ
+ childPrim[1] = prim[6] + c1o8 * ( zeroPrim + prim[0] + prim[2] - prim[4] - prim[1] - prim[3] + prim[5] ); // PX PY MZ
+ childPrim[2] = prim[6] + c1o8 * ( zeroPrim + prim[0] - prim[2] + prim[4] - prim[1] + prim[3] - prim[5] ); // PX MY PZ
+ childPrim[3] = prim[6] + c1o8 * ( zeroPrim + prim[0] - prim[2] - prim[4] - prim[1] + prim[3] + prim[5] ); // PX MY MZ
+ childPrim[4] = prim[6] + c1o8 * ( zeroPrim - prim[0] + prim[2] + prim[4] + prim[1] - prim[3] - prim[5] ); // MX PY PZ
+ childPrim[5] = prim[6] + c1o8 * ( zeroPrim - prim[0] + prim[2] - prim[4] + prim[1] - prim[3] + prim[5] ); // MX PY MZ
+ childPrim[6] = prim[6] + c1o8 * ( zeroPrim - prim[0] - prim[2] + prim[4] + prim[1] + prim[3] - prim[5] ); // MX MY PZ
+ childPrim[7] = prim[6] + c1o8 * ( zeroPrim - prim[0] - prim[2] - prim[4] + prim[1] + prim[3] + prim[5] ); // MX MY MZ
#pragma unroll
for( uint childIndex = 0; childIndex < 8; childIndex++ ){
uint childCellIndex = dataBase.coarseToFine[ COARSE_TO_FINE( index, ( 1 + childIndex ), dataBase.numberOfFineGhostCells ) ];
- dataBase.data[ RHO__(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rho ;
- dataBase.data[ RHO_U(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoU;
- dataBase.data[ RHO_V(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoV;
- dataBase.data[ RHO_W(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoW;
- dataBase.data[ RHO_E(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoE;
- #ifdef USE_PASSIVE_SCALAR
- dataBase.data[ RHO_S_1(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoS_1;
- dataBase.data[ RHO_S_2(childCellIndex, dataBase.numberOfCells) ] = childCons[childIndex].rhoS_2;
- #endif // USE_PASSIVE_SCALAR
+ ConservedVariables childCons = toConservedVariables(childPrim[childIndex], two);
+
+ writeCellData(childCellIndex, dataBase, childCons);
}
}
diff --git a/src/GksGpu/Interface/FineToCoarseKernel.cu b/src/GksGpu/Interface/FineToCoarseKernel.cu
index 5a6a742d5..aaaeff5da 100644
--- a/src/GksGpu/Interface/FineToCoarseKernel.cu
+++ b/src/GksGpu/Interface/FineToCoarseKernel.cu
@@ -9,6 +9,9 @@
#include "DataBase/DataBaseStruct.h"
#include "FlowStateData/FlowStateData.cuh"
+#include "FlowStateData/FlowStateDataConversion.cuh"
+
+#include "FlowStateData/AccessDeviceData.cuh"
#include "Definitions/PassiveScalar.h"
#include "Definitions/MemoryAccessPattern.h"
@@ -19,7 +22,8 @@
__global__ void fineToCoarseKernel ( DataBaseStruct dataBase, uint startIndex, uint numberOfEntities );
-__host__ __device__ inline void fineToCoarseFunction( DataBaseStruct dataBase, uint startIndex, uint index );
+__host__ __device__ inline void fineToCoarseFunction ( DataBaseStruct dataBase, uint startIndex, uint index );
+__host__ __device__ inline void fineToCoarseFunctionPrimitiveInterpolation( DataBaseStruct dataBase, uint startIndex, uint index );
//////////////////////////////////////////////////////////////////////////
//////////////////////////////////////////////////////////////////////////
@@ -51,39 +55,57 @@ __global__ void fineToCoarseKernel( DataBaseStruct dataBase, uint startIndex, ui
if( index >= numberOfEntities ) return;
fineToCoarseFunction( dataBase, startIndex, index );
+ //fineToCoarseFunctionPrimitiveInterpolation( dataBase, startIndex, index );
}
__host__ __device__ inline void fineToCoarseFunction( DataBaseStruct dataBase, uint startIndex, uint index )
{
index += startIndex;
- ConservedVariables cons;
+ ConservedVariables parentCons;
#pragma unroll
for( uint childIdx = 1; childIdx < LENGTH_FINE_TO_COARSE; childIdx++ ){
uint cellIdx = dataBase.fineToCoarse[ FINE_TO_COARSE( index, childIdx, dataBase.numberOfCoarseGhostCells ) ];
- cons.rho += c1o8 * dataBase.data[ RHO__(cellIdx, dataBase.numberOfCells) ];
- cons.rhoU += c1o8 * dataBase.data[ RHO_U(cellIdx, dataBase.numberOfCells) ];
- cons.rhoV += c1o8 * dataBase.data[ RHO_V(cellIdx, dataBase.numberOfCells) ];
- cons.rhoW += c1o8 * dataBase.data[ RHO_W(cellIdx, dataBase.numberOfCells) ];
- cons.rhoE += c1o8 * dataBase.data[ RHO_E(cellIdx, dataBase.numberOfCells) ];
- #ifdef USE_PASSIVE_SCALAR
- cons.rhoS_1 += c1o8 * dataBase.data[ RHO_S_1(cellIdx, dataBase.numberOfCells) ];
- cons.rhoS_2 += c1o8 * dataBase.data[ RHO_S_2(cellIdx, dataBase.numberOfCells) ];
- #endif // USE_PASSIVE_SCALAR
+ ConservedVariables cons;
+
+ readCellData( cellIdx, dataBase, cons );
+
+ parentCons = parentCons + c1o8 * cons;
+ }
+
+ {
+ uint cellIdx = dataBase.fineToCoarse[FINE_TO_COARSE(index, 0, dataBase.numberOfCoarseGhostCells)];
+
+ writeCellData(cellIdx, dataBase, parentCons);
}
+}
+
+__host__ __device__ inline void fineToCoarseFunctionPrimitiveInterpolation( DataBaseStruct dataBase, uint startIndex, uint index )
+{
+ index += startIndex;
+
+ PrimitiveVariables parentPrim;
- uint cellIdx = dataBase.fineToCoarse[ FINE_TO_COARSE( index, 0, dataBase.numberOfCoarseGhostCells ) ];
-
- dataBase.data[ RHO__(cellIdx, dataBase.numberOfCells) ] = cons.rho ;
- dataBase.data[ RHO_U(cellIdx, dataBase.numberOfCells) ] = cons.rhoU;
- dataBase.data[ RHO_V(cellIdx, dataBase.numberOfCells) ] = cons.rhoV;
- dataBase.data[ RHO_W(cellIdx, dataBase.numberOfCells) ] = cons.rhoW;
- dataBase.data[ RHO_E(cellIdx, dataBase.numberOfCells) ] = cons.rhoE;
-#ifdef USE_PASSIVE_SCALAR
- dataBase.data[ RHO_S_1(cellIdx, dataBase.numberOfCells) ] = cons.rhoS_1;
- dataBase.data[ RHO_S_2(cellIdx, dataBase.numberOfCells) ] = cons.rhoS_2;
-#endif // USE_PASSIVE_SCALAR
+#pragma unroll
+ for( uint childIdx = 1; childIdx < LENGTH_FINE_TO_COARSE; childIdx++ ){
+
+ uint cellIdx = dataBase.fineToCoarse[ FINE_TO_COARSE( index, childIdx, dataBase.numberOfCoarseGhostCells ) ];
+
+ ConservedVariables cons;
+
+ readCellData( cellIdx, dataBase, cons );
+
+ parentPrim = parentPrim + c1o8 * toPrimitiveVariables(cons, two);
+ }
+
+ {
+ uint cellIdx = dataBase.fineToCoarse[FINE_TO_COARSE(index, 0, dataBase.numberOfCoarseGhostCells)];
+
+ ConservedVariables parentCons = toConservedVariables(parentPrim, two);
+
+ writeCellData(cellIdx, dataBase, parentCons);
+ }
}
diff --git a/targets/apps/GKS/PoolFire/PoolFire.cpp b/targets/apps/GKS/PoolFire/PoolFire.cpp
index be457327c..320e47acc 100644
--- a/targets/apps/GKS/PoolFire/PoolFire.cpp
+++ b/targets/apps/GKS/PoolFire/PoolFire.cpp
@@ -127,7 +127,7 @@ void thermalCavity( std::string path, std::string simulationName )
////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////
- bool threeDimensional = true;
+ bool threeDimensional = false;
if( threeDimensional )
{
--
GitLab