waLBerla 7.2
Loading...
Searching...
No Matches
Classes | Typedefs | Functions | Variables
walberla::lbm_mesapd_coupling::psm::gpu Namespace Reference

Classes

class  BoxFractionMappingSweep
 
struct  ParticleAndVolumeFractionSoA_T
 
class  PSMSweepCollection
 
class  ReduceParticleForcesSweep
 
class  SetParticleVelocitiesSweep
 
class  SphereFractionMappingSweep
 

Typedefs

using nOverlappingParticlesField_T = GhostLayerField< uint_t, 1 >
 
using BsField_T = GhostLayerField< real_t, MaxParticlesPerCell >
 
using idxField_T = GhostLayerField< size_t, MaxParticlesPerCell >
 
using BField_T = GhostLayerField< real_t, 1 >
 
using particleVelocitiesField_T = GhostLayerField< real_t, MaxParticlesPerCell * 3 >
 
using particleForcesField_T = GhostLayerField< real_t, MaxParticlesPerCell * 3 >
 

Functions

template<int Weighting_T>
__global__ void superSampling (walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField, walberla::gpu::FieldAccessor< real_t > BsField, walberla::gpu::FieldAccessor< id_t > idxField, walberla::gpu::FieldAccessor< real_t > BField, const real_t omega, const real_t *__restrict__ const spherePositions, const real_t *__restrict__ const sphereRadii, const double3 blockStart, const real_t dx, const int3 nSamples, const size_t *__restrict__ const numParticlesSubBlocks, const size_t *__restrict__ const particleIDsSubBlocks, const uint3 subBlocksPerDim)
 
template<int Weighting_T>
__global__ void linearApproximation (walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField, walberla::gpu::FieldAccessor< real_t > BsField, walberla::gpu::FieldAccessor< id_t > idxField, walberla::gpu::FieldAccessor< real_t > BField, const real_t omega, const real_t *__restrict__ const spherePositions, const real_t *__restrict__ const sphereRadii, const real_t *__restrict__ const f_rs, const double3 blockStart, const real_t dx, const size_t *__restrict__ const numParticlesSubBlocks, const size_t *__restrict__ const particleIDsSubBlocks, const uint3 subBlocksPerDim)
 
template<int Weighting_T>
__global__ void boxMapping (walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField, walberla::gpu::FieldAccessor< real_t > BsField, walberla::gpu::FieldAccessor< id_t > idxField, walberla::gpu::FieldAccessor< real_t > BField, const real_t omega, const double3 boxPositionMin, const double3 boxPositionMax, const double3 blockStart, const real_t dx, const id_t idxMapped)
 
template<int Weighting_T>
void calculateWeighting (real_t *const, const real_t &, const real_t &)
 
template<>
void calculateWeighting< 1 > (real_t *const weighting, const real_t &epsilon, const real_t &)
 
template<>
void calculateWeighting< 2 > (real_t *const weighting, const real_t &epsilon, const real_t &tau)
 
template<int Weighting_T>
void mapParticles (IBlock &blockIt, const ParticleAndVolumeFractionSoA_T< Weighting_T > &particleAndVolumeFractionSoA, const real_t *const spherePositions, const real_t *const sphereRadii, const real_t *const f_rs, const size_t *const numParticlesSubBlocks, const size_t *const particleIDsSubBlocks, const Vector3< uint_t > subBlocksPerDim)
 
template<int Weighting_T>
void mapParticles (const IBlock &blockIt, const ParticleAndVolumeFractionSoA_T< Weighting_T > &particleAndVolumeFractionSoA, const real_t *const spherePositions, const real_t *const sphereRadii, const real_t *const f_rs, const size_t *const numParticlesSubBlocks, const size_t *const particleIDsSubBlocks, const Vector3< uint_t > subBlocksPerDim)
 
template<typename SweepCollection , typename PSMSweep >
void addPSMSweepsToTimeloop (SweepTimeloop &timeloop, SweepCollection &psmSweepCollection, PSMSweep &psmSweep, bool synchronize=true)
 
template<typename SweepCollection , typename PSMSweep , typename Communication >
void addPSMSweepsToTimeloops (SweepTimeloop &commTimeloop, SweepTimeloop &timeloop, Communication &comm, SweepCollection &psmSweepCollection, PSMSweep &psmSweep, bool synchronize=true)
 
__device__ void cross (real_t *__restrict__ const crossResult, const real_t *__restrict__ const lhs, const real_t *__restrict__ const rhs)
 
__device__ void getVelocityAtWFPoint (real_t *__restrict__ const velocityAtWFPoint, const real_t *__restrict__ const linearVelocity, const real_t *__restrict__ const angularVelocity, const real_t *__restrict__ const position, const real_t *__restrict__ const wf_pt)
 
__device__ void addHydrodynamicForceTorqueAtWFPosAtomic (real_t *__restrict__ const particleForce, real_t *__restrict__ const particleTorque, const real_t *__restrict__ const f, const real_t *__restrict__ const pos, const real_t *__restrict__ const wf_pt)
 
__global__ void SetParticleVelocities (walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField, walberla::gpu::FieldAccessor< uint_t > idxField, walberla::gpu::FieldAccessor< real_t > particleVelocitiesField, real_t *__restrict__ const linearVelocities, real_t *__restrict__ const angularVelocities, real_t *__restrict__ const positions, const double3 blockStart, const real_t dx)
 
__global__ void ReduceParticleForces (walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField, walberla::gpu::FieldAccessor< id_t > idxField, walberla::gpu::FieldAccessor< real_t > particleForcesField, real_t *__restrict__ const hydrodynamicForces, real_t *__restrict__ const hydrodynamicTorques, real_t *__restrict__ const positions, const double3 blockStart, const real_t dx, const real_t forceScalingFactor)
 

Variables

const uint_t MaxParticlesPerCell = MAX_PARTICLES_PER_CELL
 
auto deviceSyncWrapper
 

Typedef Documentation

◆ BField_T

using walberla::lbm_mesapd_coupling::psm::gpu::BField_T = GhostLayerField< real_t, 1 >

◆ BsField_T

using walberla::lbm_mesapd_coupling::psm::gpu::BsField_T = GhostLayerField< real_t, MaxParticlesPerCell >

◆ idxField_T

using walberla::lbm_mesapd_coupling::psm::gpu::idxField_T = GhostLayerField< size_t, MaxParticlesPerCell >

◆ nOverlappingParticlesField_T

using walberla::lbm_mesapd_coupling::psm::gpu::nOverlappingParticlesField_T = GhostLayerField< uint_t, 1 >

◆ particleForcesField_T

using walberla::lbm_mesapd_coupling::psm::gpu::particleForcesField_T = GhostLayerField< real_t, MaxParticlesPerCell * 3 >

◆ particleVelocitiesField_T

using walberla::lbm_mesapd_coupling::psm::gpu::particleVelocitiesField_T = GhostLayerField< real_t, MaxParticlesPerCell * 3 >

Function Documentation

◆ addHydrodynamicForceTorqueAtWFPosAtomic()

__device__ void walberla::lbm_mesapd_coupling::psm::gpu::addHydrodynamicForceTorqueAtWFPosAtomic ( real_t *__restrict__ const particleForce,
real_t *__restrict__ const particleTorque,
const real_t *__restrict__ const f,
const real_t *__restrict__ const pos,
const real_t *__restrict__ const wf_pt )
inline

◆ addPSMSweepsToTimeloop()

template<typename SweepCollection , typename PSMSweep >
void walberla::lbm_mesapd_coupling::psm::gpu::addPSMSweepsToTimeloop ( SweepTimeloop & timeloop,
SweepCollection & psmSweepCollection,
PSMSweep & psmSweep,
bool synchronize = true )

◆ addPSMSweepsToTimeloops()

template<typename SweepCollection , typename PSMSweep , typename Communication >
void walberla::lbm_mesapd_coupling::psm::gpu::addPSMSweepsToTimeloops ( SweepTimeloop & commTimeloop,
SweepTimeloop & timeloop,
Communication & comm,
SweepCollection & psmSweepCollection,
PSMSweep & psmSweep,
bool synchronize = true )

◆ boxMapping()

template<int Weighting_T>
__global__ void walberla::lbm_mesapd_coupling::psm::gpu::boxMapping ( walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField,
walberla::gpu::FieldAccessor< real_t > BsField,
walberla::gpu::FieldAccessor< id_t > idxField,
walberla::gpu::FieldAccessor< real_t > BField,
const real_t omega,
const double3 boxPositionMin,
const double3 boxPositionMax,
const double3 blockStart,
const real_t dx,
const id_t idxMapped )

◆ calculateWeighting()

template<int Weighting_T>
void walberla::lbm_mesapd_coupling::psm::gpu::calculateWeighting ( real_t * const ,
const real_t & ,
const real_t &  )
inline

◆ calculateWeighting< 1 >()

template<>
void walberla::lbm_mesapd_coupling::psm::gpu::calculateWeighting< 1 > ( real_t *const weighting,
const real_t & epsilon,
const real_t &  )
inline

◆ calculateWeighting< 2 >()

template<>
void walberla::lbm_mesapd_coupling::psm::gpu::calculateWeighting< 2 > ( real_t *const weighting,
const real_t & epsilon,
const real_t & tau )
inline

◆ cross()

__device__ void walberla::lbm_mesapd_coupling::psm::gpu::cross ( real_t *__restrict__ const crossResult,
const real_t *__restrict__ const lhs,
const real_t *__restrict__ const rhs )
inline

◆ getVelocityAtWFPoint()

__device__ void walberla::lbm_mesapd_coupling::psm::gpu::getVelocityAtWFPoint ( real_t *__restrict__ const velocityAtWFPoint,
const real_t *__restrict__ const linearVelocity,
const real_t *__restrict__ const angularVelocity,
const real_t *__restrict__ const position,
const real_t *__restrict__ const wf_pt )
inline

◆ linearApproximation()

template<int Weighting_T>
__global__ void walberla::lbm_mesapd_coupling::psm::gpu::linearApproximation ( walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField,
walberla::gpu::FieldAccessor< real_t > BsField,
walberla::gpu::FieldAccessor< id_t > idxField,
walberla::gpu::FieldAccessor< real_t > BField,
const real_t omega,
const real_t *__restrict__ const spherePositions,
const real_t *__restrict__ const sphereRadii,
const real_t *__restrict__ const f_rs,
const double3 blockStart,
const real_t dx,
const size_t *__restrict__ const numParticlesSubBlocks,
const size_t *__restrict__ const particleIDsSubBlocks,
const uint3 subBlocksPerDim )

◆ mapParticles() [1/2]

template<int Weighting_T>
void walberla::lbm_mesapd_coupling::psm::gpu::mapParticles ( const IBlock & blockIt,
const ParticleAndVolumeFractionSoA_T< Weighting_T > & particleAndVolumeFractionSoA,
const real_t *const spherePositions,
const real_t *const sphereRadii,
const real_t *const f_rs,
const size_t *const numParticlesSubBlocks,
const size_t *const particleIDsSubBlocks,
const Vector3< uint_t > subBlocksPerDim )

◆ mapParticles() [2/2]

template<int Weighting_T>
void walberla::lbm_mesapd_coupling::psm::gpu::mapParticles ( IBlock & blockIt,
const ParticleAndVolumeFractionSoA_T< Weighting_T > & particleAndVolumeFractionSoA,
const real_t *const spherePositions,
const real_t *const sphereRadii,
const real_t *const f_rs,
const size_t *const numParticlesSubBlocks,
const size_t *const particleIDsSubBlocks,
const Vector3< uint_t > subBlocksPerDim )

◆ ReduceParticleForces()

__global__ void walberla::lbm_mesapd_coupling::psm::gpu::ReduceParticleForces ( walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField,
walberla::gpu::FieldAccessor< id_t > idxField,
walberla::gpu::FieldAccessor< real_t > particleForcesField,
real_t *__restrict__ const hydrodynamicForces,
real_t *__restrict__ const hydrodynamicTorques,
real_t *__restrict__ const positions,
const double3 blockStart,
const real_t dx,
const real_t forceScalingFactor )

◆ SetParticleVelocities()

__global__ void walberla::lbm_mesapd_coupling::psm::gpu::SetParticleVelocities ( walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField,
walberla::gpu::FieldAccessor< uint_t > idxField,
walberla::gpu::FieldAccessor< real_t > particleVelocitiesField,
real_t *__restrict__ const linearVelocities,
real_t *__restrict__ const angularVelocities,
real_t *__restrict__ const positions,
const double3 blockStart,
const real_t dx )

◆ superSampling()

template<int Weighting_T>
__global__ void walberla::lbm_mesapd_coupling::psm::gpu::superSampling ( walberla::gpu::FieldAccessor< uint_t > nOverlappingParticlesField,
walberla::gpu::FieldAccessor< real_t > BsField,
walberla::gpu::FieldAccessor< id_t > idxField,
walberla::gpu::FieldAccessor< real_t > BField,
const real_t omega,
const real_t *__restrict__ const spherePositions,
const real_t *__restrict__ const sphereRadii,
const double3 blockStart,
const real_t dx,
const int3 nSamples,
const size_t *__restrict__ const numParticlesSubBlocks,
const size_t *__restrict__ const particleIDsSubBlocks,
const uint3 subBlocksPerDim )

Variable Documentation

◆ deviceSyncWrapper

auto walberla::lbm_mesapd_coupling::psm::gpu::deviceSyncWrapper
inline
Initial value:
= [](std::function< void(IBlock*) > sweep) {
return [sweep](IBlock* b) {
sweep(b);
};
}

◆ MaxParticlesPerCell

const uint_t walberla::lbm_mesapd_coupling::psm::gpu::MaxParticlesPerCell = MAX_PARTICLES_PER_CELL