diff --git a/CMake/cmake_config_files/ARAGORNUBUNTU.config.cmake b/CMake/cmake_config_files/ARAGORNUBUNTU.config.cmake index 90debb1a4c270109d4dfbb455f21253b3a6754b8..cd50d08b1cadf0f54db1a4554dc3c534b3df6993 100644 --- a/CMake/cmake_config_files/ARAGORNUBUNTU.config.cmake +++ b/CMake/cmake_config_files/ARAGORNUBUNTU.config.cmake @@ -13,5 +13,6 @@ set(GPU_APP "apps/gpu/LBM/") list(APPEND USER_APPS "${GPU_APP}DrivenCavityMultiGPU" "${GPU_APP}SphereScaling" + "${GPU_APP}ActuatorLine" # "${GPU_APP}MusselOyster" ) \ No newline at end of file diff --git a/CMake/cmake_config_files/PHOENIX.config.cmake b/CMake/cmake_config_files/PHOENIX.config.cmake index d31d8684a53a769e48408ad5febe7d2c6b22c623..43f29d9d02a5de241905de9d4b739ec1c620b4ad 100644 --- a/CMake/cmake_config_files/PHOENIX.config.cmake +++ b/CMake/cmake_config_files/PHOENIX.config.cmake @@ -31,4 +31,5 @@ list(APPEND USER_APPS # "${GPU_APP}DrivenCavityMultiGPU" # "${GPU_APP}SphereScaling" # "${GPU_APP}MusselOyster" + "${GPU_APP}ActuatorLine" ) diff --git a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp index 9d982ebac0059b4512041194100f6e1fdfa61924..8b5502329a542a8ce8259a08f6e67f4e491ab85f 100644 --- a/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp +++ b/apps/gpu/LBM/ActuatorLine/ActuatorLine.cpp @@ -39,6 +39,7 @@ #include <fstream> #include <exception> #include <memory> +#include <filesystem> ////////////////////////////////////////////////////////////////////////// @@ -60,6 +61,8 @@ #include "GridGenerator/grid/GridBuilder/MultipleGridBuilder.h" #include "GridGenerator/grid/BoundaryConditions/Side.h" #include "GridGenerator/grid/BoundaryConditions/BoundaryCondition.h" +#include "GridGenerator/geometries/TriangularMesh/TriangularMesh.h" + #include "GridGenerator/grid/GridFactory.h" @@ -96,24 +99,20 @@ //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// -LbmOrGks lbmOrGks = LBM; - -std::string path("."); - -std::string simulationName("ActuatorLine"); - -//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// -//////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - void multipleLevel(const std::string& configPath) { + const LbmOrGks lbmOrGks = LBM; + const std::string path("."); + const std::string simulationName("ActuatorLine"); + logging::Logger::addStream(&std::cout); logging::Logger::setDebugLevel(logging::Logger::Level::INFO_LOW); logging::Logger::timeStamp(logging::Logger::ENABLE); logging::Logger::enablePrintedRankNumbers(logging::Logger::ENABLE); + VF_LOG_INFO("Start Preprocessing"); + vf::gpu::Communicator& communicator = vf::gpu::Communicator::getInstance(); auto gridFactory = GridFactory::make(); @@ -127,9 +126,9 @@ void multipleLevel(const std::string& configPath) const real velocity = config.getValue<real>("Velocity"); - const real L_x = 24*reference_diameter; - const real L_y = 6*reference_diameter; - const real L_z = 6*reference_diameter; + const real L_x = 3.5*reference_diameter; + const real L_y = 2.5*reference_diameter; + const real L_z = 3*reference_diameter; const real viscosity = 1.56e-5; @@ -140,36 +139,40 @@ void multipleLevel(const std::string& configPath) const float tOut = config.getValue<real>("tOut"); const float tEnd = config.getValue<real>("tEnd"); // total time of simulation - const float tStartAveraging = config.getValue<real>("tStartAveraging"); - const float tStartTmpAveraging = config.getValue<real>("tStartTmpAveraging"); - const float tAveraging = config.getValue<real>("tAveraging"); - const float tStartOutProbe = config.getValue<real>("tStartOutProbe"); - const float tOutProbe = config.getValue<real>("tOutProbe"); - + // const float tStartAveraging = config.getValue<real>("tStartAveraging"); + // const float tStartTmpAveraging = config.getValue<real>("tStartTmpAveraging"); + // const float tAveraging = config.getValue<real>("tAveraging"); + // const float tStartOutProbe = config.getValue<real>("tStartOutProbe"); + // const float tOutProbe = config.getValue<real>("tOutProbe"); + SPtr<Parameter> para = std::make_shared<Parameter>(communicator.getNummberOfProcess(), communicator.getPID(), &config); BoundaryConditionFactory bcFactory = BoundaryConditionFactory(); GridScalingFactory scalingFactory = GridScalingFactory(); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - const real dx = reference_diameter/real(nodes_per_diameter); + const real dx = reference_diameter/real(nodes_per_diameter); - real turbPos[3] = {3*reference_diameter, 3*reference_diameter, 3*reference_diameter}; + real turbPos[3] = {0.7f*reference_diameter, reference_diameter, 1.5f*reference_diameter}; - gridBuilder->addCoarseGrid(0.0, 0.0, 0.0, - L_x, L_y, L_z, dx); + gridBuilder->addCoarseGrid(0.0, 0.0, 0.0, + L_x, L_y, L_z, dx); gridBuilder->setNumberOfLayers(4,0); - gridBuilder->addGrid( new Cuboid( turbPos[0]-1.5*reference_diameter, turbPos[1]-1.5*reference_diameter, turbPos[2]-1.5*reference_diameter, - turbPos[0]+10.0*reference_diameter, turbPos[1]+1.5*reference_diameter, turbPos[2]+1.5*reference_diameter) , 1 ); + gridBuilder->addGrid( new Cuboid( turbPos[0]-0.3*reference_diameter, turbPos[1]-1*reference_diameter, turbPos[2]-0.7*reference_diameter, + turbPos[0]+2.0*reference_diameter, turbPos[1]+0.7*reference_diameter, turbPos[2]+0.7*reference_diameter) , 1 ); para->setMaxLevel(2); scalingFactory.setScalingFactory(GridScalingFactory::GridScaling::ScaleCompressible); - gridBuilder->setPeriodicBoundaryCondition(false, false, false); + // std::string stlPath = "./VirtualFluids_dev/apps/gpu/LBM/ActuatorLine/Pole.stl"; + // Object *sphere = TriangularMesh::make(stlPath); + // gridBuilder->addGeometry(sphere); + + gridBuilder->setPeriodicBoundaryCondition(false, false, false); - gridBuilder->buildGrids(lbmOrGks, false); // buildGrids() has to be called before setting the BCs!!!! + gridBuilder->buildGrids(lbmOrGks, false); // buildGrids() has to be called before setting the BCs!!!! - //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// + //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// const real dt = dx * mach / (sqrt(3) * velocity); @@ -177,8 +180,16 @@ void multipleLevel(const std::string& configPath) const real viscosityLB = viscosity * dt / (dx * dx); // LB units + VF_LOG_INFO("dx (coarse grid) [m] = {}", dx); + VF_LOG_INFO("dt [s] = {}", dt); + VF_LOG_INFO("Lx [m] = {}", L_x); + VF_LOG_INFO("Ly [m] = {}", L_y); + VF_LOG_INFO("Lz [m] = {}", L_z); + VF_LOG_INFO("velocity [m/s] = {}", velocity); VF_LOG_INFO("velocity [dx/dt] = {}", velocityLB); VF_LOG_INFO("viscosity [10^8 dx^2/dt] = {}", viscosityLB*1e8); + VF_LOG_INFO("Ma = {}", mach); + VF_LOG_INFO("nodes/turbine diameter = {}", reference_diameter/dx); //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// @@ -192,7 +203,8 @@ void multipleLevel(const std::string& configPath) para->setViscosityLB(viscosityLB); para->setVelocityRatio( dx / dt ); para->setViscosityRatio( dx*dx/dt ); - para->setMainKernel("CumulantK17"); + para->setDensityRatio( 1.0 ); + para->setMainKernel("CumulantK17Sponge"); para->setInitialCondition([&](real coordX, real coordY, real coordZ, real &rho, real &vx, real &vy, real &vz) { rho = (real)0.0; @@ -201,39 +213,48 @@ void multipleLevel(const std::string& configPath) vz = (real)0.0; }); - para->setTimestepStartOut( uint(tStartOut/dt) ); - para->setTimestepOut( uint(tOut/dt) ); - para->setTimestepEnd( uint(tEnd/dt) ); + para->setTimestepStartOut( uint(tStartOut) ); + para->setTimestepOut( uint(tOut) ); + para->setTimestepEnd( uint(tEnd) ); + + // para->setTimestepStartOut( uint(tStartOut/dt) ); + // para->setTimestepOut( uint(tOut/dt) ); + // para->setTimestepEnd( uint(tEnd/dt) ); para->setIsBodyForce( true ); para->setUseStreams( true ); ///////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// - gridBuilder->setVelocityBoundaryCondition(SideType::MX, velocityLB, 0.0, 0.0); - - gridBuilder->setVelocityBoundaryCondition(SideType::MY, velocityLB, 0.0, 0.0); - gridBuilder->setVelocityBoundaryCondition(SideType::PY, velocityLB, 0.0, 0.0); - gridBuilder->setVelocityBoundaryCondition(SideType::MZ, velocityLB, 0.0, 0.0); - gridBuilder->setVelocityBoundaryCondition(SideType::PZ, velocityLB, 0.0, 0.0); + gridBuilder->setVelocityBoundaryCondition(SideType::MX, velocityLB, 0.0, 0.0); + gridBuilder->setSlipBoundaryCondition(SideType::MY,0.0, 0.0, 0.0 ); + gridBuilder->setVelocityBoundaryCondition(SideType::PY, velocityLB, 0.0, 0.0); + gridBuilder->setVelocityBoundaryCondition(SideType::MZ, velocityLB, 0.0, 0.0); + gridBuilder->setVelocityBoundaryCondition(SideType::PZ, velocityLB, 0.0, 0.0); gridBuilder->setPressureBoundaryCondition(SideType::PX, 0.0); + // gridBuilder->setVelocityBoundaryCondition(SideType::GEOMETRY, 0.0, 0.0, 0.0); + // bcFactory.setGeometryBoundaryCondition(BoundaryConditionFactory::NoSlipBC::NoSlipCompressible); + + bcFactory.setSlipBoundaryCondition(BoundaryConditionFactory::SlipBC::SlipCompressible); bcFactory.setVelocityBoundaryCondition(BoundaryConditionFactory::VelocityBC::VelocityAndPressureCompressible); bcFactory.setPressureBoundaryCondition(BoundaryConditionFactory::PressureBC::OutflowNonReflective); SPtr<TurbulenceModelFactory> tmFactory = std::make_shared<TurbulenceModelFactory>(para); tmFactory->readConfigFile(config); + // gridBuilder->writeGridsToVtk(para->getOutputPath() + "Grid" + "_"); + //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// int level = 1; // grid level at which the turbine samples velocities and distributes forces - const real epsilon = dx*exp2(-level)*1.5; // width of gaussian smearing + const real epsilon = dx*exp2(-level)*2.0; // width of gaussian smearing const real density = 1.225f; const uint nBlades = 3; - const uint nBladeNodes = 32; + const uint nBladeNodes = reference_diameter * 4 + 3; // passt zu auflösung von 105 Knoten + VF_LOG_INFO("number of blade nodes ALM = {}", nBladeNodes); const real tipspeed_ratio = 7.5f; // tipspeed ratio = angular vel * radius / inflow vel const real omega = 2*tipspeed_ratio*velocity/reference_diameter; - SPtr<ActuatorFarm> actuator_farm = std::make_shared<ActuatorFarm>(nBlades, density, nBladeNodes, epsilon, level, dt, dx, true); std::vector<real> bladeRadii; @@ -242,7 +263,6 @@ void multipleLevel(const std::string& configPath) actuator_farm->addTurbine(turbPos[0], turbPos[1], turbPos[2], reference_diameter, omega, 0, 0, bladeRadii); para->addActuator( actuator_farm ); - // SPtr<PointProbe> pointProbe = std::make_shared<PointProbe>("pointProbe", para->getOutputPath(), 100, 1, 500, 100); // std::vector<real> probeCoordsX = {reference_diameter,2*reference_diameter,5*reference_diameter}; // std::vector<real> probeCoordsY = {3*reference_diameter,3*reference_diameter,3*reference_diameter}; @@ -274,11 +294,20 @@ int main( int argc, char* argv[]) { try { - vf::logging::Logger::initalizeLogger(); + ////////////////////////////////////////////////////////////////////////// + // assuming that a config files is stored parallel to this file. + std::filesystem::path configPath = __FILE__; + + // the config file's default name can be replaced by passing a command line argument + std::string configName("configActuatorLine.txt"); + if (argc == 2) { + configName = argv[1]; + std::cout << "Using configFile command line argument: " << configName << std::endl; + } - if( argc > 1){ path = argv[1]; } + configPath.replace_filename(configName); - multipleLevel(path + "/configActuatorLine.txt"); + multipleLevel(configPath); } catch (const spdlog::spdlog_ex &ex) { std::cout << "Log initialization failed: " << ex.what() << std::endl; diff --git a/apps/gpu/LBM/ActuatorLine/Pole.stl b/apps/gpu/LBM/ActuatorLine/Pole.stl new file mode 100644 index 0000000000000000000000000000000000000000..f8e00efcec53eda3abf3e95e171a899447482c41 --- /dev/null +++ b/apps/gpu/LBM/ActuatorLine/Pole.stl @@ -0,0 +1,3166 @@ +solid Visualization Toolkit generated SLA File + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 105 126 189 + vertex 105 0 189 + vertex 104.989 126 188.614 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.989 0 188.614 + vertex 104.989 126 188.614 + vertex 105 0 189 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.989 126 188.614 + vertex 104.989 0 188.614 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.958 126 188.23 + vertex 104.989 0 188.614 + vertex 104.958 0 188.23 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.958 126 188.23 + vertex 104.958 0 188.23 + vertex 104.905 126 187.848 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.905 0 187.848 + vertex 104.905 126 187.848 + vertex 104.958 0 188.23 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.905 126 187.848 + vertex 104.905 0 187.848 + vertex 104.831 126 187.469 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.831 0 187.469 + vertex 104.831 126 187.469 + vertex 104.905 0 187.848 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.831 126 187.469 + vertex 104.831 0 187.469 + vertex 104.736 126 187.095 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.736 126 187.095 + vertex 104.831 0 187.469 + vertex 104.736 0 187.095 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.736 126 187.095 + vertex 104.736 0 187.095 + vertex 104.621 126 186.727 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.621 0 186.727 + vertex 104.621 126 186.727 + vertex 104.736 0 187.095 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.621 126 186.727 + vertex 104.621 0 186.727 + vertex 104.485 126 186.366 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.485 126 186.366 + vertex 104.621 0 186.727 + vertex 104.485 0 186.366 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.485 126 186.366 + vertex 104.485 0 186.366 + vertex 104.33 126 186.013 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.33 126 186.013 + vertex 104.485 0 186.366 + vertex 104.33 0 186.013 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.33 126 186.013 + vertex 104.33 0 186.013 + vertex 104.156 126 185.668 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.156 0 185.668 + vertex 104.156 126 185.668 + vertex 104.33 0 186.013 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.156 126 185.668 + vertex 104.156 0 185.668 + vertex 103.963 126 185.334 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.963 0 185.334 + vertex 103.963 126 185.334 + vertex 104.156 0 185.668 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.963 126 185.334 + vertex 103.963 0 185.334 + vertex 103.752 126 185.011 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.752 126 185.011 + vertex 103.963 0 185.334 + vertex 103.752 0 185.011 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.752 126 185.011 + vertex 103.752 0 185.011 + vertex 103.524 126 184.701 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.524 126 184.701 + vertex 103.752 0 185.011 + vertex 103.524 0 184.701 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.524 126 184.701 + vertex 103.524 0 184.701 + vertex 103.279 126 184.403 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.279 0 184.403 + vertex 103.279 126 184.403 + vertex 103.524 0 184.701 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.279 126 184.403 + vertex 103.279 0 184.403 + vertex 103.017 126 184.119 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.017 0 184.119 + vertex 103.017 126 184.119 + vertex 103.279 0 184.403 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.017 126 184.119 + vertex 103.017 0 184.119 + vertex 102.741 126 183.85 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.741 0 183.85 + vertex 102.741 126 183.85 + vertex 103.017 0 184.119 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.741 126 183.85 + vertex 102.741 0 183.85 + vertex 102.45 126 183.597 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.45 0 183.597 + vertex 102.45 126 183.597 + vertex 102.741 0 183.85 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.45 126 183.597 + vertex 102.45 0 183.597 + vertex 102.146 126 183.36 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.146 0 183.36 + vertex 102.146 126 183.36 + vertex 102.45 0 183.597 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.146 126 183.36 + vertex 102.146 0 183.36 + vertex 101.829 126 183.14 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.829 0 183.14 + vertex 101.829 126 183.14 + vertex 102.146 0 183.36 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.829 126 183.14 + vertex 101.829 0 183.14 + vertex 101.5 126 182.938 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.5 0 182.938 + vertex 101.5 126 182.938 + vertex 101.829 0 183.14 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.5 126 182.938 + vertex 101.5 0 182.938 + vertex 101.161 126 182.754 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.161 126 182.754 + vertex 101.5 0 182.938 + vertex 101.161 0 182.754 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.161 126 182.754 + vertex 101.161 0 182.754 + vertex 100.812 126 182.59 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.812 0 182.59 + vertex 100.812 126 182.59 + vertex 101.161 0 182.754 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.812 126 182.59 + vertex 100.812 0 182.59 + vertex 100.454 126 182.444 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.454 0 182.444 + vertex 100.454 126 182.444 + vertex 100.812 0 182.59 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.454 126 182.444 + vertex 100.454 0 182.444 + vertex 100.09 126 182.319 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.09 0 182.319 + vertex 100.09 126 182.319 + vertex 100.454 0 182.444 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.09 126 182.319 + vertex 100.09 0 182.319 + vertex 99.7184 126 182.214 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 99.7184 0 182.214 + vertex 99.7184 126 182.214 + vertex 100.09 0 182.319 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 99.7184 126 182.214 + vertex 99.7184 0 182.214 + vertex 99.342 126 182.13 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 99.342 0 182.13 + vertex 99.342 126 182.13 + vertex 99.7184 0 182.214 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 99.342 126 182.13 + vertex 99.342 0 182.13 + vertex 98.9615 126 182.066 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.9615 0 182.066 + vertex 98.9615 126 182.066 + vertex 99.342 0 182.13 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.9615 126 182.066 + vertex 98.9615 0 182.066 + vertex 98.5781 126 182.024 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.5781 0 182.024 + vertex 98.5781 126 182.024 + vertex 98.9615 0 182.066 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.5781 126 182.024 + vertex 98.5781 0 182.024 + vertex 98.1929 126 182.003 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.1929 0 182.003 + vertex 98.1929 126 182.003 + vertex 98.5781 0 182.024 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.1929 126 182.003 + vertex 98.1929 0 182.003 + vertex 97.8071 126 182.003 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.8071 0 182.003 + vertex 97.8071 126 182.003 + vertex 98.1929 0 182.003 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.8071 126 182.003 + vertex 97.8071 0 182.003 + vertex 97.4219 126 182.024 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.4219 0 182.024 + vertex 97.4219 126 182.024 + vertex 97.8071 0 182.003 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.4219 126 182.024 + vertex 97.4219 0 182.024 + vertex 97.0385 126 182.066 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.0385 0 182.066 + vertex 97.0385 126 182.066 + vertex 97.4219 0 182.024 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.0385 126 182.066 + vertex 97.0385 0 182.066 + vertex 96.658 126 182.13 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 96.658 0 182.13 + vertex 96.658 126 182.13 + vertex 97.0385 0 182.066 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 96.658 126 182.13 + vertex 96.658 0 182.13 + vertex 96.2816 126 182.214 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 96.2816 0 182.214 + vertex 96.2816 126 182.214 + vertex 96.658 0 182.13 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 96.2816 126 182.214 + vertex 96.2816 0 182.214 + vertex 95.9104 126 182.319 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.9104 0 182.319 + vertex 95.9104 126 182.319 + vertex 96.2816 0 182.214 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.9104 126 182.319 + vertex 95.9104 0 182.319 + vertex 95.5455 126 182.444 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.5455 0 182.444 + vertex 95.5455 126 182.444 + vertex 95.9104 0 182.319 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.5455 126 182.444 + vertex 95.5455 0 182.444 + vertex 95.1881 126 182.59 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.1881 0 182.59 + vertex 95.1881 126 182.59 + vertex 95.5455 0 182.444 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.1881 126 182.59 + vertex 95.1881 0 182.59 + vertex 94.8393 126 182.754 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.8393 0 182.754 + vertex 94.8393 126 182.754 + vertex 95.1881 0 182.59 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.8393 126 182.754 + vertex 94.8393 0 182.754 + vertex 94.5 126 182.938 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.5 126 182.938 + vertex 94.8393 0 182.754 + vertex 94.5 0 182.938 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.5 126 182.938 + vertex 94.5 0 182.938 + vertex 94.1714 126 183.14 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.1714 0 183.14 + vertex 94.1714 126 183.14 + vertex 94.5 0 182.938 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.1714 126 183.14 + vertex 94.1714 0 183.14 + vertex 93.8544 126 183.36 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.8544 0 183.36 + vertex 93.8544 126 183.36 + vertex 94.1714 0 183.14 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.8544 126 183.36 + vertex 93.8544 0 183.36 + vertex 93.5499 126 183.597 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.5499 0 183.597 + vertex 93.5499 126 183.597 + vertex 93.8544 0 183.36 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.5499 126 183.597 + vertex 93.5499 0 183.597 + vertex 93.259 126 183.85 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.259 0 183.85 + vertex 93.259 126 183.85 + vertex 93.5499 0 183.597 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.259 126 183.85 + vertex 93.259 0 183.85 + vertex 92.9825 126 184.119 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.9825 0 184.119 + vertex 92.9825 126 184.119 + vertex 93.259 0 183.85 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.9825 126 184.119 + vertex 92.9825 0 184.119 + vertex 92.7213 126 184.403 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.7213 0 184.403 + vertex 92.7213 126 184.403 + vertex 92.9825 0 184.119 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.7213 126 184.403 + vertex 92.7213 0 184.403 + vertex 92.476 126 184.701 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.476 0 184.701 + vertex 92.476 126 184.701 + vertex 92.7213 0 184.403 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.476 126 184.701 + vertex 92.476 0 184.701 + vertex 92.2476 126 185.011 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.2476 126 185.011 + vertex 92.476 0 184.701 + vertex 92.2476 0 185.011 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.2476 126 185.011 + vertex 92.2476 0 185.011 + vertex 92.0366 126 185.334 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.0366 126 185.334 + vertex 92.2476 0 185.011 + vertex 92.0366 0 185.334 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.0366 126 185.334 + vertex 92.0366 0 185.334 + vertex 91.8437 126 185.668 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.8437 0 185.668 + vertex 91.8437 126 185.668 + vertex 92.0366 0 185.334 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.8437 126 185.668 + vertex 91.8437 0 185.668 + vertex 91.6695 126 186.013 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.6695 0 186.013 + vertex 91.6695 126 186.013 + vertex 91.8437 0 185.668 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.6695 126 186.013 + vertex 91.6695 0 186.013 + vertex 91.5145 126 186.366 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.5145 126 186.366 + vertex 91.6695 0 186.013 + vertex 91.5145 0 186.366 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.5145 126 186.366 + vertex 91.5145 0 186.366 + vertex 91.3793 126 186.727 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.3793 126 186.727 + vertex 91.5145 0 186.366 + vertex 91.3793 0 186.727 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.3793 126 186.727 + vertex 91.3793 0 186.727 + vertex 91.2641 126 187.095 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.2641 0 187.095 + vertex 91.2641 126 187.095 + vertex 91.3793 0 186.727 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.2641 126 187.095 + vertex 91.2641 0 187.095 + vertex 91.1694 126 187.469 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.1694 126 187.469 + vertex 91.2641 0 187.095 + vertex 91.1694 0 187.469 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.1694 126 187.469 + vertex 91.1694 0 187.469 + vertex 91.0955 126 187.848 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0955 0 187.848 + vertex 91.0955 126 187.848 + vertex 91.1694 0 187.469 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0955 126 187.848 + vertex 91.0955 0 187.848 + vertex 91.0425 126 188.23 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0425 0 188.23 + vertex 91.0425 126 188.23 + vertex 91.0955 0 187.848 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0425 126 188.23 + vertex 91.0425 0 188.23 + vertex 91.0106 126 188.614 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0106 126 188.614 + vertex 91.0425 0 188.23 + vertex 91.0106 0 188.614 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0106 126 188.614 + vertex 91.0106 0 188.614 + vertex 91 126 189 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91 0 189 + vertex 91 126 189 + vertex 91.0106 0 188.614 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91 126 189 + vertex 91 0 189 + vertex 91.0106 126 189.386 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0106 0 189.386 + vertex 91.0106 126 189.386 + vertex 91 0 189 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0106 126 189.386 + vertex 91.0106 0 189.386 + vertex 91.0425 126 189.77 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0425 126 189.77 + vertex 91.0106 0 189.386 + vertex 91.0425 0 189.77 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0425 126 189.77 + vertex 91.0425 0 189.77 + vertex 91.0955 126 190.152 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.0955 0 190.152 + vertex 91.0955 126 190.152 + vertex 91.0425 0 189.77 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.0955 126 190.152 + vertex 91.0955 0 190.152 + vertex 91.1694 126 190.531 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.1694 0 190.531 + vertex 91.1694 126 190.531 + vertex 91.0955 0 190.152 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.1694 126 190.531 + vertex 91.1694 0 190.531 + vertex 91.2641 126 190.905 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.2641 126 190.905 + vertex 91.1694 0 190.531 + vertex 91.2641 0 190.905 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.2641 126 190.905 + vertex 91.2641 0 190.905 + vertex 91.3793 126 191.273 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.3793 0 191.273 + vertex 91.3793 126 191.273 + vertex 91.2641 0 190.905 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.3793 126 191.273 + vertex 91.3793 0 191.273 + vertex 91.5145 126 191.634 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.5145 126 191.634 + vertex 91.3793 0 191.273 + vertex 91.5145 0 191.634 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.5145 126 191.634 + vertex 91.5145 0 191.634 + vertex 91.6695 126 191.987 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.6695 126 191.987 + vertex 91.5145 0 191.634 + vertex 91.6695 0 191.987 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.6695 126 191.987 + vertex 91.6695 0 191.987 + vertex 91.8437 126 192.332 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 91.8437 0 192.332 + vertex 91.8437 126 192.332 + vertex 91.6695 0 191.987 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 91.8437 126 192.332 + vertex 91.8437 0 192.332 + vertex 92.0366 126 192.666 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.0366 0 192.666 + vertex 92.0366 126 192.666 + vertex 91.8437 0 192.332 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.0366 126 192.666 + vertex 92.0366 0 192.666 + vertex 92.2476 126 192.989 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.2476 126 192.989 + vertex 92.0366 0 192.666 + vertex 92.2476 0 192.989 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.2476 126 192.989 + vertex 92.2476 0 192.989 + vertex 92.476 126 193.299 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.476 126 193.299 + vertex 92.2476 0 192.989 + vertex 92.476 0 193.299 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.476 126 193.299 + vertex 92.476 0 193.299 + vertex 92.7213 126 193.597 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.7213 0 193.597 + vertex 92.7213 126 193.597 + vertex 92.476 0 193.299 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.7213 126 193.597 + vertex 92.7213 0 193.597 + vertex 92.9825 126 193.881 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 92.9825 0 193.881 + vertex 92.9825 126 193.881 + vertex 92.7213 0 193.597 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 92.9825 126 193.881 + vertex 92.9825 0 193.881 + vertex 93.259 126 194.15 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.259 0 194.15 + vertex 93.259 126 194.15 + vertex 92.9825 0 193.881 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.259 126 194.15 + vertex 93.259 0 194.15 + vertex 93.5499 126 194.403 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.5499 0 194.403 + vertex 93.5499 126 194.403 + vertex 93.259 0 194.15 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.5499 126 194.403 + vertex 93.5499 0 194.403 + vertex 93.8544 126 194.64 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 93.8544 0 194.64 + vertex 93.8544 126 194.64 + vertex 93.5499 0 194.403 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 93.8544 126 194.64 + vertex 93.8544 0 194.64 + vertex 94.1714 126 194.86 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.1714 0 194.86 + vertex 94.1714 126 194.86 + vertex 93.8544 0 194.64 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.1714 126 194.86 + vertex 94.1714 0 194.86 + vertex 94.5 126 195.062 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.5 0 195.062 + vertex 94.5 126 195.062 + vertex 94.1714 0 194.86 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.5 126 195.062 + vertex 94.5 0 195.062 + vertex 94.8393 126 195.246 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 94.8393 126 195.246 + vertex 94.5 0 195.062 + vertex 94.8393 0 195.246 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 94.8393 126 195.246 + vertex 94.8393 0 195.246 + vertex 95.1881 126 195.41 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.1881 0 195.41 + vertex 95.1881 126 195.41 + vertex 94.8393 0 195.246 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.1881 126 195.41 + vertex 95.1881 0 195.41 + vertex 95.5455 126 195.556 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.5455 0 195.556 + vertex 95.5455 126 195.556 + vertex 95.1881 0 195.41 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.5455 126 195.556 + vertex 95.5455 0 195.556 + vertex 95.9104 126 195.681 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 95.9104 0 195.681 + vertex 95.9104 126 195.681 + vertex 95.5455 0 195.556 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 95.9104 126 195.681 + vertex 95.9104 0 195.681 + vertex 96.2816 126 195.786 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 96.2816 0 195.786 + vertex 96.2816 126 195.786 + vertex 95.9104 0 195.681 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 96.2816 126 195.786 + vertex 96.2816 0 195.786 + vertex 96.658 126 195.87 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 96.658 0 195.87 + vertex 96.658 126 195.87 + vertex 96.2816 0 195.786 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 96.658 126 195.87 + vertex 96.658 0 195.87 + vertex 97.0385 126 195.934 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.0385 0 195.934 + vertex 97.0385 126 195.934 + vertex 96.658 0 195.87 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.0385 126 195.934 + vertex 97.0385 0 195.934 + vertex 97.4219 126 195.976 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.4219 0 195.976 + vertex 97.4219 126 195.976 + vertex 97.0385 0 195.934 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.4219 126 195.976 + vertex 97.4219 0 195.976 + vertex 97.8071 126 195.997 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 97.8071 0 195.997 + vertex 97.8071 126 195.997 + vertex 97.4219 0 195.976 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 97.8071 126 195.997 + vertex 97.8071 0 195.997 + vertex 98.1929 126 195.997 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.1929 0 195.997 + vertex 98.1929 126 195.997 + vertex 97.8071 0 195.997 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.1929 126 195.997 + vertex 98.1929 0 195.997 + vertex 98.5781 126 195.976 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.5781 0 195.976 + vertex 98.5781 126 195.976 + vertex 98.1929 0 195.997 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.5781 126 195.976 + vertex 98.5781 0 195.976 + vertex 98.9615 126 195.934 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 98.9615 0 195.934 + vertex 98.9615 126 195.934 + vertex 98.5781 0 195.976 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 98.9615 126 195.934 + vertex 98.9615 0 195.934 + vertex 99.342 126 195.87 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 99.342 0 195.87 + vertex 99.342 126 195.87 + vertex 98.9615 0 195.934 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 99.342 126 195.87 + vertex 99.342 0 195.87 + vertex 99.7184 126 195.786 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 99.7184 0 195.786 + vertex 99.7184 126 195.786 + vertex 99.342 0 195.87 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 99.7184 126 195.786 + vertex 99.7184 0 195.786 + vertex 100.09 126 195.681 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.09 0 195.681 + vertex 100.09 126 195.681 + vertex 99.7184 0 195.786 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.09 126 195.681 + vertex 100.09 0 195.681 + vertex 100.454 126 195.556 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.454 0 195.556 + vertex 100.454 126 195.556 + vertex 100.09 0 195.681 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.454 126 195.556 + vertex 100.454 0 195.556 + vertex 100.812 126 195.41 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 100.812 0 195.41 + vertex 100.812 126 195.41 + vertex 100.454 0 195.556 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 100.812 126 195.41 + vertex 100.812 0 195.41 + vertex 101.161 126 195.246 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.161 0 195.246 + vertex 101.161 126 195.246 + vertex 100.812 0 195.41 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.161 126 195.246 + vertex 101.161 0 195.246 + vertex 101.5 126 195.062 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.5 126 195.062 + vertex 101.161 0 195.246 + vertex 101.5 0 195.062 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.5 126 195.062 + vertex 101.5 0 195.062 + vertex 101.829 126 194.86 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 101.829 0 194.86 + vertex 101.829 126 194.86 + vertex 101.5 0 195.062 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 101.829 126 194.86 + vertex 101.829 0 194.86 + vertex 102.146 126 194.64 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.146 0 194.64 + vertex 102.146 126 194.64 + vertex 101.829 0 194.86 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.146 126 194.64 + vertex 102.146 0 194.64 + vertex 102.45 126 194.403 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.45 0 194.403 + vertex 102.45 126 194.403 + vertex 102.146 0 194.64 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.45 126 194.403 + vertex 102.45 0 194.403 + vertex 102.741 126 194.15 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 102.741 0 194.15 + vertex 102.741 126 194.15 + vertex 102.45 0 194.403 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 102.741 126 194.15 + vertex 102.741 0 194.15 + vertex 103.017 126 193.881 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.017 0 193.881 + vertex 103.017 126 193.881 + vertex 102.741 0 194.15 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.017 126 193.881 + vertex 103.017 0 193.881 + vertex 103.279 126 193.597 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.279 0 193.597 + vertex 103.279 126 193.597 + vertex 103.017 0 193.881 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.279 126 193.597 + vertex 103.279 0 193.597 + vertex 103.524 126 193.299 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.524 0 193.299 + vertex 103.524 126 193.299 + vertex 103.279 0 193.597 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.524 126 193.299 + vertex 103.524 0 193.299 + vertex 103.752 126 192.989 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.752 126 192.989 + vertex 103.524 0 193.299 + vertex 103.752 0 192.989 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.752 126 192.989 + vertex 103.752 0 192.989 + vertex 103.963 126 192.666 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 103.963 126 192.666 + vertex 103.752 0 192.989 + vertex 103.963 0 192.666 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 103.963 126 192.666 + vertex 103.963 0 192.666 + vertex 104.156 126 192.332 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.156 0 192.332 + vertex 104.156 126 192.332 + vertex 103.963 0 192.666 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.156 126 192.332 + vertex 104.156 0 192.332 + vertex 104.33 126 191.987 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.33 0 191.987 + vertex 104.33 126 191.987 + vertex 104.156 0 192.332 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.33 126 191.987 + vertex 104.33 0 191.987 + vertex 104.485 126 191.634 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.485 126 191.634 + vertex 104.33 0 191.987 + vertex 104.485 0 191.634 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.485 126 191.634 + vertex 104.485 0 191.634 + vertex 104.621 126 191.273 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.621 126 191.273 + vertex 104.485 0 191.634 + vertex 104.621 0 191.273 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.621 126 191.273 + vertex 104.621 0 191.273 + vertex 104.736 126 190.905 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.736 0 190.905 + vertex 104.736 126 190.905 + vertex 104.621 0 191.273 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.736 126 190.905 + vertex 104.736 0 190.905 + vertex 104.831 126 190.531 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.831 126 190.531 + vertex 104.736 0 190.905 + vertex 104.831 0 190.531 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.831 126 190.531 + vertex 104.831 0 190.531 + vertex 104.905 126 190.152 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.905 0 190.152 + vertex 104.905 126 190.152 + vertex 104.831 0 190.531 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.905 126 190.152 + vertex 104.905 0 190.152 + vertex 104.958 126 189.77 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.958 0 189.77 + vertex 104.958 126 189.77 + vertex 104.905 0 190.152 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.958 126 189.77 + vertex 104.958 0 189.77 + vertex 104.989 126 189.386 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 104.989 126 189.386 + vertex 104.958 0 189.77 + vertex 104.989 0 189.386 + endloop + endfacet + facet normal 0.99962 0 -0.0275497 + outer loop + vertex 104.989 126 189.386 + vertex 104.989 0 189.386 + vertex 105 126 189 + endloop + endfacet + facet normal -0.99962 0 0.0275497 + outer loop + vertex 105 0 189 + vertex 105 126 189 + vertex 104.989 0 189.386 + endloop + endfacet + facet normal 0.915772 0 -0.401698 + outer loop + vertex 102.741 126 183.85 + vertex 102.45 126 183.597 + vertex 103.017 126 184.119 + endloop + endfacet + facet normal -0.915772 0.000168664 0.401698 + outer loop + vertex 103.017 126 184.119 + vertex 102.45 126 183.597 + vertex 103.279 126 184.403 + endloop + endfacet + facet normal -0.926493 0 0.376311 + outer loop + vertex 102.45 126 183.597 + vertex 102.146 126 183.36 + vertex 103.279 126 184.403 + endloop + endfacet + facet normal 0.926493 0 -0.376311 + outer loop + vertex 103.279 126 184.403 + vertex 102.146 126 183.36 + vertex 103.524 126 184.701 + endloop + endfacet + facet normal -0.936511 0.00016866 0.350639 + outer loop + vertex 103.524 126 184.701 + vertex 102.146 126 183.36 + vertex 103.752 126 185.011 + endloop + endfacet + facet normal 0.936511 0 -0.350639 + outer loop + vertex 103.752 126 185.011 + vertex 102.146 126 183.36 + vertex 103.963 126 185.334 + endloop + endfacet + facet normal -0.945818 0.000252854 0.324698 + outer loop + vertex 103.963 126 185.334 + vertex 102.146 126 183.36 + vertex 104.156 126 185.668 + endloop + endfacet + facet normal 0.945818 0 -0.324698 + outer loop + vertex 104.156 126 185.668 + vertex 102.146 126 183.36 + vertex 104.33 126 186.013 + endloop + endfacet + facet normal -0.954405 0.000336805 0.298513 + outer loop + vertex 104.33 126 186.013 + vertex 102.146 126 183.36 + vertex 104.485 126 186.366 + endloop + endfacet + facet normal 0.954406 0 -0.298513 + outer loop + vertex 104.485 126 186.366 + vertex 102.146 126 183.36 + vertex 104.621 126 186.727 + endloop + endfacet + facet normal -0.962268 0.000420524 0.272102 + outer loop + vertex 104.621 126 186.727 + vertex 102.146 126 183.36 + vertex 104.736 126 187.095 + endloop + endfacet + facet normal 0.962269 0 -0.272102 + outer loop + vertex 104.736 126 187.095 + vertex 102.146 126 183.36 + vertex 104.831 126 187.469 + endloop + endfacet + facet normal -0.9694 0.000503865 0.245487 + outer loop + vertex 104.831 126 187.469 + vertex 102.146 126 183.36 + vertex 104.905 126 187.848 + endloop + endfacet + facet normal 0.9694 0 -0.245487 + outer loop + vertex 104.905 126 187.848 + vertex 102.146 126 183.36 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.962268 0.000586912 -0.272103 + outer loop + vertex 102.146 126 183.36 + vertex 101.829 126 183.14 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.962268 0 0.272103 + outer loop + vertex 101.829 126 183.14 + vertex 101.5 126 182.938 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.954404 0.000669558 -0.298517 + outer loop + vertex 101.5 126 182.938 + vertex 101.161 126 182.754 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.954404 0 0.298517 + outer loop + vertex 101.161 126 182.754 + vertex 100.812 126 182.59 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.945817 0.000751555 -0.3247 + outer loop + vertex 100.812 126 182.59 + vertex 100.454 126 182.444 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.945817 0 0.324701 + outer loop + vertex 100.454 126 182.444 + vertex 100.09 126 182.319 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.936511 0.00083305 -0.350638 + outer loop + vertex 100.09 126 182.319 + vertex 99.7184 126 182.214 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.936511 0 0.350638 + outer loop + vertex 99.7184 126 182.214 + vertex 99.342 126 182.13 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.926493 0.000913942 -0.37631 + outer loop + vertex 99.342 126 182.13 + vertex 98.9615 126 182.066 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.926494 0 0.37631 + outer loop + vertex 98.9615 126 182.066 + vertex 98.5781 126 182.024 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.915773 0.000994094 -0.401696 + outer loop + vertex 98.5781 126 182.024 + vertex 98.1929 126 182.003 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.915773 0 0.401696 + outer loop + vertex 98.1929 126 182.003 + vertex 97.8071 126 182.003 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.904357 0.00107347 -0.426776 + outer loop + vertex 97.8071 126 182.003 + vertex 97.4219 126 182.024 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.904357 0 0.426776 + outer loop + vertex 97.4219 126 182.024 + vertex 97.0385 126 182.066 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.892253 0.00115216 -0.451534 + outer loop + vertex 97.0385 126 182.066 + vertex 96.658 126 182.13 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.892254 0 0.451534 + outer loop + vertex 96.658 126 182.13 + vertex 96.2816 126 182.214 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.879472 0.00122983 -0.475948 + outer loop + vertex 96.2816 126 182.214 + vertex 95.9104 126 182.319 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.879473 0 0.475949 + outer loop + vertex 95.9104 126 182.319 + vertex 95.5455 126 182.444 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.866024 0.00130657 -0.5 + outer loop + vertex 95.5455 126 182.444 + vertex 95.1881 126 182.59 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.866025 0 0.5 + outer loop + vertex 95.1881 126 182.59 + vertex 94.8393 126 182.754 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.851918 0.00138241 -0.523673 + outer loop + vertex 94.8393 126 182.754 + vertex 94.5 126 182.938 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.851919 0 0.523673 + outer loop + vertex 94.5 126 182.938 + vertex 94.1714 126 183.14 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.837165 0.00145717 -0.546948 + outer loop + vertex 94.1714 126 183.14 + vertex 93.8544 126 183.36 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.837166 0 0.546949 + outer loop + vertex 93.8544 126 183.36 + vertex 93.5499 126 183.597 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.821777 0.00153078 -0.569808 + outer loop + vertex 93.5499 126 183.597 + vertex 93.259 126 183.85 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.821777 0 0.569809 + outer loop + vertex 93.259 126 183.85 + vertex 92.9825 126 184.119 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.805763 0.00160332 -0.592236 + outer loop + vertex 92.9825 126 184.119 + vertex 92.7213 126 184.403 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.805764 0 0.592236 + outer loop + vertex 92.7213 126 184.403 + vertex 92.476 126 184.701 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.789139 0.00167447 -0.614212 + outer loop + vertex 92.476 126 184.701 + vertex 92.2476 126 185.011 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.78914 0 0.614213 + outer loop + vertex 92.2476 126 185.011 + vertex 92.0366 126 185.334 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.771915 0.00174454 -0.635723 + outer loop + vertex 92.0366 126 185.334 + vertex 91.8437 126 185.668 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.771916 0 0.635724 + outer loop + vertex 91.8437 126 185.668 + vertex 91.6695 126 186.013 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.754105 0.00181316 -0.656751 + outer loop + vertex 91.6695 126 186.013 + vertex 91.5145 126 186.366 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.754106 0 0.656752 + outer loop + vertex 91.5145 126 186.366 + vertex 91.3793 126 186.727 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.735722 0.00188048 -0.677281 + outer loop + vertex 91.3793 126 186.727 + vertex 91.2641 126 187.095 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.735723 0 0.677282 + outer loop + vertex 91.2641 126 187.095 + vertex 91.1694 126 187.469 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal 0.716781 0.00194631 -0.697296 + outer loop + vertex 91.1694 126 187.469 + vertex 91.0955 126 187.848 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.716782 0 0.697297 + outer loop + vertex 91.0955 126 187.848 + vertex 91.0425 126 188.23 + vertex 104.958 126 188.23 + endloop + endfacet + facet normal -0.735722 0.00201074 0.67728 + outer loop + vertex 104.958 126 188.23 + vertex 91.0425 126 188.23 + vertex 104.989 126 188.614 + endloop + endfacet + facet normal 0.735724 0 -0.677282 + outer loop + vertex 104.989 126 188.614 + vertex 91.0425 126 188.23 + vertex 105 126 189 + endloop + endfacet + facet normal -0.654901 0.075013 -0.751982 + outer loop + vertex 105 126 189 + vertex 91.0425 126 188.23 + vertex 104.989 126 189.386 + endloop + endfacet + facet normal 0.656751 0 0.754107 + outer loop + vertex 104.989 126 189.386 + vertex 91.0425 126 188.23 + vertex 104.958 126 189.77 + endloop + endfacet + facet normal -0.635723 0.00219439 -0.771914 + outer loop + vertex 104.958 126 189.77 + vertex 91.0425 126 188.23 + vertex 104.905 126 190.152 + endloop + endfacet + facet normal 0.635724 0 0.771916 + outer loop + vertex 104.905 126 190.152 + vertex 91.0425 126 188.23 + vertex 104.831 126 190.531 + endloop + endfacet + facet normal -0.614211 0.00213487 -0.789139 + outer loop + vertex 104.831 126 190.531 + vertex 91.0425 126 188.23 + vertex 104.736 126 190.905 + endloop + endfacet + facet normal 0.614213 0 0.78914 + outer loop + vertex 104.736 126 190.905 + vertex 91.0425 126 188.23 + vertex 104.621 126 191.273 + endloop + endfacet + facet normal -0.592234 0.00207354 -0.805763 + outer loop + vertex 104.621 126 191.273 + vertex 91.0425 126 188.23 + vertex 104.485 126 191.634 + endloop + endfacet + facet normal 0.592236 0 0.805765 + outer loop + vertex 104.485 126 191.634 + vertex 91.0425 126 188.23 + vertex 104.33 126 191.987 + endloop + endfacet + facet normal -0.569807 0.00201075 -0.821776 + outer loop + vertex 104.33 126 191.987 + vertex 91.0425 126 188.23 + vertex 104.156 126 192.332 + endloop + endfacet + facet normal 0.569808 0 0.821778 + outer loop + vertex 104.156 126 192.332 + vertex 91.0425 126 188.23 + vertex 103.963 126 192.666 + endloop + endfacet + facet normal -0.546946 0.00194636 -0.837165 + outer loop + vertex 103.963 126 192.666 + vertex 91.0425 126 188.23 + vertex 103.752 126 192.989 + endloop + endfacet + facet normal 0.546947 0.00181322 0.837166 + outer loop + vertex 91.0425 126 188.23 + vertex 91.0106 126 188.614 + vertex 103.752 126 192.989 + endloop + endfacet + facet normal -0.569808 0 -0.821778 + outer loop + vertex 91.0106 126 188.614 + vertex 91 126 189 + vertex 103.752 126 192.989 + endloop + endfacet + facet normal 0.569807 0.0017445 0.821777 + outer loop + vertex 91 126 189 + vertex 91.0106 126 189.386 + vertex 103.752 126 192.989 + endloop + endfacet + facet normal 0.592235 0 0.805765 + outer loop + vertex 103.752 126 192.989 + vertex 91.0106 126 189.386 + vertex 103.524 126 193.299 + endloop + endfacet + facet normal -0.592234 0.00167452 -0.805764 + outer loop + vertex 103.524 126 193.299 + vertex 91.0106 126 189.386 + vertex 103.279 126 193.597 + endloop + endfacet + facet normal 0.569808 0 0.821778 + outer loop + vertex 103.279 126 193.597 + vertex 91.0106 126 189.386 + vertex 103.017 126 193.881 + endloop + endfacet + facet normal -0.569807 0.00160328 -0.821777 + outer loop + vertex 103.017 126 193.881 + vertex 91.0106 126 189.386 + vertex 102.741 126 194.15 + endloop + endfacet + facet normal 0.546948 0 0.837167 + outer loop + vertex 102.741 126 194.15 + vertex 91.0106 126 189.386 + vertex 102.45 126 194.403 + endloop + endfacet + facet normal -0.546947 0.00153079 -0.837166 + outer loop + vertex 102.45 126 194.403 + vertex 91.0106 126 189.386 + vertex 102.146 126 194.64 + endloop + endfacet + facet normal 0.523673 0 0.85192 + outer loop + vertex 102.146 126 194.64 + vertex 91.0106 126 189.386 + vertex 101.829 126 194.86 + endloop + endfacet + facet normal -0.523672 0.00145711 -0.851919 + outer loop + vertex 101.829 126 194.86 + vertex 91.0106 126 189.386 + vertex 101.5 126 195.062 + endloop + endfacet + facet normal 0.5 0 0.866025 + outer loop + vertex 101.5 126 195.062 + vertex 91.0106 126 189.386 + vertex 101.161 126 195.246 + endloop + endfacet + facet normal -0.5 0.00138245 -0.866024 + outer loop + vertex 101.161 126 195.246 + vertex 91.0106 126 189.386 + vertex 100.812 126 195.41 + endloop + endfacet + facet normal 0.475947 0 0.879474 + outer loop + vertex 100.812 126 195.41 + vertex 91.0106 126 189.386 + vertex 100.454 126 195.556 + endloop + endfacet + facet normal -0.475947 0.00130656 -0.879473 + outer loop + vertex 100.454 126 195.556 + vertex 91.0106 126 189.386 + vertex 100.09 126 195.681 + endloop + endfacet + facet normal 0.451534 0 0.892254 + outer loop + vertex 100.09 126 195.681 + vertex 91.0106 126 189.386 + vertex 99.7184 126 195.786 + endloop + endfacet + facet normal -0.451534 0.00122982 -0.892253 + outer loop + vertex 99.7184 126 195.786 + vertex 91.0106 126 189.386 + vertex 99.342 126 195.87 + endloop + endfacet + facet normal 0.426777 0 0.904357 + outer loop + vertex 99.342 126 195.87 + vertex 91.0106 126 189.386 + vertex 98.9615 126 195.934 + endloop + endfacet + facet normal -0.426777 0.0011521 -0.904356 + outer loop + vertex 98.9615 126 195.934 + vertex 91.0106 126 189.386 + vertex 98.5781 126 195.976 + endloop + endfacet + facet normal 0.401696 0 0.915773 + outer loop + vertex 98.5781 126 195.976 + vertex 91.0106 126 189.386 + vertex 98.1929 126 195.997 + endloop + endfacet + facet normal -0.401696 0.00107357 -0.915772 + outer loop + vertex 98.1929 126 195.997 + vertex 91.0106 126 189.386 + vertex 97.8071 126 195.997 + endloop + endfacet + facet normal 0.376309 0 0.926494 + outer loop + vertex 97.8071 126 195.997 + vertex 91.0106 126 189.386 + vertex 97.4219 126 195.976 + endloop + endfacet + facet normal -0.376309 0.000994031 -0.926494 + outer loop + vertex 97.4219 126 195.976 + vertex 91.0106 126 189.386 + vertex 97.0385 126 195.934 + endloop + endfacet + facet normal 0.350639 0 0.936511 + outer loop + vertex 97.0385 126 195.934 + vertex 91.0106 126 189.386 + vertex 96.658 126 195.87 + endloop + endfacet + facet normal -0.350639 0.000913918 -0.93651 + outer loop + vertex 96.658 126 195.87 + vertex 91.0106 126 189.386 + vertex 96.2816 126 195.786 + endloop + endfacet + facet normal 0.324701 0 0.945817 + outer loop + vertex 96.2816 126 195.786 + vertex 91.0106 126 189.386 + vertex 95.9104 126 195.681 + endloop + endfacet + facet normal -0.324701 0.00083313 -0.945816 + outer loop + vertex 95.9104 126 195.681 + vertex 91.0106 126 189.386 + vertex 95.5455 126 195.556 + endloop + endfacet + facet normal 0.298515 0 0.954405 + outer loop + vertex 95.5455 126 195.556 + vertex 91.0106 126 189.386 + vertex 95.1881 126 195.41 + endloop + endfacet + facet normal -0.298515 0.000751569 -0.954405 + outer loop + vertex 95.1881 126 195.41 + vertex 91.0106 126 189.386 + vertex 94.8393 126 195.246 + endloop + endfacet + facet normal 0.272104 0 0.962268 + outer loop + vertex 94.8393 126 195.246 + vertex 91.0106 126 189.386 + vertex 94.5 126 195.062 + endloop + endfacet + facet normal -0.272104 0.000669456 -0.962268 + outer loop + vertex 94.5 126 195.062 + vertex 91.0106 126 189.386 + vertex 94.1714 126 194.86 + endloop + endfacet + facet normal 0.245488 0 0.9694 + outer loop + vertex 94.1714 126 194.86 + vertex 91.0106 126 189.386 + vertex 93.8544 126 194.64 + endloop + endfacet + facet normal -0.245488 0.000586992 -0.969399 + outer loop + vertex 93.8544 126 194.64 + vertex 91.0106 126 189.386 + vertex 93.5499 126 194.403 + endloop + endfacet + facet normal 0.218681 0 0.975796 + outer loop + vertex 93.5499 126 194.403 + vertex 91.0106 126 189.386 + vertex 93.259 126 194.15 + endloop + endfacet + facet normal -0.218681 0.000503852 -0.975796 + outer loop + vertex 93.259 126 194.15 + vertex 91.0106 126 189.386 + vertex 92.9825 126 193.881 + endloop + endfacet + facet normal 0.191714 0 0.981451 + outer loop + vertex 92.9825 126 193.881 + vertex 91.0106 126 189.386 + vertex 92.7213 126 193.597 + endloop + endfacet + facet normal -0.191714 0.000420542 -0.981451 + outer loop + vertex 92.7213 126 193.597 + vertex 91.0106 126 189.386 + vertex 92.476 126 193.299 + endloop + endfacet + facet normal 0.164597 0 0.986361 + outer loop + vertex 92.476 126 193.299 + vertex 91.0106 126 189.386 + vertex 92.2476 126 192.989 + endloop + endfacet + facet normal -0.164597 0.000336776 -0.986361 + outer loop + vertex 92.2476 126 192.989 + vertex 91.0106 126 189.386 + vertex 92.0366 126 192.666 + endloop + endfacet + facet normal 0.137358 0 0.990521 + outer loop + vertex 92.0366 126 192.666 + vertex 91.0106 126 189.386 + vertex 91.8437 126 192.332 + endloop + endfacet + facet normal -0.137358 0.000252864 -0.990521 + outer loop + vertex 91.8437 126 192.332 + vertex 91.0106 126 189.386 + vertex 91.6695 126 191.987 + endloop + endfacet + facet normal 0.110009 0 0.993931 + outer loop + vertex 91.6695 126 191.987 + vertex 91.0106 126 189.386 + vertex 91.5145 126 191.634 + endloop + endfacet + facet normal -0.110009 0.000168636 -0.993931 + outer loop + vertex 91.5145 126 191.634 + vertex 91.0106 126 189.386 + vertex 91.3793 126 191.273 + endloop + endfacet + facet normal 0.0825829 0 0.996584 + outer loop + vertex 91.3793 126 191.273 + vertex 91.0106 126 189.386 + vertex 91.2641 126 190.905 + endloop + endfacet + facet normal -0.0825829 8.43349e-05 -0.996584 + outer loop + vertex 91.2641 126 190.905 + vertex 91.0106 126 189.386 + vertex 91.1694 126 190.531 + endloop + endfacet + facet normal 0.0550998 0 0.998481 + outer loop + vertex 91.1694 126 190.531 + vertex 91.0106 126 189.386 + vertex 91.0955 126 190.152 + endloop + endfacet + facet normal -0.0550998 0 -0.998481 + outer loop + vertex 91.0425 126 189.77 + vertex 91.0955 126 190.152 + vertex 91.0106 126 189.386 + endloop + endfacet + facet normal -0.936511 0 0.350638 + outer loop + vertex 102.741 0 194.15 + vertex 102.45 0 194.403 + vertex 103.017 0 193.881 + endloop + endfacet + facet normal 0.936511 0 -0.350638 + outer loop + vertex 103.017 0 193.881 + vertex 102.45 0 194.403 + vertex 103.279 0 193.597 + endloop + endfacet + facet normal 0.926493 8.43648e-05 -0.376311 + outer loop + vertex 102.45 0 194.403 + vertex 102.146 0 194.64 + vertex 103.279 0 193.597 + endloop + endfacet + facet normal -0.926493 0.000252832 0.376311 + outer loop + vertex 103.279 0 193.597 + vertex 102.146 0 194.64 + vertex 103.524 0 193.299 + endloop + endfacet + facet normal 0.936511 0 -0.350639 + outer loop + vertex 103.524 0 193.299 + vertex 102.146 0 194.64 + vertex 103.752 0 192.989 + endloop + endfacet + facet normal -0.936511 0.00033684 0.350639 + outer loop + vertex 103.752 0 192.989 + vertex 102.146 0 194.64 + vertex 103.963 0 192.666 + endloop + endfacet + facet normal 0.945818 0 -0.324698 + outer loop + vertex 103.963 0 192.666 + vertex 102.146 0 194.64 + vertex 104.156 0 192.332 + endloop + endfacet + facet normal -0.945818 0.000420525 0.324698 + outer loop + vertex 104.156 0 192.332 + vertex 102.146 0 194.64 + vertex 104.33 0 191.987 + endloop + endfacet + facet normal 0.954406 0 -0.298513 + outer loop + vertex 104.33 0 191.987 + vertex 102.146 0 194.64 + vertex 104.485 0 191.634 + endloop + endfacet + facet normal -0.954405 0.000503926 0.298513 + outer loop + vertex 104.485 0 191.634 + vertex 102.146 0 194.64 + vertex 104.621 0 191.273 + endloop + endfacet + facet normal 0.962269 0 -0.272102 + outer loop + vertex 104.621 0 191.273 + vertex 102.146 0 194.64 + vertex 104.736 0 190.905 + endloop + endfacet + facet normal -0.962268 0.000586871 0.272102 + outer loop + vertex 104.736 0 190.905 + vertex 102.146 0 194.64 + vertex 104.831 0 190.531 + endloop + endfacet + facet normal 0.9694 0 -0.245487 + outer loop + vertex 104.831 0 190.531 + vertex 102.146 0 194.64 + vertex 104.905 0 190.152 + endloop + endfacet + facet normal -0.9694 0.00066954 0.245487 + outer loop + vertex 104.905 0 190.152 + vertex 102.146 0 194.64 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.975796 0 0.218681 + outer loop + vertex 102.146 0 194.64 + vertex 101.829 0 194.86 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.975796 0.000751546 -0.218681 + outer loop + vertex 101.829 0 194.86 + vertex 101.5 0 195.062 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.9694 0 0.245485 + outer loop + vertex 101.5 0 195.062 + vertex 101.161 0 195.246 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.9694 0.000833124 -0.245485 + outer loop + vertex 101.161 0 195.246 + vertex 100.812 0 195.41 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.962268 0 0.272104 + outer loop + vertex 100.812 0 195.41 + vertex 100.454 0 195.556 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.962267 0.000913913 -0.272104 + outer loop + vertex 100.454 0 195.556 + vertex 100.09 0 195.681 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.954405 0 0.298515 + outer loop + vertex 100.09 0 195.681 + vertex 99.7184 0 195.786 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.954404 0.000994073 -0.298515 + outer loop + vertex 99.7184 0 195.786 + vertex 99.342 0 195.87 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.945817 0 0.324699 + outer loop + vertex 99.342 0 195.87 + vertex 98.9615 0 195.934 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.945817 0.00107352 -0.324699 + outer loop + vertex 98.9615 0 195.934 + vertex 98.5781 0 195.976 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.936511 0 0.350638 + outer loop + vertex 98.5781 0 195.976 + vertex 98.1929 0 195.997 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.936511 0.0011521 -0.350637 + outer loop + vertex 98.1929 0 195.997 + vertex 97.8071 0 195.997 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.926494 0 0.376309 + outer loop + vertex 97.8071 0 195.997 + vertex 97.4219 0 195.976 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.926493 0.00122979 -0.376309 + outer loop + vertex 97.4219 0 195.976 + vertex 97.0385 0 195.934 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.915774 0 0.401695 + outer loop + vertex 97.0385 0 195.934 + vertex 96.658 0 195.87 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.915773 0.00130667 -0.401694 + outer loop + vertex 96.658 0 195.87 + vertex 96.2816 0 195.786 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.904357 0 0.426777 + outer loop + vertex 96.2816 0 195.786 + vertex 95.9104 0 195.681 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.904356 0.00138242 -0.426776 + outer loop + vertex 95.9104 0 195.681 + vertex 95.5455 0 195.556 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.892254 0 0.451534 + outer loop + vertex 95.5455 0 195.556 + vertex 95.1881 0 195.41 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.892253 0.00145712 -0.451533 + outer loop + vertex 95.1881 0 195.41 + vertex 94.8393 0 195.246 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.879474 0 0.475947 + outer loop + vertex 94.8393 0 195.246 + vertex 94.5 0 195.062 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.879473 0.00153079 -0.475947 + outer loop + vertex 94.5 0 195.062 + vertex 94.1714 0 194.86 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.866025 0 0.5 + outer loop + vertex 94.1714 0 194.86 + vertex 93.8544 0 194.64 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.866024 0.00160329 -0.499999 + outer loop + vertex 93.8544 0 194.64 + vertex 93.5499 0 194.403 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.851919 0 0.523673 + outer loop + vertex 93.5499 0 194.403 + vertex 93.259 0 194.15 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.851918 0.00167452 -0.523672 + outer loop + vertex 93.259 0 194.15 + vertex 92.9825 0 193.881 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.837166 0 0.546948 + outer loop + vertex 92.9825 0 193.881 + vertex 92.7213 0 193.597 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.837165 0.00174457 -0.546947 + outer loop + vertex 92.7213 0 193.597 + vertex 92.476 0 193.299 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.821777 0 0.569809 + outer loop + vertex 92.476 0 193.299 + vertex 92.2476 0 192.989 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.821776 0.00181312 -0.569808 + outer loop + vertex 92.2476 0 192.989 + vertex 92.0366 0 192.666 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.805765 0 0.592235 + outer loop + vertex 92.0366 0 192.666 + vertex 91.8437 0 192.332 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.805764 0.00188049 -0.592234 + outer loop + vertex 91.8437 0 192.332 + vertex 91.6695 0 191.987 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.78914 0 0.614213 + outer loop + vertex 91.6695 0 191.987 + vertex 91.5145 0 191.634 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.789139 0.00194631 -0.614212 + outer loop + vertex 91.5145 0 191.634 + vertex 91.3793 0 191.273 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.771917 0 0.635724 + outer loop + vertex 91.3793 0 191.273 + vertex 91.2641 0 190.905 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.771915 0.00201072 -0.635722 + outer loop + vertex 91.2641 0 190.905 + vertex 91.1694 0 190.531 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal -0.754106 0 0.656752 + outer loop + vertex 91.1694 0 190.531 + vertex 91.0955 0 190.152 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.754105 0.00207354 -0.656751 + outer loop + vertex 91.0955 0 190.152 + vertex 91.0425 0 189.77 + vertex 104.958 0 189.77 + endloop + endfacet + facet normal 0.735724 0 -0.677282 + outer loop + vertex 104.958 0 189.77 + vertex 91.0425 0 189.77 + vertex 104.989 0 189.386 + endloop + endfacet + facet normal -0.733403 0.0793611 0.675145 + outer loop + vertex 104.989 0 189.386 + vertex 91.0425 0 189.77 + vertex 105 0 189 + endloop + endfacet + facet normal 0.656751 0 0.754107 + outer loop + vertex 105 0 189 + vertex 91.0425 0 189.77 + vertex 104.989 0 188.614 + endloop + endfacet + facet normal -0.65675 0.00213473 -0.754105 + outer loop + vertex 104.989 0 188.614 + vertex 91.0425 0 189.77 + vertex 104.958 0 188.23 + endloop + endfacet + facet normal 0.635724 0 0.771916 + outer loop + vertex 104.958 0 188.23 + vertex 91.0425 0 189.77 + vertex 104.905 0 187.848 + endloop + endfacet + facet normal -0.635723 0.00207359 -0.771915 + outer loop + vertex 104.905 0 187.848 + vertex 91.0425 0 189.77 + vertex 104.831 0 187.469 + endloop + endfacet + facet normal 0.614213 0 0.78914 + outer loop + vertex 104.831 0 187.469 + vertex 91.0425 0 189.77 + vertex 104.736 0 187.095 + endloop + endfacet + facet normal -0.614212 0.00201069 -0.789139 + outer loop + vertex 104.736 0 187.095 + vertex 91.0425 0 189.77 + vertex 104.621 0 186.727 + endloop + endfacet + facet normal 0.592236 0 0.805765 + outer loop + vertex 104.621 0 186.727 + vertex 91.0425 0 189.77 + vertex 104.485 0 186.366 + endloop + endfacet + facet normal -0.592234 0.00194637 -0.805763 + outer loop + vertex 104.485 0 186.366 + vertex 91.0425 0 189.77 + vertex 104.33 0 186.013 + endloop + endfacet + facet normal 0.569808 0 0.821778 + outer loop + vertex 104.33 0 186.013 + vertex 91.0425 0 189.77 + vertex 104.156 0 185.668 + endloop + endfacet + facet normal -0.569807 0.0018805 -0.821777 + outer loop + vertex 104.156 0 185.668 + vertex 91.0425 0 189.77 + vertex 103.963 0 185.334 + endloop + endfacet + facet normal 0.546947 0 0.837167 + outer loop + vertex 103.963 0 185.334 + vertex 91.0425 0 189.77 + vertex 103.752 0 185.011 + endloop + endfacet + facet normal -0.546947 0 -0.837167 + outer loop + vertex 91.0425 0 189.77 + vertex 91.0106 0 189.386 + vertex 103.752 0 185.011 + endloop + endfacet + facet normal 0.569807 0.00188051 0.821777 + outer loop + vertex 91.0106 0 189.386 + vertex 91 0 189 + vertex 103.752 0 185.011 + endloop + endfacet + facet normal -0.569808 0 -0.821778 + outer loop + vertex 91 0 189 + vertex 91.0106 0 188.614 + vertex 103.752 0 185.011 + endloop + endfacet + facet normal -0.546947 0.00181317 -0.837165 + outer loop + vertex 103.752 0 185.011 + vertex 91.0106 0 188.614 + vertex 103.524 0 184.701 + endloop + endfacet + facet normal 0.546948 0 0.837167 + outer loop + vertex 103.524 0 184.701 + vertex 91.0106 0 188.614 + vertex 103.279 0 184.403 + endloop + endfacet + facet normal -0.523672 0.00174452 -0.851918 + outer loop + vertex 103.279 0 184.403 + vertex 91.0106 0 188.614 + vertex 103.017 0 184.119 + endloop + endfacet + facet normal 0.523673 0 0.85192 + outer loop + vertex 103.017 0 184.119 + vertex 91.0106 0 188.614 + vertex 102.741 0 183.85 + endloop + endfacet + facet normal -0.499999 0.00167452 -0.866024 + outer loop + vertex 102.741 0 183.85 + vertex 91.0106 0 188.614 + vertex 102.45 0 183.597 + endloop + endfacet + facet normal 0.5 0 0.866025 + outer loop + vertex 102.45 0 183.597 + vertex 91.0106 0 188.614 + vertex 102.146 0 183.36 + endloop + endfacet + facet normal -0.475947 0.00160323 -0.879472 + outer loop + vertex 102.146 0 183.36 + vertex 91.0106 0 188.614 + vertex 101.829 0 183.14 + endloop + endfacet + facet normal 0.475948 0 0.879473 + outer loop + vertex 101.829 0 183.14 + vertex 91.0106 0 188.614 + vertex 101.5 0 182.938 + endloop + endfacet + facet normal -0.451533 0.00153084 -0.892253 + outer loop + vertex 101.5 0 182.938 + vertex 91.0106 0 188.614 + vertex 101.161 0 182.754 + endloop + endfacet + facet normal 0.451533 0 0.892254 + outer loop + vertex 101.161 0 182.754 + vertex 91.0106 0 188.614 + vertex 100.812 0 182.59 + endloop + endfacet + facet normal -0.426777 0.0014571 -0.904356 + outer loop + vertex 100.812 0 182.59 + vertex 91.0106 0 188.614 + vertex 100.454 0 182.444 + endloop + endfacet + facet normal 0.426777 0 0.904357 + outer loop + vertex 100.454 0 182.444 + vertex 91.0106 0 188.614 + vertex 100.09 0 182.319 + endloop + endfacet + facet normal -0.401696 0.0013824 -0.915772 + outer loop + vertex 100.09 0 182.319 + vertex 91.0106 0 188.614 + vertex 99.7184 0 182.214 + endloop + endfacet + facet normal 0.401696 0 0.915773 + outer loop + vertex 99.7184 0 182.214 + vertex 91.0106 0 188.614 + vertex 99.342 0 182.13 + endloop + endfacet + facet normal -0.37631 0.00130661 -0.926493 + outer loop + vertex 99.342 0 182.13 + vertex 91.0106 0 188.614 + vertex 98.9615 0 182.066 + endloop + endfacet + facet normal 0.376311 0 0.926494 + outer loop + vertex 98.9615 0 182.066 + vertex 91.0106 0 188.614 + vertex 98.5781 0 182.024 + endloop + endfacet + facet normal -0.350637 0.00122989 -0.936511 + outer loop + vertex 98.5781 0 182.024 + vertex 91.0106 0 188.614 + vertex 98.1929 0 182.003 + endloop + endfacet + facet normal 0.350637 0 0.936511 + outer loop + vertex 98.1929 0 182.003 + vertex 91.0106 0 188.614 + vertex 97.8071 0 182.003 + endloop + endfacet + facet normal -0.324701 0.00115203 -0.945816 + outer loop + vertex 97.8071 0 182.003 + vertex 91.0106 0 188.614 + vertex 97.4219 0 182.024 + endloop + endfacet + facet normal 0.324701 0 0.945817 + outer loop + vertex 97.4219 0 182.024 + vertex 91.0106 0 188.614 + vertex 97.0385 0 182.066 + endloop + endfacet + facet normal -0.298516 0.0010735 -0.954404 + outer loop + vertex 97.0385 0 182.066 + vertex 91.0106 0 188.614 + vertex 96.658 0 182.13 + endloop + endfacet + facet normal 0.298517 0 0.954404 + outer loop + vertex 96.658 0 182.13 + vertex 91.0106 0 188.614 + vertex 96.2816 0 182.214 + endloop + endfacet + facet normal -0.272103 0.000994163 -0.962267 + outer loop + vertex 96.2816 0 182.214 + vertex 91.0106 0 188.614 + vertex 95.9104 0 182.319 + endloop + endfacet + facet normal 0.272104 0 0.962268 + outer loop + vertex 95.9104 0 182.319 + vertex 91.0106 0 188.614 + vertex 95.5455 0 182.444 + endloop + endfacet + facet normal -0.245486 0.000913921 -0.9694 + outer loop + vertex 95.5455 0 182.444 + vertex 91.0106 0 188.614 + vertex 95.1881 0 182.59 + endloop + endfacet + facet normal 0.245486 0 0.9694 + outer loop + vertex 95.1881 0 182.59 + vertex 91.0106 0 188.614 + vertex 94.8393 0 182.754 + endloop + endfacet + facet normal -0.218684 0.000833004 -0.975795 + outer loop + vertex 94.8393 0 182.754 + vertex 91.0106 0 188.614 + vertex 94.5 0 182.938 + endloop + endfacet + facet normal 0.218684 0 0.975796 + outer loop + vertex 94.5 0 182.938 + vertex 91.0106 0 188.614 + vertex 94.1714 0 183.14 + endloop + endfacet + facet normal -0.191711 0.000751638 -0.981451 + outer loop + vertex 94.1714 0 183.14 + vertex 91.0106 0 188.614 + vertex 93.8544 0 183.36 + endloop + endfacet + facet normal 0.191711 0 0.981451 + outer loop + vertex 93.8544 0 183.36 + vertex 91.0106 0 188.614 + vertex 93.5499 0 183.597 + endloop + endfacet + facet normal -0.164598 0.000669433 -0.98636 + outer loop + vertex 93.5499 0 183.597 + vertex 91.0106 0 188.614 + vertex 93.259 0 183.85 + endloop + endfacet + facet normal 0.164598 0 0.986361 + outer loop + vertex 93.259 0 183.85 + vertex 91.0106 0 188.614 + vertex 92.9825 0 184.119 + endloop + endfacet + facet normal -0.137356 0.000586964 -0.990522 + outer loop + vertex 92.9825 0 184.119 + vertex 91.0106 0 188.614 + vertex 92.7213 0 184.403 + endloop + endfacet + facet normal 0.137356 0 0.990522 + outer loop + vertex 92.7213 0 184.403 + vertex 91.0106 0 188.614 + vertex 92.476 0 184.701 + endloop + endfacet + facet normal -0.110012 0.000503894 -0.99393 + outer loop + vertex 92.476 0 184.701 + vertex 91.0106 0 188.614 + vertex 92.2476 0 185.011 + endloop + endfacet + facet normal 0.110012 0 0.99393 + outer loop + vertex 92.2476 0 185.011 + vertex 91.0106 0 188.614 + vertex 92.0366 0 185.334 + endloop + endfacet + facet normal -0.0825797 0.000420575 -0.996584 + outer loop + vertex 92.0366 0 185.334 + vertex 91.0106 0 188.614 + vertex 91.8437 0 185.668 + endloop + endfacet + facet normal 0.0825797 0 0.996584 + outer loop + vertex 91.8437 0 185.668 + vertex 91.0106 0 188.614 + vertex 91.6695 0 186.013 + endloop + endfacet + facet normal -0.0550903 0.000336774 -0.998481 + outer loop + vertex 91.6695 0 186.013 + vertex 91.0106 0 188.614 + vertex 91.5145 0 186.366 + endloop + endfacet + facet normal 0.0550903 0 0.998481 + outer loop + vertex 91.5145 0 186.366 + vertex 91.0106 0 188.614 + vertex 91.3793 0 186.727 + endloop + endfacet + facet normal -0.0275606 0.000252789 -0.99962 + outer loop + vertex 91.3793 0 186.727 + vertex 91.0106 0 188.614 + vertex 91.2641 0 187.095 + endloop + endfacet + facet normal 0.0275606 0 0.99962 + outer loop + vertex 91.2641 0 187.095 + vertex 91.0106 0 188.614 + vertex 91.1694 0 187.469 + endloop + endfacet + facet normal 0 0.000168694 -1 + outer loop + vertex 91.1694 0 187.469 + vertex 91.0106 0 188.614 + vertex 91.0955 0 187.848 + endloop + endfacet + facet normal 0 0 1 + outer loop + vertex 91.0425 0 188.23 + vertex 91.0955 0 187.848 + vertex 91.0106 0 188.614 + endloop + endfacet +endsolid diff --git a/apps/gpu/LBM/ActuatorLine/configActuatorLine.txt b/apps/gpu/LBM/ActuatorLine/configActuatorLine.txt index 5799f24716777295b2f835ab00561ff767ba87b9..c7b9c51788f24fc7452a3a8bf16acffe3a6e88ff 100644 --- a/apps/gpu/LBM/ActuatorLine/configActuatorLine.txt +++ b/apps/gpu/LBM/ActuatorLine/configActuatorLine.txt @@ -1,31 +1,31 @@ ################################################## #informations for Writing ################################################## -Path = . +Path = /work/y0078217/Results/ActuatorLine105 ################################################## #informations for reading ################################################## GridPath=. ################################################## ReferenceDiameter=126 -NodesPerDiameter=32 +NodesPerDiameter=105 Velocity=9 ################################################## -tStartOut=100 -tOut=100 -tEnd=1000 +tStartOut=0 +tOut=1 +tEnd=1 ################################################## -tStartTmpAveraging=100 -tStartAveraging=100 -tAveraging=100 -tTmpAveraging=100 -tStartOutProbe=100 -tOutProbe=100 +# tStartTmpAveraging=100 +# tStartAveraging=100 +# tAveraging=100 +# tTmpAveraging=100 +# tStartOutProbe=100 +# tOutProbe=100 ################################################## -#TurbulenceModel = QR -#SGSconstant = 0.3333333 +TurbulenceModel = None +# SGSconstant = 0.3333333 # #QuadricLimiterP = 100000.0 #QuadricLimiterM = 100000.0 diff --git a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp index 4ef9786cdd9130d4da7c459b4f83b468cc0d581a..ed2bd3a9c5be2c7847799ad0a207f8900230eba7 100644 --- a/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp +++ b/apps/gpu/LBM/DrivenCavity/DrivenCavity.cpp @@ -85,7 +85,7 @@ int main() const real L = 1.0; const real Re = 1000.0; const real velocity = 1.0; - const real dt = (real)0.5e-3; + const real velocityLB = 0.05; // LB units const uint nx = 64; const uint timeStepOut = 1000; @@ -112,7 +112,7 @@ int main() // create grid ////////////////////////////////////////////////////////////////////////// - real dx = L / real(nx); + const real dx = L / real(nx); gridBuilder->addCoarseGrid(-0.5 * L, -0.5 * L, -0.5 * L, 0.5 * L, 0.5 * L, 0.5 * L, dx); @@ -128,7 +128,7 @@ int main() // compute parameters in lattice units ////////////////////////////////////////////////////////////////////////// - const real velocityLB = velocity * dt / dx; // LB units + const real dt = velocityLB / velocity * dx; const real vxLB = velocityLB / sqrt(2.0); // LB units const real vyLB = velocityLB / sqrt(2.0); // LB units diff --git a/src/basics/tests/testUtilities.h b/src/basics/tests/testUtilities.h index c70d9cc5c11633ded6b696d92692e3d4edf8d2ca..57606edc130b0471b957202420cb12859a9cde84 100644 --- a/src/basics/tests/testUtilities.h +++ b/src/basics/tests/testUtilities.h @@ -1,6 +1,8 @@ #ifndef TESTUTILITIES_H #define TESTUTILITIES_H +#include <gmock/gmock.h> + inline auto RealEq = [](auto value) { #ifdef VF_DOUBLE_ACCURACY return testing::DoubleEq(value); diff --git a/src/gpu/GridGenerator/grid/Grid.h b/src/gpu/GridGenerator/grid/Grid.h index 85b19bedd470c9856954a1ca20fb446c3d875da2..bc46e6b23fdcd75a65202afd34865e3d08fe3754 100644 --- a/src/gpu/GridGenerator/grid/Grid.h +++ b/src/gpu/GridGenerator/grid/Grid.h @@ -175,7 +175,6 @@ public: virtual uint getNumberOfFluidNodes() const = 0; virtual void getFluidNodeIndices(uint *fluidNodeIndices) const = 0; - virtual void findFluidNodeIndicesBorder() = 0; virtual uint getNumberOfFluidNodesBorder() const = 0; virtual void getFluidNodeIndicesBorder(uint *fluidNodeIndicesBorder) const = 0; diff --git a/src/gpu/GridGenerator/grid/GridImp.h b/src/gpu/GridGenerator/grid/GridImp.h index 8283bf569e266b84f020334a306d93756b01c394..d518868b6bf57d094b33596921d66fa7f87eb328 100644 --- a/src/gpu/GridGenerator/grid/GridImp.h +++ b/src/gpu/GridGenerator/grid/GridImp.h @@ -343,6 +343,8 @@ private: void allocateQs(); + void findFluidNodeIndicesBorder(); + public: void findCommunicationIndices(int direction, SPtr<BoundingBox> subDomainBox, LbmOrGks lbmOrGks) override; void findCommunicationIndex(uint index, real coordinate, real limit, int direction); @@ -359,7 +361,6 @@ public: void repairCommunicationIndices(int direction) override; void findFluidNodeIndices(bool splitDomain) override; - void findFluidNodeIndicesBorder() override; uint getNumberOfFluidNodes() const override; void getFluidNodeIndices(uint *fluidNodeIndices) const override; diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp index 4136614dfbfc9e0d2fc1bf7f4b01624f94eabb6f..26632f18dec369f9b437d0c4549f15a97f15d2d6 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp +++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.cpp @@ -13,8 +13,13 @@ #include "CollisionStrategy.h" #include "RefinementStrategy.h" +#include "Output/Timer.h" + void UpdateGrid27::updateGrid(int level, unsigned int t) { + + timer->startTimer(); + ////////////////////////////////////////////////////////////////////////// if (level != para->getFine()) { @@ -22,18 +27,27 @@ void UpdateGrid27::updateGrid(int level, unsigned int t) updateGrid(level + 1, t); } + ////////////////////////////////////////////////////////////////////////// + std::cout << "updateGrid: level = " << level << ", t = " << t << std::endl; + interactWithProbes(level, t); + std::cout << " interactWithProbes, " << timer->startStopGetElapsed() << std::endl; ////////////////////////////////////////////////////////////////////////// collision(this, para.get(), level, t); + std::cout << " collision, " << timer->startStopGetElapsed() << std::endl; + ////////////////////////////////////////////////////////////////////////// postCollisionBC(level, t); + std::cout << " postCollisionBC, " << timer->startStopGetElapsed() << std::endl; + + ////////////////////////////////////////////////////////////////////////// swapBetweenEvenAndOddTimestep(level); @@ -45,20 +59,35 @@ void UpdateGrid27::updateGrid(int level, unsigned int t) calcTurbulentViscosity(level); + + std::cout << " calcTurbulentViscosity, " << timer->startStopGetElapsed() << std::endl; + ////////////////////////////////////////////////////////////////////////// this->preCollisionBC(level, t); + std::cout << " preCollisionBC, " << timer->startStopGetElapsed() << std::endl; + + ////////////////////////////////////////////////////////////////////////// if( level != para->getFine() ) { refinement(this, para.get(), level); } + + std::cout << " refinement, " << timer->startStopGetElapsed() << std::endl; + ////////////////////////////////////////////////////////////////////////// interactWithActuators(level, t); + std::cout << " interactWithActuators, " << timer->startStopGetElapsed() << std::endl; + std::cout << "total time, " << timer->getTotalElapsedTime() << std::endl; + timer->resetTimer(); + + + } void UpdateGrid27::collisionAllNodes(int level, unsigned int t) @@ -391,4 +420,6 @@ UpdateGrid27::UpdateGrid27(SPtr<Parameter> para, vf::gpu::Communicator &comm, SP this->bcKernelManager = std::make_shared<BCKernelManager>(para, bcFactory); this->adKernelManager = std::make_shared<ADKernelManager>(para); this->gridScalingKernelManager = std::make_shared<GridScalingKernelManager>(para, scalingFactory); + timer = new Timer("ALM blade performance"); + timer->initTimer(); } diff --git a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h index 8ce2cf5bfd72f9f53cdb35bc92502ee9ca0d3ad8..da6d53c174c3178123a78ac59738bb390d0c036e 100644 --- a/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h +++ b/src/gpu/VirtualFluids_GPU/Calculation/UpdateGrid27.h @@ -17,6 +17,8 @@ class BoundaryConditionFactory; class GridScalingFactory; class TurbulenceModelFactory; class UpdateGrid27; +class Timer; + using CollisionStrategy = std::function<void (UpdateGrid27* updateGrid, Parameter* para, int level, unsigned int t)>; using RefinementStrategy = std::function<void (UpdateGrid27* updateGrid, Parameter* para, int level)>; @@ -84,6 +86,8 @@ private: std::shared_ptr<GridScalingKernelManager> gridScalingKernelManager; //! \property tmFactory is a shared pointer to an object of TurbulenceModelFactory std::shared_ptr<TurbulenceModelFactory> tmFactory; + + Timer* timer = nullptr; }; #endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu new file mode 100644 index 0000000000000000000000000000000000000000..b6b6b78f5f2957c787f1cce1fabd7039f07f0dac --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.cu @@ -0,0 +1,152 @@ +#include "CumulantK17Sponge.h" +#include <logger/Logger.h> +#include "Parameter/Parameter.h" +#include "Parameter/CudaStreamManager.h" +#include "CumulantK17Sponge_Device.cuh" + +#include <cuda.h> + +template<TurbulenceModel turbulenceModel> +std::shared_ptr< CumulantK17Sponge<turbulenceModel> > CumulantK17Sponge<turbulenceModel>::getNewInstance(std::shared_ptr<Parameter> para, int level) +{ + return std::shared_ptr<CumulantK17Sponge<turbulenceModel> >(new CumulantK17Sponge<turbulenceModel>(para,level)); +} + +template<TurbulenceModel turbulenceModel> +void CumulantK17Sponge<turbulenceModel>::run() +{ + LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads >>>( + para->getParD(level)->omega, + para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->rho, + para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ, + para->getParD(level)->turbViscosity, + para->getSGSConstant(), + (unsigned long)para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP, + para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ, + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep, + para->getParD(level)->taggedFluidNodeIndices[CollisionTemplate::Default], + para->getParD(level)->numberOfTaggedFluidNodes[CollisionTemplate::Default]); + + getLastCudaError("LB_Kernel_CumulantK17Sponge execution failed"); +} + +template<TurbulenceModel turbulenceModel> +void CumulantK17Sponge<turbulenceModel>::runOnIndices( const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex ) +{ + cudaStream_t stream = para->getStreamManager()->getStream(streamIndex); + + switch (collisionTemplate) + { + case CollisionTemplate::Default: + LB_Kernel_CumulantK17Sponge < turbulenceModel, false, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( + para->getParD(level)->omega, + para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->rho, + para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ, + para->getParD(level)->turbViscosity, + para->getSGSConstant(), + (unsigned long)para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP, + para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ, + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep, + indices, + size_indices); + break; + + case CollisionTemplate::WriteMacroVars: + LB_Kernel_CumulantK17Sponge < turbulenceModel, true, false > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( + para->getParD(level)->omega, + para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->rho, + para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ, + para->getParD(level)->turbViscosity, + para->getSGSConstant(), + (unsigned long)para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP, + para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ, + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep, + indices, + size_indices); + break; + + case CollisionTemplate::SubDomainBorder: + case CollisionTemplate::AllFeatures: + LB_Kernel_CumulantK17Sponge < turbulenceModel, true, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( + para->getParD(level)->omega, + para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->rho, + para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ, + para->getParD(level)->turbViscosity, + para->getSGSConstant(), + (unsigned long)para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP, + para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ, + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep, + indices, + size_indices); + break; + case CollisionTemplate::ApplyBodyForce: + LB_Kernel_CumulantK17Sponge < turbulenceModel, false, true > <<< cudaGrid.grid, cudaGrid.threads, 0, stream >>>( + para->getParD(level)->omega, + para->getParD(level)->neighborX, para->getParD(level)->neighborY, para->getParD(level)->neighborZ, + para->getParD(level)->distributions.f[0], + para->getParD(level)->rho, + para->getParD(level)->velocityX, para->getParD(level)->velocityY, para->getParD(level)->velocityZ, + para->getParD(level)->turbViscosity, + para->getSGSConstant(), + (unsigned long)para->getParD(level)->numberOfNodes, + level, + para->getForcesDev(), + para->getParD(level)->forceX_SP, para->getParD(level)->forceY_SP, para->getParD(level)->forceZ_SP, + para->getParD(level)->coordinateX, para->getParD(level)->coordinateY, para->getParD(level)->coordinateZ, + para->getQuadricLimitersDev(), + para->getParD(level)->isEvenTimestep, + indices, + size_indices); + break; + default: + throw std::runtime_error("Invalid CollisionTemplate in CumulantK17Sponge::runOnIndices()"); + break; + } + + getLastCudaError("LB_Kernel_CumulantK17Sponge execution failed"); +} + +template<TurbulenceModel turbulenceModel> +CumulantK17Sponge<turbulenceModel>::CumulantK17Sponge(std::shared_ptr<Parameter> para, int level) +{ + this->para = para; + this->level = level; + + myPreProcessorTypes.push_back(InitCompSP27); + + myKernelGroup = BasicKernel; + + this->cudaGrid = vf::cuda::CudaGrid(para->getParD(level)->numberofthreads, para->getParD(level)->numberOfNodes); + this->kernelUsesFluidNodeIndices = true; + + VF_LOG_INFO("Using turbulence model: {}", turbulenceModel); +} + +template class CumulantK17Sponge<TurbulenceModel::AMD>; +template class CumulantK17Sponge<TurbulenceModel::Smagorinsky>; +template class CumulantK17Sponge<TurbulenceModel::QR>; +template class CumulantK17Sponge<TurbulenceModel::None>; diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.h b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.h new file mode 100644 index 0000000000000000000000000000000000000000..60c4a70dc3ed8c7b1f9ad7c23c183356d9899410 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.h @@ -0,0 +1,20 @@ +#ifndef CUMULANT_K17_S_H +#define CUMULANT_K17_S_H + +#include "Kernel/KernelImp.h" +#include "Parameter/Parameter.h" + +template<TurbulenceModel turbulenceModel> +class CumulantK17Sponge : public KernelImp +{ +public: + static std::shared_ptr< CumulantK17Sponge<turbulenceModel> > getNewInstance(std::shared_ptr< Parameter> para, int level); + void run() override; + void runOnIndices(const unsigned int *indices, unsigned int size_indices, CollisionTemplate collisionTemplate, CudaStreamIndex streamIndex) override; + +private: + CumulantK17Sponge(); + CumulantK17Sponge(std::shared_ptr<Parameter> para, int level); +}; + +#endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu new file mode 100644 index 0000000000000000000000000000000000000000..c4598415b118e0d170b49cb285ffd0bbf1ad01cc --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cu @@ -0,0 +1,753 @@ + +// ____ ____ __ ______ __________ __ __ __ __ +// \ \ | | | | | _ \ |___ ___| | | | | / \ | | +// \ \ | | | | | |_) | | | | | | | / \ | | +// \ \ | | | | | _ / | | | | | | / /\ \ | | +// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ +// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| +// \ \ | | ________________________________________________________________ +// \ \ | | | ______________________________________________________________| +// \ \| | | | __ __ __ __ ______ _______ +// \ | | |_____ | | | | | | | | | _ \ / _____) +// \ | | _____| | | | | | | | | | | \ \ \_______ +// \ | | | | |_____ | \_/ | | | | |_/ / _____ | +// \ _____| |__| |________| \_______/ |__| |______/ (_______/ +// +// This file is part of VirtualFluids. VirtualFluids is free software: you can +// redistribute it and/or modify it under the terms of the GNU General Public +// License as published by the Free Software Foundation, either version 3 of +// the License, or (at your option) any later version. +// +// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT +// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or +// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License +// for more details. +// +// You should have received a copy of the GNU General Public License along +// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. +// +//! \file CumlantK17_Device.cu +//! \author Anna Wellmann, Martin Schönherr, Henry Korb, Henrik Asmuth +//! \date 05/12/2022 +//! \brief Kernel for CumulantK17Sponge including different turbulence models and options for local body forces and writing macroscopic variables +//! +//! CumulantK17Sponge kernel using chimera transformations and quartic limiters as present in Geier et al. (2017). Additional options are three different +//! eddy-viscosity turbulence models (Smagorinsky, AMD, QR) that can be set via the template parameter turbulenceModel (with default +//! TurbulenceModel::None). +//! The kernel is executed separately for each subset of fluid node indices with a different tag CollisionTemplate. For each subset, only the locally +//! required options are switched on ( \param writeMacroscopicVariables and/or \param applyBodyForce) in order to minimize memory accesses. The default +//! refers to the plain cumlant kernel (CollisionTemplate::Default). +//! Nodes are added to subsets (taggedFluidNodes) in Simulation::init using a corresponding tag with different values of CollisionTemplate. These subsets +//! are provided by the utilized PostCollisionInteractiors depending on they specifc requirements (e.g. writeMacroscopicVariables for probes). + +//======================================================================================= +/* Device code */ +#include "LBM/LB.h" +#include "lbm/constants/D3Q27.h" +#include <lbm/constants/NumericConstants.h> +#include "Kernel/Utilities/DistributionHelper.cuh" + +#include "GPU/TurbulentViscosityInlines.cuh" + +using namespace vf::lbm::constant; +using namespace vf::lbm::dir; +#include "Kernel/Utilities/ChimeraTransformation.h" + + +//////////////////////////////////////////////////////////////////////////////// +template<TurbulenceModel turbulenceModel, bool writeMacroscopicVariables, bool applyBodyForce> +__global__ void LB_Kernel_CumulantK17Sponge( + real omega_in, + uint* neighborX, + uint* neighborY, + uint* neighborZ, + real* distributions, + real* rho, + real* vx, + real* vy, + real* vz, + real* turbulentViscosity, + real SGSconstant, + unsigned long numberOfLBnodes, + int level, + real* forces, + real* bodyForceX, + real* bodyForceY, + real* bodyForceZ, + real* coordX, + real* coordY, + real* coordZ, + real* quadricLimiters, + bool isEvenTimestep, + const uint *fluidNodeIndices, + uint numberOfFluidNodes) +{ + ////////////////////////////////////////////////////////////////////////// + //! Cumulant K17 Kernel is based on \ref + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 + //! ]</b></a> and \ref <a href="https://doi.org/10.1016/j.jcp.2017.07.004"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.07.004 ]</b></a> + //! + //! The cumulant kernel is executed in the following steps + //! + //////////////////////////////////////////////////////////////////////////////// + //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim. + //! + const unsigned kThread = vf::gpu::getNodeIndex(); + + ////////////////////////////////////////////////////////////////////////// + // run for all indices in size_Mat and fluid nodes + if (kThread >= numberOfFluidNodes) + return; + //////////////////////////////////////////////////////////////////////////////// + //! - Get the node index from the array containing all indices of fluid nodes + //! + const unsigned k_000 = fluidNodeIndices[kThread]; + + ////////////////////////////////////////////////////////////////////////// + //! - Read distributions: style of reading and writing the distributions from/to stored arrays dependent on + //! timestep is based on the esoteric twist algorithm \ref <a + //! href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), + //! DOI:10.3390/computation5020019 ]</b></a> + //! + Distributions27 dist = vf::gpu::getDistributionReferences27(distributions, numberOfLBnodes, isEvenTimestep); + + //////////////////////////////////////////////////////////////////////////////// + //! - Set neighbor indices (necessary for indirect addressing) + uint k_M00 = neighborX[k_000]; + uint k_0M0 = neighborY[k_000]; + uint k_00M = neighborZ[k_000]; + uint k_MM0 = neighborY[k_M00]; + uint k_M0M = neighborZ[k_M00]; + uint k_0MM = neighborZ[k_0M0]; + uint k_MMM = neighborZ[k_MM0]; + //////////////////////////////////////////////////////////////////////////////////// + //! - Set local distributions + //! + real f_000 = (dist.f[DIR_000])[k_000]; + real f_P00 = (dist.f[DIR_P00])[k_000]; + real f_M00 = (dist.f[DIR_M00])[k_M00]; + real f_0P0 = (dist.f[DIR_0P0])[k_000]; + real f_0M0 = (dist.f[DIR_0M0])[k_0M0]; + real f_00P = (dist.f[DIR_00P])[k_000]; + real f_00M = (dist.f[DIR_00M])[k_00M]; + real f_PP0 = (dist.f[DIR_PP0])[k_000]; + real f_MM0 = (dist.f[DIR_MM0])[k_MM0]; + real f_PM0 = (dist.f[DIR_PM0])[k_0M0]; + real f_MP0 = (dist.f[DIR_MP0])[k_M00]; + real f_P0P = (dist.f[DIR_P0P])[k_000]; + real f_M0M = (dist.f[DIR_M0M])[k_M0M]; + real f_P0M = (dist.f[DIR_P0M])[k_00M]; + real f_M0P = (dist.f[DIR_M0P])[k_M00]; + real f_0PP = (dist.f[DIR_0PP])[k_000]; + real f_0MM = (dist.f[DIR_0MM])[k_0MM]; + real f_0PM = (dist.f[DIR_0PM])[k_00M]; + real f_0MP = (dist.f[DIR_0MP])[k_0M0]; + real f_PPP = (dist.f[DIR_PPP])[k_000]; + real f_MPP = (dist.f[DIR_MPP])[k_M00]; + real f_PMP = (dist.f[DIR_PMP])[k_0M0]; + real f_MMP = (dist.f[DIR_MMP])[k_MM0]; + real f_PPM = (dist.f[DIR_PPM])[k_00M]; + real f_MPM = (dist.f[DIR_MPM])[k_M0M]; + real f_PMM = (dist.f[DIR_PMM])[k_0MM]; + real f_MMM = (dist.f[DIR_MMM])[k_MMM]; + + //////////////////////////////////////////////////////////////////////////////////// + //! - Define aliases to use the same variable for the moments (m's): + //! + real& m_111 = f_000; + real& m_211 = f_P00; + real& m_011 = f_M00; + real& m_121 = f_0P0; + real& m_101 = f_0M0; + real& m_112 = f_00P; + real& m_110 = f_00M; + real& m_221 = f_PP0; + real& m_001 = f_MM0; + real& m_201 = f_PM0; + real& m_021 = f_MP0; + real& m_212 = f_P0P; + real& m_010 = f_M0M; + real& m_210 = f_P0M; + real& m_012 = f_M0P; + real& m_122 = f_0PP; + real& m_100 = f_0MM; + real& m_120 = f_0PM; + real& m_102 = f_0MP; + real& m_222 = f_PPP; + real& m_022 = f_MPP; + real& m_202 = f_PMP; + real& m_002 = f_MMP; + real& m_220 = f_PPM; + real& m_020 = f_MPM; + real& m_200 = f_PMM; + real& m_000 = f_MMM; + + //////////////////////////////////////////////////////(unsigned long)////////////////////////////// + //! - Calculate density and velocity using pyramid summation for low round-off errors as in Eq. (J1)-(J3) \ref + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> + //! + real drho = ((((f_PPP + f_MMM) + (f_MPM + f_PMP)) + ((f_MPP + f_PMM) + (f_MMP + f_PPM))) + + (((f_0MP + f_0PM) + (f_0MM + f_0PP)) + ((f_M0P + f_P0M) + (f_M0M + f_P0P)) + + ((f_MP0 + f_PM0) + (f_MM0 + f_PP0))) + + ((f_M00 + f_P00) + (f_0M0 + f_0P0) + (f_00M + f_00P))) + + f_000; + + real oneOverRho = c1o1 / (c1o1 + drho); + + real vvx = ((((f_PPP - f_MMM) + (f_PMP - f_MPM)) + ((f_PMM - f_MPP) + (f_PPM - f_MMP))) + + (((f_P0M - f_M0P) + (f_P0P - f_M0M)) + ((f_PM0 - f_MP0) + (f_PP0 - f_MM0))) + (f_P00 - f_M00)) * + oneOverRho; + real vvy = ((((f_PPP - f_MMM) + (f_MPM - f_PMP)) + ((f_MPP - f_PMM) + (f_PPM - f_MMP))) + + (((f_0PM - f_0MP) + (f_0PP - f_0MM)) + ((f_MP0 - f_PM0) + (f_PP0 - f_MM0))) + (f_0P0 - f_0M0)) * + oneOverRho; + real vvz = ((((f_PPP - f_MMM) + (f_PMP - f_MPM)) + ((f_MPP - f_PMM) + (f_MMP - f_PPM))) + + (((f_0MP - f_0PM) + (f_0PP - f_0MM)) + ((f_M0P - f_P0M) + (f_P0P - f_M0M))) + (f_00P - f_00M)) * + oneOverRho; + + //////////////////////////////////////////////////////////////////////////////////// + //! - Add half of the acceleration (body force) to the velocity as in Eq. (42) \ref + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> + //! + real factor = c1o1; + for (size_t i = 1; i <= level; i++) { + factor *= c2o1; + } + + real fx = forces[0]; + real fy = forces[1]; + real fz = forces[2]; + + if( applyBodyForce ){ + fx += bodyForceX[k_000]; + fy += bodyForceY[k_000]; + fz += bodyForceZ[k_000]; + + // real vx = vvx; + // real vy = vvy; + // real vz = vvz; + real acc_x = fx * c1o2 / factor; + real acc_y = fy * c1o2 / factor; + real acc_z = fz * c1o2 / factor; + + vvx += acc_x; + vvy += acc_y; + vvz += acc_z; + + // Reset body force. To be used when not using round-off correction. + bodyForceX[k_000] = 0.0f; + bodyForceY[k_000] = 0.0f; + bodyForceZ[k_000] = 0.0f; + + //////////////////////////////////////////////////////////////////////////////////// + //!> Round-off correction + //! + //!> Similar to Kahan summation algorithm (https://en.wikipedia.org/wiki/Kahan_summation_algorithm) + //!> Essentially computes the round-off error of the applied force and adds it in the next time step as a compensation. + //!> Seems to be necesseary at very high Re boundary layers, where the forcing and velocity can + //!> differ by several orders of magnitude. + //!> \note 16/05/2022: Testing, still ongoing! + //! + // bodyForceX[k_000] = (acc_x-(vvx-vx))*factor*c2o1; + // bodyForceY[k_000] = (acc_y-(vvy-vy))*factor*c2o1; + // bodyForceZ[k_000] = (acc_z-(vvz-vz))*factor*c2o1; + } + else{ + vvx += fx * c1o2 / factor; + vvy += fy * c1o2 / factor; + vvz += fz * c1o2 / factor; + } + + + //////////////////////////////////////////////////////////////////////////////////// + // calculate the square of velocities for this lattice node + real vx2 = vvx * vvx; + real vy2 = vvy * vvy; + real vz2 = vvz * vvz; + //////////////////////////////////////////////////////////////////////////////////// + //! - Set relaxation limiters for third order cumulants to default value \f$ \lambda=0.001 \f$ according to + //! section 6 in \ref <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + real quadricLimitP = quadricLimiters[0]; + real quadricLimitM = quadricLimiters[1]; + real quadricLimitD = quadricLimiters[2]; + //////////////////////////////////////////////////////////////////////////////////// + //! - Chimera transform from well conditioned distributions to central moments as defined in Appendix J in \ref + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> see also Eq. (6)-(14) in \ref <a + //! href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 + //! ]</b></a> + //! + //////////////////////////////////////////////////////////////////////////////////// + // Z - Dir + forwardInverseChimeraWithK(f_MMM, f_MM0, f_MMP, vvz, vz2, c36o1, c1o36); + forwardInverseChimeraWithK(f_M0M, f_M00, f_M0P, vvz, vz2, c9o1, c1o9); + forwardInverseChimeraWithK(f_MPM, f_MP0, f_MPP, vvz, vz2, c36o1, c1o36); + forwardInverseChimeraWithK(f_0MM, f_0M0, f_0MP, vvz, vz2, c9o1, c1o9); + forwardInverseChimeraWithK(f_00M, f_000, f_00P, vvz, vz2, c9o4, c4o9); + forwardInverseChimeraWithK(f_0PM, f_0P0, f_0PP, vvz, vz2, c9o1, c1o9); + forwardInverseChimeraWithK(f_PMM, f_PM0, f_PMP, vvz, vz2, c36o1, c1o36); + forwardInverseChimeraWithK(f_P0M, f_P00, f_P0P, vvz, vz2, c9o1, c1o9); + forwardInverseChimeraWithK(f_PPM, f_PP0, f_PPP, vvz, vz2, c36o1, c1o36); + + //////////////////////////////////////////////////////////////////////////////////// + // Y - Dir + forwardInverseChimeraWithK(f_MMM, f_M0M, f_MPM, vvy, vy2, c6o1, c1o6); + forwardChimera( f_MM0, f_M00, f_MP0, vvy, vy2); + forwardInverseChimeraWithK(f_MMP, f_M0P, f_MPP, vvy, vy2, c18o1, c1o18); + forwardInverseChimeraWithK(f_0MM, f_00M, f_0PM, vvy, vy2, c3o2, c2o3); + forwardChimera( f_0M0, f_000, f_0P0, vvy, vy2); + forwardInverseChimeraWithK(f_0MP, f_00P, f_0PP, vvy, vy2, c9o2, c2o9); + forwardInverseChimeraWithK(f_PMM, f_P0M, f_PPM, vvy, vy2, c6o1, c1o6); + forwardChimera( f_PM0, f_P00, f_PP0, vvy, vy2); + forwardInverseChimeraWithK(f_PMP, f_P0P, f_PPP, vvy, vy2, c18o1, c1o18); + + //////////////////////////////////////////////////////////////////////////////////// + // X - Dir + forwardInverseChimeraWithK(f_MMM, f_0MM, f_PMM, vvx, vx2, c1o1, c1o1); + forwardChimera( f_M0M, f_00M, f_P0M, vvx, vx2); + forwardInverseChimeraWithK(f_MPM, f_0PM, f_PPM, vvx, vx2, c3o1, c1o3); + forwardChimera( f_MM0, f_0M0, f_PM0, vvx, vx2); + forwardChimera( f_M00, f_000, f_P00, vvx, vx2); + forwardChimera( f_MP0, f_0P0, f_PP0, vvx, vx2); + forwardInverseChimeraWithK(f_MMP, f_0MP, f_PMP, vvx, vx2, c3o1, c1o3); + forwardChimera( f_M0P, f_00P, f_P0P, vvx, vx2); + forwardInverseChimeraWithK(f_MPP, f_0PP, f_PPP, vvx, vx2, c3o1, c1o9); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Setting relaxation rates for non-hydrodynamic cumulants (default values). Variable names and equations + //! according to <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! => [NAME IN PAPER]=[NAME IN CODE]=[DEFAULT VALUE]. + //! - Trace of second order cumulants \f$ C_{200}+C_{020}+C_{002} \f$ used to adjust bulk + //! viscosity:\f$\omega_2=OxxPyyPzz=1.0 \f$. + //! - Third order cumulants \f$ C_{120}+C_{102}, C_{210}+C_{012}, C_{201}+C_{021} \f$: \f$ \omega_3=OxyyPxzz + //! \f$ set according to Eq. (111) with simplifications assuming \f$ \omega_2=1.0\f$. + //! - Third order cumulants \f$ C_{120}-C_{102}, C_{210}-C_{012}, C_{201}-C_{021} \f$: \f$ \omega_4 = OxyyMxzz + //! \f$ set according to Eq. (112) with simplifications assuming \f$ \omega_2 = 1.0\f$. + //! - Third order cumulants \f$ C_{111} \f$: \f$ \omega_5 = Oxyz \f$ set according to Eq. (113) with + //! simplifications assuming \f$ \omega_2 = 1.0\f$ (modify for different bulk viscosity). + //! - Fourth order cumulants \f$ C_{220}, C_{202}, C_{022}, C_{211}, C_{121}, C_{112} \f$: for simplification + //! all set to the same default value \f$ \omega_6=\omega_7=\omega_8=O4=1.0 \f$. + //! - Fifth order cumulants \f$ C_{221}, C_{212}, C_{122}\f$: \f$\omega_9=O5=1.0\f$. + //! - Sixth order cumulant \f$ C_{222}\f$: \f$\omega_{10}=O6=1.0\f$. + //! + //////////////////////////////////////////////////////////////////////////////////// + //! - Calculate modified omega with turbulent viscosity + //! + real omega = omega_in; + if(turbulenceModel != TurbulenceModel::None){ omega /= (c1o1 + c3o1*omega_in*turbulentViscosity[k_000]); } + + //////////////////////////////////////////////////////////////////////////////////// + //! - A and DIR_00M: parameters for fourth order convergence of the diffusion term according to Eq. (115) and (116) + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> with simplifications assuming \f$ \omega_2 = 1.0 \f$ (modify for + //! different bulk viscosity). + //! + real factorA = (c4o1 + c2o1 * omega - c3o1 * omega * omega) / (c2o1 - c7o1 * omega + c5o1 * omega * omega); + real factorB = (c4o1 + c28o1 * omega - c14o1 * omega * omega) / (c6o1 - c21o1 * omega + c15o1 * omega * omega); + + //////////////////////////////////////////////////////////// + // 2. + real OxxPyyPzz = c1o1; + + //////////////////////////////////////////////////////////// + // 3. + + // Calculate modified omega for sponge bob layer + real startXsponge = 441.0f; + real endXsponge = 381.0f; + real sizeSponge = endXsponge - startXsponge; + + real OxyyPxzz = c8o1 * (-c2o1 + omega) * (c1o1 + c2o1 * omega) / (-c8o1 - c14o1 * omega + c7o1 * omega * omega); + real OxyyMxzz = + c8o1 * (-c2o1 + omega) * (-c7o1 + c4o1 * omega) / (c56o1 - c50o1 * omega + c9o1 * omega * omega); + real Oxyz = c24o1 * (-c2o1 + omega) * (-c2o1 - c7o1 * omega + c3o1 * omega * omega) / + (c48o1 + c152o1 * omega - c130o1 * omega * omega + c29o1 * omega * omega * omega); + + if (coordX[k_000] > startXsponge) { + real spongeFactor = (((endXsponge - coordX[k_000]) / sizeSponge) * c1o2) + c1o2; + omega = spongeFactor * omega; + + OxyyPxzz = c1o1; + OxyyMxzz = c1o1; + Oxyz = c1o1; + + quadricLimitM = c1o100; + quadricLimitP = c1o100; + quadricLimitD = c1o100; + + factorA = c0o1; + factorB = c0o1; + } + //////////////////////////////////////////////////////////// + // 4. + real O4 = c1o1; + //////////////////////////////////////////////////////////// + // 5. + real O5 = c1o1; + //////////////////////////////////////////////////////////// + // 6. + real O6 = c1o1; + + //////////////////////////////////////////////////////////////////////////////////// + //! - Compute cumulants from central moments according to Eq. (20)-(23) in + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + //////////////////////////////////////////////////////////// + // 4. + real c_211 = m_211 - ((m_200 + c1o3) * m_011 + c2o1 * m_110 * m_101) * oneOverRho; + real c_121 = m_121 - ((m_020 + c1o3) * m_101 + c2o1 * m_110 * m_011) * oneOverRho; + real c_112 = m_112 - ((m_002 + c1o3) * m_110 + c2o1 * m_101 * m_011) * oneOverRho; + + real c_220 = m_220 - (((m_200 * m_020 + c2o1 * m_110 * m_110) + c1o3 * (m_200 + m_020)) * oneOverRho - c1o9 * (drho * oneOverRho)); + real c_202 = m_202 - (((m_200 * m_002 + c2o1 * m_101 * m_101) + c1o3 * (m_200 + m_002)) * oneOverRho - c1o9 * (drho * oneOverRho)); + real c_022 = m_022 - (((m_002 * m_020 + c2o1 * m_011 * m_011) + c1o3 * (m_002 + m_020)) * oneOverRho - c1o9 * (drho * oneOverRho)); + //////////////////////////////////////////////////////////// + // 5. + real c_122 = + m_122 - ((m_002 * m_120 + m_020 * m_102 + c4o1 * m_011 * m_111 + c2o1 * (m_101 * m_021 + m_110 * m_012)) + + c1o3 * (m_120 + m_102)) * + oneOverRho; + real c_212 = + m_212 - ((m_002 * m_210 + m_200 * m_012 + c4o1 * m_101 * m_111 + c2o1 * (m_011 * m_201 + m_110 * m_102)) + + c1o3 * (m_210 + m_012)) * + oneOverRho; + real c_221 = + m_221 - ((m_200 * m_021 + m_020 * m_201 + c4o1 * m_110 * m_111 + c2o1 * (m_101 * m_120 + m_011 * m_210)) + + c1o3 * (m_021 + m_201)) * + oneOverRho; + //////////////////////////////////////////////////////////// + // 6. + real c_222 = m_222 + ((-c4o1 * m_111 * m_111 - (m_200 * m_022 + m_020 * m_202 + m_002 * m_220) - + c4o1 * (m_011 * m_211 + m_101 * m_121 + m_110 * m_112) - + c2o1 * (m_120 * m_102 + m_210 * m_012 + m_201 * m_021)) * + oneOverRho + + (c4o1 * (m_101 * m_101 * m_020 + m_011 * m_011 * m_200 + m_110 * m_110 * m_002) + + c2o1 * (m_200 * m_020 * m_002) + c16o1 * m_110 * m_101 * m_011) * + oneOverRho * oneOverRho - + c1o3 * (m_022 + m_202 + m_220) * oneOverRho - c1o9 * (m_200 + m_020 + m_002) * oneOverRho + + (c2o1 * (m_101 * m_101 + m_011 * m_011 + m_110 * m_110) + + (m_002 * m_020 + m_002 * m_200 + m_020 * m_200) + c1o3 * (m_002 + m_020 + m_200)) * + oneOverRho * oneOverRho * c2o3 + + c1o27 * ((drho * drho - drho) * oneOverRho * oneOverRho)); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Compute linear combinations of second and third order cumulants + //! + //////////////////////////////////////////////////////////// + // 2. + real mxxPyyPzz = m_200 + m_020 + m_002; + real mxxMyy = m_200 - m_020; + real mxxMzz = m_200 - m_002; + //////////////////////////////////////////////////////////// + // 3. + real mxxyPyzz = m_210 + m_012; + real mxxyMyzz = m_210 - m_012; + + real mxxzPyyz = m_201 + m_021; + real mxxzMyyz = m_201 - m_021; + + real mxyyPxzz = m_120 + m_102; + real mxyyMxzz = m_120 - m_102; + + //////////////////////////////////////////////////////////////////////////////////// + // incl. correction + //////////////////////////////////////////////////////////// + //! - Compute velocity gradients from second order cumulants according to Eq. (27)-(32) + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> Further explanations of the correction in viscosity in Appendix H of + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> Note that the division by rho is omitted here as we need rho times + //! the gradients later. + //! + real Dxy = -c3o1 * omega * m_110; + real Dxz = -c3o1 * omega * m_101; + real Dyz = -c3o1 * omega * m_011; + real dxux = c1o2 * (-omega) * (mxxMyy + mxxMzz) + c1o2 * OxxPyyPzz * (m_000 - mxxPyyPzz); + real dyuy = dxux + omega * c3o2 * mxxMyy; + real dzuz = dxux + omega * c3o2 * mxxMzz; + + //////////////////////////////////////////////////////////////////////////////////// + switch (turbulenceModel) + { + case TurbulenceModel::None: + case TurbulenceModel::AMD: //AMD is computed in separate kernel + break; + case TurbulenceModel::Smagorinsky: + turbulentViscosity[k_000] = calcTurbulentViscositySmagorinsky(SGSconstant, dxux, dyuy, dzuz, Dxy, Dxz , Dyz); + break; + case TurbulenceModel::QR: + turbulentViscosity[k_000] = calcTurbulentViscosityQR(SGSconstant, dxux, dyuy, dzuz, Dxy, Dxz , Dyz); + break; + default: + break; + } + //////////////////////////////////////////////////////////// + //! - Relaxation of second order cumulants with correction terms according to Eq. (33)-(35) in + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + mxxPyyPzz += OxxPyyPzz * (m_000 - mxxPyyPzz) - c3o1 * (c1o1 - c1o2 * OxxPyyPzz) * (vx2 * dxux + vy2 * dyuy + vz2 * dzuz); + mxxMyy += omega * (-mxxMyy) - c3o1 * (c1o1 + c1o2 * (-omega)) * (vx2 * dxux - vy2 * dyuy); + mxxMzz += omega * (-mxxMzz) - c3o1 * (c1o1 + c1o2 * (-omega)) * (vx2 * dxux - vz2 * dzuz); + + //////////////////////////////////////////////////////////////////////////////////// + ////no correction + // mxxPyyPzz += OxxPyyPzz*(mfaaa - mxxPyyPzz); + // mxxMyy += -(-omega) * (-mxxMyy); + // mxxMzz += -(-omega) * (-mxxMzz); + ////////////////////////////////////////////////////////////////////////// + m_011 += omega * (-m_011); + m_101 += omega * (-m_101); + m_110 += omega * (-m_110); + + //////////////////////////////////////////////////////////////////////////////////// + // relax + ////////////////////////////////////////////////////////////////////////// + // incl. limiter + //! - Relaxation of third order cumulants including limiter according to Eq. (116)-(123) + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + real wadjust = Oxyz + (c1o1 - Oxyz) * abs(m_111) / (abs(m_111) + quadricLimitD); + m_111 += wadjust * (-m_111); + wadjust = OxyyPxzz + (c1o1 - OxyyPxzz) * abs(mxxyPyzz) / (abs(mxxyPyzz) + quadricLimitP); + mxxyPyzz += wadjust * (-mxxyPyzz); + wadjust = OxyyMxzz + (c1o1 - OxyyMxzz) * abs(mxxyMyzz) / (abs(mxxyMyzz) + quadricLimitM); + mxxyMyzz += wadjust * (-mxxyMyzz); + wadjust = OxyyPxzz + (c1o1 - OxyyPxzz) * abs(mxxzPyyz) / (abs(mxxzPyyz) + quadricLimitP); + mxxzPyyz += wadjust * (-mxxzPyyz); + wadjust = OxyyMxzz + (c1o1 - OxyyMxzz) * abs(mxxzMyyz) / (abs(mxxzMyyz) + quadricLimitM); + mxxzMyyz += wadjust * (-mxxzMyyz); + wadjust = OxyyPxzz + (c1o1 - OxyyPxzz) * abs(mxyyPxzz) / (abs(mxyyPxzz) + quadricLimitP); + mxyyPxzz += wadjust * (-mxyyPxzz); + wadjust = OxyyMxzz + (c1o1 - OxyyMxzz) * abs(mxyyMxzz) / (abs(mxyyMxzz) + quadricLimitM); + mxyyMxzz += wadjust * (-mxyyMxzz); + ////////////////////////////////////////////////////////////////////////// + // no limiter + // mfbbb += OxyyMxzz * (-mfbbb); + // mxxyPyzz += OxyyPxzz * (-mxxyPyzz); + // mxxyMyzz += OxyyMxzz * (-mxxyMyzz); + // mxxzPyyz += OxyyPxzz * (-mxxzPyyz); + // mxxzMyyz += OxyyMxzz * (-mxxzMyyz); + // mxyyPxzz += OxyyPxzz * (-mxyyPxzz); + // mxyyMxzz += OxyyMxzz * (-mxyyMxzz); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Compute inverse linear combinations of second and third order cumulants + //! + m_200 = c1o3 * (mxxMyy + mxxMzz + mxxPyyPzz); + m_020 = c1o3 * (-c2o1 * mxxMyy + mxxMzz + mxxPyyPzz); + m_002 = c1o3 * (mxxMyy - c2o1 * mxxMzz + mxxPyyPzz); + + m_210 = ( mxxyMyzz + mxxyPyzz) * c1o2; + m_012 = (-mxxyMyzz + mxxyPyzz) * c1o2; + m_201 = ( mxxzMyyz + mxxzPyyz) * c1o2; + m_021 = (-mxxzMyyz + mxxzPyyz) * c1o2; + m_120 = ( mxyyMxzz + mxyyPxzz) * c1o2; + m_102 = (-mxyyMxzz + mxyyPxzz) * c1o2; + ////////////////////////////////////////////////////////////////////////// + + ////////////////////////////////////////////////////////////////////////// + // 4. + // no limiter + //! - Relax fourth order cumulants to modified equilibrium for fourth order convergence of diffusion according + //! to Eq. (43)-(48) <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + c_022 = -O4 * (c1o1 / omega - c1o2) * (dyuy + dzuz) * c2o3 * factorA + (c1o1 - O4) * (c_022); + c_202 = -O4 * (c1o1 / omega - c1o2) * (dxux + dzuz) * c2o3 * factorA + (c1o1 - O4) * (c_202); + c_220 = -O4 * (c1o1 / omega - c1o2) * (dyuy + dxux) * c2o3 * factorA + (c1o1 - O4) * (c_220); + c_112 = -O4 * (c1o1 / omega - c1o2) * Dxy * c1o3 * factorB + (c1o1 - O4) * (c_112); + c_121 = -O4 * (c1o1 / omega - c1o2) * Dxz * c1o3 * factorB + (c1o1 - O4) * (c_121); + c_211 = -O4 * (c1o1 / omega - c1o2) * Dyz * c1o3 * factorB + (c1o1 - O4) * (c_211); + + + ////////////////////////////////////////////////////////////////////////// + // 5. + c_122 += O5 * (-c_122); + c_212 += O5 * (-c_212); + c_221 += O5 * (-c_221); + + ////////////////////////////////////////////////////////////////////////// + // 6. + c_222 += O6 * (-c_222); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Compute central moments from post collision cumulants according to Eq. (53)-(56) in + //! <a href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), + //! DOI:10.1016/j.jcp.2017.05.040 ]</b></a> + //! + + ////////////////////////////////////////////////////////////////////////// + // 4. + m_211 = c_211 + c1o3 * ((c3o1 * m_200 + c1o1) * m_011 + c6o1 * m_110 * m_101) * oneOverRho; + m_121 = c_121 + c1o3 * ((c3o1 * m_020 + c1o1) * m_101 + c6o1 * m_110 * m_011) * oneOverRho; + m_112 = c_112 + c1o3 * ((c3o1 * m_002 + c1o1) * m_110 + c6o1 * m_101 * m_011) * oneOverRho; + + m_220 = + c_220 + (((m_200 * m_020 + c2o1 * m_110 * m_110) * c9o1 + c3o1 * (m_200 + m_020)) * oneOverRho - (drho * oneOverRho)) * c1o9; + m_202 = + c_202 + (((m_200 * m_002 + c2o1 * m_101 * m_101) * c9o1 + c3o1 * (m_200 + m_002)) * oneOverRho - (drho * oneOverRho)) * c1o9; + m_022 = + c_022 + (((m_002 * m_020 + c2o1 * m_011 * m_011) * c9o1 + c3o1 * (m_002 + m_020)) * oneOverRho - (drho * oneOverRho)) * c1o9; + + ////////////////////////////////////////////////////////////////////////// + // 5. + m_122 = c_122 + c1o3 * + (c3o1 * (m_002 * m_120 + m_020 * m_102 + c4o1 * m_011 * m_111 + c2o1 * (m_101 * m_021 + m_110 * m_012)) + + (m_120 + m_102)) * oneOverRho; + m_212 = c_212 + c1o3 * + (c3o1 * (m_002 * m_210 + m_200 * m_012 + c4o1 * m_101 * m_111 + c2o1 * (m_011 * m_201 + m_110 * m_102)) + + (m_210 + m_012)) * oneOverRho; + m_221 = c_221 + c1o3 * + (c3o1 * (m_200 * m_021 + m_020 * m_201 + c4o1 * m_110 * m_111 + c2o1 * (m_101 * m_120 + m_011 * m_210)) + + (m_021 + m_201)) * oneOverRho; + + ////////////////////////////////////////////////////////////////////////// + // 6. + m_222 = c_222 - ((-c4o1 * m_111 * m_111 - (m_200 * m_022 + m_020 * m_202 + m_002 * m_220) - + c4o1 * (m_011 * m_211 + m_101 * m_121 + m_110 * m_112) - + c2o1 * (m_120 * m_102 + m_210 * m_012 + m_201 * m_021)) * + oneOverRho + + (c4o1 * (m_101 * m_101 * m_020 + m_011 * m_011 * m_200 + m_110 * m_110 * m_002) + + c2o1 * (m_200 * m_020 * m_002) + c16o1 * m_110 * m_101 * m_011) * + oneOverRho * oneOverRho - + c1o3 * (m_022 + m_202 + m_220) * oneOverRho - c1o9 * (m_200 + m_020 + m_002) * oneOverRho + + (c2o1 * (m_101 * m_101 + m_011 * m_011 + m_110 * m_110) + + (m_002 * m_020 + m_002 * m_200 + m_020 * m_200) + c1o3 * (m_002 + m_020 + m_200)) * + oneOverRho * oneOverRho * c2o3 + + c1o27 * ((drho * drho - drho) * oneOverRho * oneOverRho)); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Add acceleration (body force) to first order cumulants according to Eq. (85)-(87) in + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> + //! + m_100 = -m_100; + m_010 = -m_010; + m_001 = -m_001; + + //Write to array here to distribute read/write + if(writeMacroscopicVariables) + { + rho[k_000] = drho; + vx[k_000] = vvx; + vy[k_000] = vvy; + vz[k_000] = vvz; + } + + //////////////////////////////////////////////////////////////////////////////////// + //! - Chimera transform from central moments to well conditioned distributions as defined in Appendix J in + //! <a href="https://doi.org/10.1016/j.camwa.2015.05.001"><b>[ M. Geier et al. (2015), + //! DOI:10.1016/j.camwa.2015.05.001 ]</b></a> see also Eq. (88)-(96) in <a + //! href="https://doi.org/10.1016/j.jcp.2017.05.040"><b>[ M. Geier et al. (2017), DOI:10.1016/j.jcp.2017.05.040 + //! ]</b></a> + //! + //////////////////////////////////////////////////////////////////////////////////// + // X - Dir + backwardInverseChimeraWithK(m_000, m_100, m_200, vvx, vx2, c1o1, c1o1); + backwardChimera( m_010, m_110, m_210, vvx, vx2); + backwardInverseChimeraWithK(m_020, m_120, m_220, vvx, vx2, c3o1, c1o3); + backwardChimera( m_001, m_101, m_201, vvx, vx2); + backwardChimera( m_011, m_111, m_211, vvx, vx2); + backwardChimera( m_021, m_121, m_221, vvx, vx2); + backwardInverseChimeraWithK(m_002, m_102, m_202, vvx, vx2, c3o1, c1o3); + backwardChimera( m_012, m_112, m_212, vvx, vx2); + backwardInverseChimeraWithK(m_022, m_122, m_222, vvx, vx2, c9o1, c1o9); + + //////////////////////////////////////////////////////////////////////////////////// + // Y - Dir + backwardInverseChimeraWithK(m_000, m_010, m_020, vvy, vy2, c6o1, c1o6); + backwardChimera( m_001, m_011, m_021, vvy, vy2); + backwardInverseChimeraWithK(m_002, m_012, m_022, vvy, vy2, c18o1, c1o18); + backwardInverseChimeraWithK(m_100, m_110, m_120, vvy, vy2, c3o2, c2o3); + backwardChimera( m_101, m_111, m_121, vvy, vy2); + backwardInverseChimeraWithK(m_102, m_112, m_122, vvy, vy2, c9o2, c2o9); + backwardInverseChimeraWithK(m_200, m_210, m_220, vvy, vy2, c6o1, c1o6); + backwardChimera( m_201, m_211, m_221, vvy, vy2); + backwardInverseChimeraWithK(m_202, m_212, m_222, vvy, vy2, c18o1, c1o18); + + //////////////////////////////////////////////////////////////////////////////////// + // Z - Dir + backwardInverseChimeraWithK(m_000, m_001, m_002, vvz, vz2, c36o1, c1o36); + backwardInverseChimeraWithK(m_010, m_011, m_012, vvz, vz2, c9o1, c1o9); + backwardInverseChimeraWithK(m_020, m_021, m_022, vvz, vz2, c36o1, c1o36); + backwardInverseChimeraWithK(m_100, m_101, m_102, vvz, vz2, c9o1, c1o9); + backwardInverseChimeraWithK(m_110, m_111, m_112, vvz, vz2, c9o4, c4o9); + backwardInverseChimeraWithK(m_120, m_121, m_122, vvz, vz2, c9o1, c1o9); + backwardInverseChimeraWithK(m_200, m_201, m_202, vvz, vz2, c36o1, c1o36); + backwardInverseChimeraWithK(m_210, m_211, m_212, vvz, vz2, c9o1, c1o9); + backwardInverseChimeraWithK(m_220, m_221, m_222, vvz, vz2, c36o1, c1o36); + + //////////////////////////////////////////////////////////////////////////////////// + //! - Write distributions: style of reading and writing the distributions from/to + //! stored arrays dependent on timestep is based on the esoteric twist algorithm + //! <a href="https://doi.org/10.3390/computation5020019"><b>[ M. Geier et al. (2017), + //! DOI:10.3390/computation5020019 ]</b></a> + //! + (dist.f[DIR_P00])[k_000] = f_M00; + (dist.f[DIR_M00])[k_M00] = f_P00; + (dist.f[DIR_0P0])[k_000] = f_0M0; + (dist.f[DIR_0M0])[k_0M0] = f_0P0; + (dist.f[DIR_00P])[k_000] = f_00M; + (dist.f[DIR_00M])[k_00M] = f_00P; + (dist.f[DIR_PP0])[k_000] = f_MM0; + (dist.f[DIR_MM0])[k_MM0] = f_PP0; + (dist.f[DIR_PM0])[k_0M0] = f_MP0; + (dist.f[DIR_MP0])[k_M00] = f_PM0; + (dist.f[DIR_P0P])[k_000] = f_M0M; + (dist.f[DIR_M0M])[k_M0M] = f_P0P; + (dist.f[DIR_P0M])[k_00M] = f_M0P; + (dist.f[DIR_M0P])[k_M00] = f_P0M; + (dist.f[DIR_0PP])[k_000] = f_0MM; + (dist.f[DIR_0MM])[k_0MM] = f_0PP; + (dist.f[DIR_0PM])[k_00M] = f_0MP; + (dist.f[DIR_0MP])[k_0M0] = f_0PM; + (dist.f[DIR_000])[k_000] = f_000; + (dist.f[DIR_PPP])[k_000] = f_MMM; + (dist.f[DIR_PMP])[k_0M0] = f_MPM; + (dist.f[DIR_PPM])[k_00M] = f_MMP; + (dist.f[DIR_PMM])[k_0MM] = f_MPP; + (dist.f[DIR_MPP])[k_M00] = f_PMM; + (dist.f[DIR_MMP])[k_MM0] = f_PPM; + (dist.f[DIR_MPM])[k_M0M] = f_PMP; + (dist.f[DIR_MMM])[k_MMM] = f_PPP; +} + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* bodyForceZ, real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, true, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, true, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, true, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, true, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, true > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::AMD, false, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::Smagorinsky, false, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::QR, false, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); + +template __global__ void LB_Kernel_CumulantK17Sponge < TurbulenceModel::None, false, false > ( real omega_in, uint* neighborX, uint* neighborY, uint* neighborZ, real* distributions, real* rho, real* vx, real* vy, real* vz, real* turbulentViscosity, real SGSconstant, unsigned long size_Mat, int level, real* forces, real* bodyForceX, real* bodyForceY, real* , real* coordX, real* coordY, real* coordZ, real* quadricLimiters, bool isEvenTimestep, const uint *fluidNodeIndices, uint numberOfFluidNodes); \ No newline at end of file diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh new file mode 100644 index 0000000000000000000000000000000000000000..236ea4b8d2b721b1491dce9f7f05b4ddd5590cb1 --- /dev/null +++ b/src/gpu/VirtualFluids_GPU/Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge_Device.cuh @@ -0,0 +1,32 @@ +#ifndef LB_Kernel_CUMULANT_K17_S_H +#define LB_Kernel_CUMULANT_K17_S_H + +#include <DataTypes.h> +#include <curand.h> + +template< TurbulenceModel turbulenceModel, bool writeMacroscopicVariables, bool applyBodyForce > __global__ void LB_Kernel_CumulantK17Sponge( + real omega_in, + uint* neighborX, + uint* neighborY, + uint* neighborZ, + real* distributions, + real* rho, + real* vx, + real* vy, + real* vz, + real* turbulentViscosity, + real SGSconstant, + unsigned long numberOfLBnodes, + int level, + real* forces, + real* bodyForceX, + real* bodyForceY, + real* bodyForceZ, + real* coordX, + real* coordY, + real* coordZ, + real* quadricLimiters, + bool isEvenTimestep, + const uint *fluidNodeIndices, + uint numberOfFluidNodes); +#endif diff --git a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp index 5a2d8c9a426e5cb23ca75f91aaf6fbff75cba72b..81e5ac06cc04f54e88088e7ac7dad7d7627a6ae0 100644 --- a/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp +++ b/src/gpu/VirtualFluids_GPU/Kernel/Utilities/KernelFactory/KernelFactoryImp.cpp @@ -11,6 +11,7 @@ #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Unified/CumulantK17Unified.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17chim/CumulantK17CompChim.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17/CumulantK17.h" +#include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Sponge/CumulantK17Sponge.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK17Bulk/CumulantK17BulkComp.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantAll4/CumulantAll4CompSP27.h" #include "Kernel/Kernels/BasicKernels/FluidFlow/Compressible/CumulantK18/CumulantK18Comp.h" @@ -141,11 +142,30 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> break; default: throw std::runtime_error("Unknown turbulence model!"); - break; - } - checkStrategy = FluidFlowCompStrategy::getInstance(); + break; + } + checkStrategy = FluidFlowCompStrategy::getInstance(); + } else if (kernel == "CumulantK17Sponge") { + switch (para->getTurbulenceModel()) { + case TurbulenceModel::AMD: + newKernel = CumulantK17Sponge<TurbulenceModel::AMD>::getNewInstance(para, level); + break; + case TurbulenceModel::Smagorinsky: + newKernel = CumulantK17Sponge<TurbulenceModel::Smagorinsky>::getNewInstance(para, level); + break; + case TurbulenceModel::QR: + newKernel = CumulantK17Sponge<TurbulenceModel::QR>::getNewInstance(para, level); + break; + case TurbulenceModel::None: + newKernel = CumulantK17Sponge<TurbulenceModel::None>::getNewInstance(para, level); + break; + default: + throw std::runtime_error("Unknown turbulence model!"); + break; + } + checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == "CumulantAll4CompSP27") { - newKernel = CumulantAll4CompSP27::getNewInstance(para, level); + newKernel = CumulantAll4CompSP27::getNewInstance(para, level); checkStrategy = FluidFlowCompStrategy::getInstance(); } else if (kernel == "CumulantK18Comp") { newKernel = CumulantK18Comp::getNewInstance(para, level); @@ -162,8 +182,8 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> } else if (kernel == "CumulantK15SpongeComp") { // /\ // newKernel = CumulantK15SpongeComp::getNewInstance(para, level); // || checkStrategy = FluidFlowCompStrategy::getInstance(); // compressible - } //=============== - else if ( kernel == "BGKIncompSP27") { // incompressible + } //=============== + else if (kernel == "BGKIncompSP27") { // incompressible newKernel = BGKIncompSP27::getNewInstance(para, level); // || checkStrategy = FluidFlowIncompStrategy::getInstance(); // \/ } else if (kernel == "BGKPlusIncompSP27") { @@ -181,11 +201,11 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> } else if (kernel == "CumulantIsoIncompSP27") { newKernel = CumulantIsoIncompSP27::getNewInstance(para, level); checkStrategy = FluidFlowIncompStrategy::getInstance(); - } else if (kernel == "CumulantK15Incomp") { // /\ // + } else if (kernel == "CumulantK15Incomp") { // /\ // newKernel = CumulantK15Incomp::getNewInstance(para, level); // || checkStrategy = FluidFlowIncompStrategy::getInstance(); // incompressible - } //=============== - else if (kernel == "PMCumulantOneCompSP27") { // porous media + } //=============== + else if (kernel == "PMCumulantOneCompSP27") { // porous media newKernel = PMCumulantOneCompSP27::getNewInstance(para, pm, level); // || checkStrategy = PMFluidFlowCompStrategy::getInstance(); // porous media } //=============== @@ -198,10 +218,10 @@ std::shared_ptr<Kernel> KernelFactoryImp::makeKernel(std::shared_ptr<Parameter> } else if (kernel == "WaleCumulantK15Comp") { newKernel = WaleCumulantK15Comp::getNewInstance(para, level); checkStrategy = WaleFluidFlowCompStrategy::getInstance(); - } else if (kernel == "WaleBySoniMalavCumulantK15Comp") { // /\ // + } else if (kernel == "WaleBySoniMalavCumulantK15Comp") { // /\ // newKernel = WaleBySoniMalavCumulantK15Comp::getNewInstance(para, level);// || checkStrategy = WaleFluidFlowCompStrategy::getInstance(); // wale model - } //=============== + } //=============== else { throw std::runtime_error("KernelFactory does not know the KernelType."); } diff --git a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp index 84ab84ff93fa7706bcc27d7e61a18f580f3c8dbe..83f39ee8248f3fc84bf45aa6c95f796ea6f85a61 100644 --- a/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp +++ b/src/gpu/VirtualFluids_GPU/LBM/Simulation.cpp @@ -370,8 +370,8 @@ void Simulation::init(GridProvider &gridProvider, BoundaryConditionFactory *bcFa ////////////////////////////////////////////////////////////////////////// // Write Initialized Files ////////////////////////////////////////////////////////////////////////// - VF_LOG_INFO("Write initialized Files ..."); - dataWriter->writeInit(para, cudaMemoryManager); + // VF_LOG_INFO("Write initialized Files ..."); + // dataWriter->writeInit(para, cudaMemoryManager); if (para->getCalcParticles()) copyAndPrintParticles(para.get(), cudaMemoryManager.get(), 0, true); VF_LOG_INFO("... done."); @@ -461,6 +461,10 @@ void Simulation::run() //////////////////////////////////////////////////////////////////////////////// // Time loop //////////////////////////////////////////////////////////////////////////////// + + puts("\n\n\n\n\n\n"); + VF_LOG_INFO("Start Simulation \n"); + for(timestep=para->getTimestepStart();timestep<=para->getTimestepEnd();timestep++) { this->updateGrid27->updateGrid(0, timestep); diff --git a/src/gpu/VirtualFluids_GPU/Output/Timer.cpp b/src/gpu/VirtualFluids_GPU/Output/Timer.cpp index 74a706165489a86cace40047beb09996aa0aa8db..b9ec5d9ba90a4beac5bc3d5201133cdee2e6ff23 100644 --- a/src/gpu/VirtualFluids_GPU/Output/Timer.cpp +++ b/src/gpu/VirtualFluids_GPU/Output/Timer.cpp @@ -25,6 +25,13 @@ void Timer::stopTimer() this->totalElapsedTime += this->elapsedTime; } +float Timer::startStopGetElapsed() +{ + this->stopTimer(); + this->startTimer(); + return this->elapsedTime; +} + void Timer::resetTimer() { this->elapsedTime = 0.0; diff --git a/src/gpu/VirtualFluids_GPU/Output/Timer.h b/src/gpu/VirtualFluids_GPU/Output/Timer.h index d035cbb6cef7ea9f8edabbd2894671a868c37eec..f6b6bad3cf29e31fb32ce998cd5591352919c487 100644 --- a/src/gpu/VirtualFluids_GPU/Output/Timer.h +++ b/src/gpu/VirtualFluids_GPU/Output/Timer.h @@ -34,6 +34,8 @@ class Timer float getElapsedTime(){ return this->elapsedTime; } float getTotalElapsedTime(){ return this->totalElapsedTime; } + float startStopGetElapsed(); + private: cudaEvent_t start_t, stop_t; diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp index 133b6dccf1dc45ee1cfa4ff1f29b4914cbeb6131..e593d16d6ed1f69ca65a22606a157e7ea9e6b111 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.cpp @@ -43,6 +43,7 @@ #include <basics/config/ConfigurationFile.h> +#include "Logger.h" #include "Parameter/CudaStreamManager.h" Parameter::Parameter() : Parameter(1, 0, {}) {} @@ -588,6 +589,30 @@ void Parameter::initLBMSimulationParameter() parD[i]->distY = parH[i]->distY; parD[i]->distZ = parH[i]->distZ; } + + checkParameterValidityCumulantK17(); +} + +void Parameter::checkParameterValidityCumulantK17() const +{ + if (this->mainKernel != "CumulantK17") + return; + + const real viscosity = this->parH[maxlevel]->vis; + const real viscosityLimit = 1.0 / 42.0; + if (viscosity > viscosityLimit) { + VF_LOG_WARNING("The viscosity (in LB units) at level {} is {:1.3g}. It is recommended to keep it smaller than {:1.3g} " + "for the CumulantK17 collision kernel.", + maxlevel, viscosity, viscosityLimit); + } + + const real velocity = this->ic.u0; + const real velocityLimit = 0.1; + if (velocity > velocityLimit) { + VF_LOG_WARNING("The velocity (in LB units) is {:1.4g}. It is recommended to keep it smaller than {:1.4g} for the " + "CumulantK17 collision kernel.", + velocity, velocityLimit); + } } void Parameter::copyMeasurePointsArrayToVector(int lev) diff --git a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h index 2e6d99fa6be200dd8c68af9371f655d7b1d827a1..fa45b1742f20e32258195c78b630ce95175af938 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h +++ b/src/gpu/VirtualFluids_GPU/Parameter/Parameter.h @@ -938,6 +938,8 @@ private: void setPathAndFilename(std::string fname); + void checkParameterValidityCumulantK17() const; + private: bool compOn{ false }; bool diffOn{ false }; diff --git a/src/gpu/VirtualFluids_GPU/Parameter/ParameterTest.cpp b/src/gpu/VirtualFluids_GPU/Parameter/ParameterTest.cpp index 4025acf7acad362e9f0f3702cb897b9c1b6dbf3b..72a12ae880556e6e257eb69dee4e806617252629 100644 --- a/src/gpu/VirtualFluids_GPU/Parameter/ParameterTest.cpp +++ b/src/gpu/VirtualFluids_GPU/Parameter/ParameterTest.cpp @@ -1,4 +1,3 @@ -#include <gmock/gmock.h> #include "basics/tests/testUtilities.h" #include <filesystem> @@ -8,7 +7,6 @@ #include "Parameter.h" #include "basics/config/ConfigurationFile.h" - TEST(ParameterTest, passingEmptyFileWithoutPath_ShouldNotThrow) { // assuming that the config files is stored parallel to this file. @@ -37,7 +35,9 @@ TEST(ParameterTest, check_all_Parameter_CanBePassedToConstructor) // test optional parameter EXPECT_THAT(para.getOutputPath(), testing::Eq("/output/path/")); - EXPECT_THAT(para.getGridPath(), testing::Eq("/path/to/grid/")); // ... all grid files (e.g. multi-gpu/ multi-level) could be tested as well + EXPECT_THAT( + para.getGridPath(), + testing::Eq("/path/to/grid/")); // ... all grid files (e.g. multi-gpu/ multi-level) could be tested as well EXPECT_THAT(para.getgeoVec(), testing::Eq("/path/to/grid/geoVec.dat")); EXPECT_THAT(para.getMaxDev(), testing::Eq(2)); EXPECT_THAT(para.getDevices(), testing::ElementsAreArray({ 2, 3 })); @@ -163,7 +163,7 @@ TEST(ParameterTest, setGridPathOverridesDefaultGridPath) Parameter para(2, 1); para.setGridPath("gridPathTest"); - EXPECT_THAT( para.getGridPath(), testing::Eq("gridPathTest/1/")); + EXPECT_THAT(para.getGridPath(), testing::Eq("gridPathTest/1/")); EXPECT_THAT(para.getConcentration(), testing::Eq("gridPathTest/1/conc.dat")); } @@ -177,9 +177,8 @@ TEST(ParameterTest, setGridPathOverridesConfigFile) auto para = Parameter(2, 0, &config); para.setGridPath("gridPathTest"); - EXPECT_THAT( para.getGridPath(), testing::Eq("gridPathTest/0/")); + EXPECT_THAT(para.getGridPath(), testing::Eq("gridPathTest/0/")); EXPECT_THAT(para.getConcentration(), testing::Eq("gridPathTest/0/conc.dat")); - } TEST(ParameterTest, userMissedSlash) @@ -189,7 +188,6 @@ TEST(ParameterTest, userMissedSlash) EXPECT_THAT(para.getGridPath(), testing::Eq("gridPathTest/")); EXPECT_THAT(para.getConcentration(), testing::Eq("gridPathTest/conc.dat")); - } TEST(ParameterTest, userMissedSlashMultiGPU) @@ -199,4 +197,87 @@ TEST(ParameterTest, userMissedSlashMultiGPU) EXPECT_THAT(para.getGridPath(), testing::Eq("gridPathTest/0/")); EXPECT_THAT(para.getConcentration(), testing::Eq("gridPathTest/0/conc.dat")); -} \ No newline at end of file +} + +class ParameterTestCumulantK17 : public testing::Test +{ +protected: + void SetUp() override + { + } + + bool stdoutContainsWarning() + { + std::string output = testing::internal::GetCapturedStdout(); + return output.find("warning") != std::string::npos; + } + + Parameter para; +}; + +TEST_F(ParameterTestCumulantK17, CumulantK17_VelocityIsTooHigh_expectWarning) +{ + + para.setVelocityLB(0.11); + para.setMainKernel("CumulantK17"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_TRUE(stdoutContainsWarning()); +} + +TEST_F(ParameterTestCumulantK17, CumulantK17_VelocityIsOk_expectNoWarning) +{ + para.setVelocityLB(0.09); + para.setMainKernel("CumulantK17"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_FALSE(stdoutContainsWarning()); +} + +TEST_F(ParameterTestCumulantK17, NotCumulantK17_VelocityIsTooHigh_expectNoWarning) +{ + para.setVelocityLB(42); + para.setMainKernel("K"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_FALSE(stdoutContainsWarning()); +} + +TEST_F(ParameterTestCumulantK17, CumulantK17_ViscosityIsTooHigh_expectWarning) +{ + para.setViscosityLB(0.024); + para.setMainKernel("CumulantK17"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_TRUE(stdoutContainsWarning()); +} + +TEST_F(ParameterTestCumulantK17, CumulantK17_ViscosityIsOk_expectNoWarning) +{ + para.setViscosityLB(0.023); + para.setMainKernel("CumulantK17"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_FALSE(stdoutContainsWarning()); +} + +TEST_F(ParameterTestCumulantK17, NotCumulantK17_ViscosityIsTooHigh_expectNoWarning) +{ + para.setViscosityLB(10); + para.setMainKernel("K"); + testing::internal::CaptureStdout(); + + para.initLBMSimulationParameter(); + + EXPECT_FALSE(stdoutContainsWarning()); +} diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.cu b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.cu index 9447a8636e801c132df9cef2feced4b5ab4e68de..75109a15c914048791102a816d6afbae7a59c420 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.cu +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.cu @@ -1,51 +1,22 @@ -//======================================================================================= -// ____ ____ __ ______ __________ __ __ __ __ -// \ \ | | | | | _ \ |___ ___| | | | | / \ | | -// \ \ | | | | | |_) | | | | | | | / \ | | -// \ \ | | | | | _ / | | | | | | / /\ \ | | -// \ \ | | | | | | \ \ | | | \__/ | / ____ \ | |____ -// \ \ | | |__| |__| \__\ |__| \________/ /__/ \__\ |_______| -// \ \ | | ________________________________________________________________ -// \ \ | | | ______________________________________________________________| -// \ \| | | | __ __ __ __ ______ _______ -// \ | | |_____ | | | | | | | | | _ \ / _____) -// \ | | _____| | | | | | | | | | | \ \ \_______ -// \ | | | | |_____ | \_/ | | | | |_/ / _____ | -// \ _____| |__| |________| \_______/ |__| |______/ (_______/ -// -// This file is part of VirtualFluids. VirtualFluids is free software: you can -// redistribute it and/or modify it under the terms of the GNU General Public -// License as published by the Free Software Foundation, either version 3 of -// the License, or (at your option) any later version. -// -// VirtualFluids is distributed in the hope that it will be useful, but WITHOUT -// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or -// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License -// for more details. -// -// You should have received a copy of the GNU General Public License along -// with VirtualFluids (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>. -// -//! \file ActuatorFarm.cu -//! \ingroup PreCollisionInteractor -//! \author Henrik Asmuth, Henry Korb -//====================================================================================== #include "ActuatorFarm.h" #include <cuda.h> #include <cuda_runtime.h> #include <helper_cuda.h> -#include "cuda/CudaGrid.h" +#include <cuda/CudaGrid.h> #include "VirtualFluids_GPU/GPU/GeometryUtils.h" -#include "LBM/GPUHelperFunctions/KernelUtilities.h" +#include "VirtualFluids_GPU/Kernel/Utilities/DistributionHelper.cuh" #include "Parameter/Parameter.h" #include "Parameter/CudaStreamManager.h" #include "DataStructureInitializer/GridProvider.h" #include "GPU/CudaMemoryManager.h" -#include "lbm/constants/NumericConstants.h" -#include "logger/Logger.h" +#include <lbm/constants/NumericConstants.h> +#include <logger/Logger.h> +#include <ostream> + +#include "Output/Timer.h" using namespace vf::lbm::constant; @@ -111,20 +82,17 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real uint* bladeIndices, real velocityRatio, real invDeltaX) { - //////////////////////////////////////////////////////////////////////////////// - //! - Get node index coordinates from threadIdx, blockIdx, blockDim and gridDim. - //! - const unsigned nodeIndex = vf::gpu::getNodeIndex(); + const uint node = vf::gpu::getNodeIndex(); - if(nodeIndex>=numberOfBladeNodes*numberOfBlades*numberOfTurbines) return; + if(node>=numberOfBladeNodes*numberOfBlades*numberOfTurbines) return; uint turbine, bladeNode, blade; - calcTurbineBladeAndBladeNode(nodeIndex, bladeNode, numberOfBladeNodes, blade, numberOfBlades, turbine, numberOfTurbines); + calcTurbineBladeAndBladeNode(node, bladeNode, numberOfBladeNodes, blade, numberOfBlades, turbine, numberOfTurbines); - real bladeCoordX_BF = bladeCoordsX[nodeIndex]; - real bladeCoordY_BF = bladeCoordsY[nodeIndex]; - real bladeCoordZ_BF = bladeCoordsZ[nodeIndex]; + real bladeCoordX_BF = bladeCoordsX[node]; + real bladeCoordY_BF = bladeCoordsY[node]; + real bladeCoordZ_BF = bladeCoordsZ[node]; real bladeCoordX_GF, bladeCoordY_GF, bladeCoordZ_GF; @@ -143,12 +111,12 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real uint k, ke, kn, kt; uint kne, kte, ktn, ktne; - k = findNearestCellBSW(bladeIndices[nodeIndex], + k = findNearestCellBSW(bladeIndices[node], gridCoordsX, gridCoordsY, gridCoordsZ, bladeCoordX_GF, bladeCoordY_GF, bladeCoordZ_GF, neighborsX, neighborsY, neighborsZ, neighborsWSB); - bladeIndices[nodeIndex] = k; + bladeIndices[node] = k; getNeighborIndicesOfBSW(k, ke, kn, kt, kne, kte, ktn, ktne, neighborsX, neighborsY, neighborsZ); @@ -170,9 +138,9 @@ __global__ void interpolateVelocities(real* gridCoordsX, real* gridCoordsY, real bladeVelX_GF, bladeVelY_GF, bladeVelZ_GF, localAzimuth, yaw); - bladeVelocitiesX[nodeIndex] = bladeVelX_BF; - bladeVelocitiesY[nodeIndex] = bladeVelY_BF+omegas[turbine]*bladeCoordZ_BF; - bladeVelocitiesZ[nodeIndex] = bladeVelZ_BF; + bladeVelocitiesX[node] = bladeVelX_BF; + bladeVelocitiesY[node] = bladeVelY_BF+omegas[turbine]*bladeCoordZ_BF; + bladeVelocitiesZ[node] = bladeVelZ_BF; } @@ -302,26 +270,38 @@ void ActuatorFarm::addTurbine(real posX, real posY, real posZ, real diameter, re void ActuatorFarm::init(Parameter* para, GridProvider* gridProvider, CudaMemoryManager* cudaMemoryManager) { - if(!para->getIsBodyForce()) throw std::runtime_error("try to allocate ActuatorFarm but BodyForce is not set in Parameter."); + if (!para->getIsBodyForce()) + throw std::runtime_error("try to allocate ActuatorFarm but BodyForce is not set in Parameter."); this->forceRatio = para->getForceRatio(); this->initTurbineGeometries(cudaMemoryManager); - this->initBladeCoords(cudaMemoryManager); + this->initBladeCoords(cudaMemoryManager); this->initBladeIndices(para, cudaMemoryManager); this->initBladeVelocities(cudaMemoryManager); - this->initBladeForces(cudaMemoryManager); - this->initBoundingSpheres(para, cudaMemoryManager); + this->initBladeForces(cudaMemoryManager); + this->initBoundingSpheres(para, cudaMemoryManager); this->streamIndex = 0; + + bladeTimer = new Timer("ALM blade performance"); + bladeTimer->initTimer(); } -void ActuatorFarm::interact(Parameter* para, CudaMemoryManager* cudaMemoryManager, int level, unsigned int t) +void ActuatorFarm::interact(Parameter* para, CudaMemoryManager* cudaMemoryManager, int currentLevel, unsigned int t) { - if (level != this->level) return; + if (currentLevel != this->level) return; + bool useTimer = false; cudaStream_t stream = para->getStreamManager()->getStream(CudaStreamIndex::ActuatorFarm, this->streamIndex); + if (useTimer) + std::cout << "ActuatorFarm::interact: level = " << currentLevel << ", t = " << t << " useHostArrays = " << useHostArrays <<std::endl; + bladeTimer->startTimer(); + if(useHostArrays) cudaMemoryManager->cudaCopyBladeCoordsHtoD(this); - vf::cuda::CudaGrid bladeGrid = vf::cuda::CudaGrid(para->getParH(level)->numberofthreads, this->numberOfNodes); + vf::cuda::CudaGrid bladeGrid = vf::cuda::CudaGrid(para->getParH(currentLevel)->numberofthreads, this->numberOfNodes); + + if (useTimer) + std::cout << " cudaCopyBladeCoordsHtoD, " << bladeTimer->startStopGetElapsed() << std::endl; interpolateVelocities<<< bladeGrid.grid, bladeGrid.threads, 0, stream >>>( para->getParD(this->level)->coordinateX, para->getParD(this->level)->coordinateY, para->getParD(this->level)->coordinateZ, @@ -334,14 +314,28 @@ void ActuatorFarm::interact(Parameter* para, CudaMemoryManager* cudaMemoryManage this->turbinePosXD, this->turbinePosYD, this->turbinePosZD, this->bladeIndicesD, para->getVelocityRatio(), this->invDeltaX); + if (useTimer) + std::cout << " interpolateVelocities, " << bladeTimer->startStopGetElapsed() << std::endl; + cudaStreamSynchronize(stream); if(useHostArrays) cudaMemoryManager->cudaCopyBladeVelocitiesDtoH(this); + + if (useTimer) + std::cout << " cudaCopyBladeVelocitiesDtoH, " << bladeTimer->startStopGetElapsed() << std::endl; + this->calcBladeForces(); + + if (useTimer) + std::cout << " calcBladeForces, " << bladeTimer->startStopGetElapsed() << std::endl; + this->swapDeviceArrays(); if(useHostArrays) cudaMemoryManager->cudaCopyBladeForcesHtoD(this); - vf::cuda::CudaGrid sphereGrid = vf::cuda::CudaGrid(para->getParH(level)->numberofthreads, this->numberOfIndices); + if (useTimer) + std::cout << " cudaCopyBladeForcesHtoD, " << bladeTimer->startStopGetElapsed() << std::endl; + + vf::cuda::CudaGrid sphereGrid = vf::cuda::CudaGrid(para->getParH(currentLevel)->numberofthreads, this->numberOfIndices); applyBodyForces<<<sphereGrid.grid, sphereGrid.threads, 0, stream>>>( para->getParD(this->level)->coordinateX, para->getParD(this->level)->coordinateY, para->getParD(this->level)->coordinateZ, @@ -353,12 +347,25 @@ void ActuatorFarm::interact(Parameter* para, CudaMemoryManager* cudaMemoryManage this->turbinePosXD, this->turbinePosYD, this->turbinePosZD, this->boundingSphereIndicesD, this->numberOfIndices, this->invEpsilonSqrd, this->factorGaussian); + + if (useTimer) + std::cout << " applyBodyForces, " << bladeTimer->startStopGetElapsed() << std::endl; + cudaMemoryManager->cudaCopyBladeOrientationsHtoD(this); + + if (useTimer) + std::cout << " cudaCopyBladeOrientationsHtoD, " << bladeTimer->startStopGetElapsed() << std::endl; + if (useTimer) + std::cout << "total time, " << bladeTimer->getTotalElapsedTime() << std::endl; + bladeTimer->resetTimer(); + cudaStreamSynchronize(stream); + + } -void ActuatorFarm::free(Parameter* para, CudaMemoryManager* cudaMemoryManager) +void ActuatorFarm::free(Parameter* /*para*/, CudaMemoryManager* cudaMemoryManager) { cudaMemoryManager->cudaFreeBladeGeometries(this); cudaMemoryManager->cudaFreeBladeOrientations(this); @@ -543,11 +550,11 @@ void ActuatorFarm::initBoundingSpheres(Parameter* para, CudaMemoryManager* cudaM } } - if(nodesInThisSphere<minimumNumberOfNodesPerSphere) - { - VF_LOG_CRITICAL("Found only {} nodes in bounding sphere of turbine no. {}, expected at least {}!", nodesInThisSphere, turbine, minimumNumberOfNodesPerSphere); - throw std::runtime_error("ActuatorFarm::initBoundingSpheres: Turbine bounding sphere partially out of domain."); - } + // if(nodesInThisSphere<minimumNumberOfNodesPerSphere) + // { + // VF_LOG_CRITICAL("Found only {} nodes in bounding sphere of turbine no. {}, expected at least {}!", nodesInThisSphere, turbine, minimumNumberOfNodesPerSphere); + // throw std::runtime_error("ActuatorFarm::initBoundingSpheres: Turbine bounding sphere partially out of domain."); + // } } this->numberOfIndices = uint(nodesInSpheres.size()); diff --git a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.h b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.h index 8e21cdb6b21efd323f6723e21d6b28614109f1ec..e6066e28620581324cf105e384b939c8deda075b 100644 --- a/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.h +++ b/src/gpu/VirtualFluids_GPU/PreCollisionInteractor/ActuatorFarm.h @@ -10,6 +10,7 @@ using namespace vf::lbm::constant; class Parameter; class GridProvider; +class Timer; using namespace vf::lbm::constant; class ActuatorFarm : public PreCollisionInteractor @@ -192,6 +193,8 @@ private: uint numberOfNodes; real forceRatio, factorGaussian, invEpsilonSqrd, invDeltaX; int streamIndex; + + Timer* bladeTimer = nullptr; }; #endif