implements overloaded operators for FlowStateData und simplifies the interface interpolations

b50e017c · LEGOLAS\lenz · 1364a0f5 · b50e017c · b50e017c · b50e017c
Commit b50e017c authored 5 years ago by $LEGOLAS\lenz's avatar$ LEGOLAS\lenz
--- 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;
            }
        }
    }

--- 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

--- a/src/GksGpu/Interface/CoarseToFineKernel.cu
+++ b/src/GksGpu/Interface/CoarseToFineKernel.cu
--- 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);
+    }
 }
--- 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 )
    {