OpenLB 1.8.1
Loading...
Searching...
No Matches
Public Member Functions | Protected Member Functions | List of all members
olb::gpu::cuda::Cell< T, DESCRIPTOR > Class Template Referencefinal

Device-side implementation of the Cell concept for post processors. More...

#include <context.hh>

+ Inheritance diagram for olb::gpu::cuda::Cell< T, DESCRIPTOR >:
+ Collaboration diagram for olb::gpu::cuda::Cell< T, DESCRIPTOR >:

Public Member Functions

 Cell (DeviceBlockLattice< T, DESCRIPTOR > &data, CellID iCell) __device__
 
Cell< T, DESCRIPTOR > neighbor (LatticeR< DESCRIPTOR::d > offset) __device__
 
Dynamics< T, DESCRIPTOR > & getDynamics () __device__
 
std::size_t getCellId () const __device__
 
void setCellId (std::size_t iCell) __device__
 
void setLatticeR (LatticeR< DESCRIPTOR::d > latticeR) __device__
 
bool isPadding () const __device__
 
T computeRho () __device__
 
void computeU (T *u) __device__
 
void computeJ (T *j) __device__
 
void computeRhoU (T &rho, T *u) __device__
 
void computeStress (T *pi) __device__
 
void computeAllMomenta (T &rho, T *u, T *pi) __device__
 
void defineRho (T &rho) __device__
 
void defineU (T *u) __device__
 
void defineRhoU (T &rho, T *u) __device__
 
void defineAllMomenta (T &rho, T *u, T *pi) __device__
 
void definePopulations (const T *f) __device__
 
void iniEquilibrium (T rho, T *u) __device__
 
void iniRegularized (T rho, T *u, T *pi) __device__
 
void inverseShiftRhoU (T &rho, T *u) __device__
 
- Public Member Functions inherited from olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >
 DataOnlyCell (DeviceContext< T, DESCRIPTOR > &data, CellID iCell) __device__
 
void setCellId (std::size_t iCell) __device__
 
template<typename FIELD >
FIELD::template value_type< T > getFieldComponent (unsigned iD) __device__
 
value_t & operator[] (int iPop) __device__
 
template<typename FIELD >
auto getField () const __device__
 
template<typename FIELD >
FieldPtr< T, DESCRIPTOR, FIELD > getFieldPointer () __device__
 
template<typename FIELD >
auto getFieldComponent (unsigned iD) const __device__
 
template<typename FIELD >
void setField (FieldD< value_t, descriptor_t, FIELD > &&value) __device__
 
template<typename FIELD >
void setField (const FieldD< value_t, descriptor_t, FIELD > &value) __device__
 

Protected Member Functions

DeviceBlockLattice< T, DESCRIPTOR > & getBlock () __device__
 
DeviceBlockLattice< T, DESCRIPTOR > & getBlock () const __device__
 

Additional Inherited Members

- Public Types inherited from olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >
using value_t = T
 
using descriptor_t = DESCRIPTOR
 
- Protected Attributes inherited from olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >
DeviceContext< T, DESCRIPTOR > & _data
 
CellID _iCell
 

Detailed Description

template<typename T, typename DESCRIPTOR>
class olb::gpu::cuda::Cell< T, DESCRIPTOR >

Device-side implementation of the Cell concept for post processors.

Adds neighborhood and dynamically-dispatched momenta access to DataOnlyCell

Definition at line 313 of file context.hh.

Constructor & Destructor Documentation

◆ Cell()

template<typename T , typename DESCRIPTOR >
olb::gpu::cuda::Cell< T, DESCRIPTOR >::Cell ( DeviceBlockLattice< T, DESCRIPTOR > & data,
CellID iCell )
inline

Definition at line 324 of file context.hh.

324 :
325 DataOnlyCell<T,DESCRIPTOR>(data, iCell)
326 { }
olb::gpu::cuda::DataOnlyCell::DataOnlyCell
DataOnlyCell(DeviceContext< T, DESCRIPTOR > &data, CellID iCell) __device__
Definition context.hh:176

Member Function Documentation

◆ computeAllMomenta()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeAllMomenta ( T & rho,
T * u,
T * pi )
inline

Definition at line 369 of file context.hh.

369 {
370 getDynamics().computeAllMomenta(*this, rho, u, pi);
371 }
olb::gpu::cuda::Cell::getDynamics
Dynamics< T, DESCRIPTOR > & getDynamics() __device__
Definition context.hh:332

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ computeJ()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeJ ( T * j)
inline

Definition at line 358 of file context.hh.

358 {
359 getDynamics().computeJ(*this, j);
360 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ computeRho()

template<typename T , typename DESCRIPTOR >
T olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeRho ( )
inline

Definition at line 352 of file context.hh.

352 {
353 return getDynamics().computeRho(*this);
354 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ computeRhoU()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeRhoU ( T & rho,
T * u )
inline

Definition at line 361 of file context.hh.

361 {
362 getDynamics().computeRhoU(*this, rho, u);
363 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ computeStress()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeStress ( T * pi)
inline

Definition at line 364 of file context.hh.

364 {
365 T rho, u[DESCRIPTOR::d] { };
366 getDynamics().computeRhoU(*this, rho, u);
367 getDynamics().computeStress(*this, rho, u, pi);
368 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ computeU()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::computeU ( T * u)
inline

Definition at line 355 of file context.hh.

355 {
356 getDynamics().computeU(*this, u);
357 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ defineAllMomenta()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::defineAllMomenta ( T & rho,
T * u,
T * pi )
inline

Definition at line 382 of file context.hh.

382 {
383 return getDynamics().defineAllMomenta(*this, rho, u, pi);
384 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ definePopulations()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::definePopulations ( const T * f)
inline

Definition at line 385 of file context.hh.

385 {
386 for (int iPop=0; iPop < descriptors::q<DESCRIPTOR>(); ++iPop) {
387 this->operator[](iPop) = f[iPop];
388 }
389 }
olb::gpu::cuda::DataOnlyCell::operator[]
value_t & operator[](int iPop) __device__
Definition context.hh:190

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::operator[]().

+ Here is the call graph for this function:

◆ defineRho()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::defineRho ( T & rho)
inline

Definition at line 373 of file context.hh.

373 {
374 return getDynamics().defineRho(*this, rho);
375 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ defineRhoU()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::defineRhoU ( T & rho,
T * u )
inline

Definition at line 379 of file context.hh.

379 {
380 return getDynamics().defineRhoU(*this, rho, u);
381 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ defineU()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::defineU ( T * u)
inline

Definition at line 376 of file context.hh.

376 {
377 return getDynamics().defineU(*this, u);
378 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ getBlock() [1/2]

template<typename T , typename DESCRIPTOR >
DeviceBlockLattice< T, DESCRIPTOR > & olb::gpu::cuda::Cell< T, DESCRIPTOR >::getBlock ( )
inlineprotected

Definition at line 315 of file context.hh.

315 {
316 return static_cast<DeviceBlockLattice<T,DESCRIPTOR>&>(this->_data);
317 }
olb::gpu::cuda::DataOnlyCell::_data
DeviceContext< T, DESCRIPTOR > & _data
Definition context.hh:169

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_data.

+ Here is the caller graph for this function:

◆ getBlock() [2/2]

template<typename T , typename DESCRIPTOR >
DeviceBlockLattice< T, DESCRIPTOR > & olb::gpu::cuda::Cell< T, DESCRIPTOR >::getBlock ( ) const
inlineprotected

Definition at line 319 of file context.hh.

319 {
320 return static_cast<DeviceBlockLattice<T,DESCRIPTOR>&>(this->_data);
321 }

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_data.

◆ getCellId()

template<typename T , typename DESCRIPTOR >
std::size_t olb::gpu::cuda::Cell< T, DESCRIPTOR >::getCellId ( ) const
inline

Definition at line 336 of file context.hh.

336 {
337 return this->_iCell;
338 }

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_iCell.

◆ getDynamics()

template<typename T , typename DESCRIPTOR >
Dynamics< T, DESCRIPTOR > & olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics ( )
inline

Definition at line 332 of file context.hh.

332 {
333 return *this->template getField<DYNAMICS<T,DESCRIPTOR>>();
334 }
olb::gpu::cuda::DataOnlyCell::getField
auto getField() const __device__
Definition context.hh:195

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::getField().

+ Here is the call graph for this function:
+ Here is the caller graph for this function:

◆ iniEquilibrium()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::iniEquilibrium ( T rho,
T * u )
inline

Definition at line 391 of file context.hh.

391 {
392 getDynamics().iniEquilibrium(*this, rho, u);
393 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ iniRegularized()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::iniRegularized ( T rho,
T * u,
T * pi )
inline

Definition at line 394 of file context.hh.

394 {
395 getDynamics().iniRegularized(*this, rho, u, pi);
396 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ inverseShiftRhoU()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::inverseShiftRhoU ( T & rho,
T * u )
inline

Definition at line 397 of file context.hh.

397 {
398 getDynamics().inverseShiftRhoU(*this, rho, u);
399 }

References olb::gpu::cuda::Cell< T, DESCRIPTOR >::getDynamics().

+ Here is the call graph for this function:

◆ isPadding()

template<typename T , typename DESCRIPTOR >
bool olb::gpu::cuda::Cell< T, DESCRIPTOR >::isPadding ( ) const
inline

Definition at line 348 of file context.hh.

348 {
349 return getBlock().isPadding(this->_iCell);
350 }
olb::gpu::cuda::Cell::getBlock
DeviceBlockLattice< T, DESCRIPTOR > & getBlock() __device__
Definition context.hh:315

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_iCell, and olb::gpu::cuda::Cell< T, DESCRIPTOR >::getBlock().

+ Here is the call graph for this function:

◆ neighbor()

template<typename T , typename DESCRIPTOR >
Cell< T, DESCRIPTOR > olb::gpu::cuda::Cell< T, DESCRIPTOR >::neighbor ( LatticeR< DESCRIPTOR::d > offset)
inline

Definition at line 328 of file context.hh.

328 {
329 return {getBlock(), CellID(this->_iCell + getBlock().getNeighborDistance(offset))};
330 }
olb::CellID
std::uint32_t CellID
Type for sequential block-local cell indices.
Definition blockStructure.h:36

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_iCell, and olb::gpu::cuda::Cell< T, DESCRIPTOR >::getBlock().

+ Here is the call graph for this function:

◆ setCellId()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::setCellId ( std::size_t iCell)
inline

Definition at line 340 of file context.hh.

340 {
341 this->_iCell = iCell;
342 }

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_iCell.

◆ setLatticeR()

template<typename T , typename DESCRIPTOR >
void olb::gpu::cuda::Cell< T, DESCRIPTOR >::setLatticeR ( LatticeR< DESCRIPTOR::d > latticeR)
inline

Definition at line 344 of file context.hh.

344 {
345 this->_iCell = getBlock().getCellId(latticeR);
346 }

References olb::gpu::cuda::DataOnlyCell< T, DESCRIPTOR >::_iCell, and olb::gpu::cuda::Cell< T, DESCRIPTOR >::getBlock().

+ Here is the call graph for this function:
The documentation for this class was generated from the following file:
Generated on Sat May 17 2025 10:48:15 for OpenLB by doxygen 1.12.0