olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS > Class Template Reference

Application of the collision step on a concrete CUDA block. More...

#include <operator.h>

Inheritance diagram for olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >:

Collaboration diagram for olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >:

Public Member Functions
	ConcreteBlockCollisionO ()
std::type_index	id () const override
std::size_t	weight () const override
	Returns number of assigned cells.
void	set (CellID iCell, bool state, bool overlap) override
	Set whether iCell is covered by the present collision step.
Dynamics< T, DESCRIPTOR > *	getDynamics () override
void	setup (ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &block) override
void	apply (ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &block, ConcreteBlockMask< T, Platform::GPU_CUDA > &subdomain, CollisionDispatchStrategy strategy) override
	Apply collision to subdomain of block using strategy.
Public Member Functions inherited from olb::BlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA >
virtual void	setup (ConcreteBlockLattice< T, DESCRIPTOR, PLATFORM > &block)=0
	Setup collision on block.
virtual void	apply (ConcreteBlockLattice< T, DESCRIPTOR, PLATFORM > &block, ConcreteBlockMask< T, PLATFORM > &subdomain, CollisionDispatchStrategy strategy)=0
	Apply collision on subdomain of block using strategy.
Public Member Functions inherited from olb::AbstractBlockO
virtual	~AbstractBlockO ()=default

Detailed Description

template<typename T, typename DESCRIPTOR, typename DYNAMICS>
class olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >

Application of the collision step on a concrete CUDA block.

Definition at line 41 of file operator.h.

Constructor & Destructor Documentation

ConcreteBlockCollisionO()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::ConcreteBlockCollisionO

(

)

Definition at line 467 of file operator.hh.

                                                                                        :
  _dynamics(new DYNAMICS()),
  _parameters(nullptr),
  _mask(nullptr),
  _cells(0),
  _modified(true),
  _stream(cudaStreamDefault)
{ }

Member Function Documentation

apply()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

void olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::apply ( ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > & block, ConcreteBlockMask< T, Platform::GPU_CUDA > & subdomain, CollisionDispatchStrategy strategy )

inlineoverride

Apply collision to subdomain of block using strategy.

The subdomain argument is currently assumed to be the core mask of BlockDynamicsMap

Definition at line 103 of file operator.h.

  {
    switch (strategy) {
    case CollisionDispatchStrategy::Dominant:
      return applyDominant(block, subdomain);
    case CollisionDispatchStrategy::Individual:
      return applyIndividual(block, subdomain);
    default:
      throw std::runtime_error("Invalid collision dispatch strategy");
    }
  }

References olb::Dominant, and olb::Individual.

getDynamics()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

Dynamics< T, DESCRIPTOR > * olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::getDynamics ( )

inlineoverridevirtual

Implements olb::AbstractCollisionO< T, DESCRIPTOR >.

Definition at line 92 of file operator.h.

  {
    return _dynamics.get();
  }

id()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

std::type_index olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::id ( ) const

inlineoverridevirtual

Implements olb::AbstractBlockO.

Definition at line 68 of file operator.h.

  {
    return typeid(DYNAMICS);
  }

set()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

void olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::set ( CellID iCell, bool state, bool overlap )

inlineoverridevirtual

Set whether iCell is covered by the present collision step.

Parameters

iCell	Cell index
state	(De)activate for this dynamics / collision
overlap	Cell index in overlap (set dynamics but do not collide)

Only unmask cells that actually do something

Implements olb::AbstractCollisionO< T, DESCRIPTOR >.

Definition at line 78 of file operator.h.

  {
    if constexpr (!std::is_same_v<DYNAMICS,NoDynamics<T,DESCRIPTOR>>) {
      if (!overlap) {
        _mask->set(iCell, state);
      }
    }
    if (state) {
      _dynamicsOfCells[iCell] = _deviceDynamics.get();
    }
    _modified = true;
  }

References olb::gpu::cuda::device::unique_ptr< T >::get().

Here is the call graph for this function:

setup()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

void olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::setup ( ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > & block )

override

Definition at line 552 of file operator.hh.

{
  // Fetch pointers to DYNAMICS-specific parameter and mask data
  _parameters = &block.template getData<OperatorParameters<DYNAMICS>>();
  _mask = &block.template getData<DynamicsMask<DYNAMICS>>();
 
  {
    // Construct on-device dynamics proxy for dynamic dispatch
    _deviceDynamics = gpu::cuda::device::malloc<gpu::cuda::ConcreteDynamics<T,DESCRIPTOR,DYNAMICS>>(1);
    gpu::cuda::kernel::construct_dynamics<T,DESCRIPTOR,DYNAMICS><<<1,1>>>(
      _deviceDynamics.get(),
      _parameters->deviceData());
    gpu::cuda::device::check();
 
    // Fetch pointer to on-device dynamic-dispatch field
    _dynamicsOfCells = block.template getField<gpu::cuda::DYNAMICS<T,DESCRIPTOR>>()[0].data();
  }
}

References olb::gpu::cuda::device::check().

Here is the call graph for this function:

weight()

template<typename T , typename DESCRIPTOR , typename DYNAMICS >

std::size_t olb::ConcreteBlockCollisionO< T, DESCRIPTOR, Platform::GPU_CUDA, DYNAMICS >::weight ( ) const

inlineoverridevirtual

Returns number of assigned cells.

Used to determine the dominant dynamics to choose e.g. which collision operator to vectorize or to prefer in GPU kernels.

Implements olb::AbstractCollisionO< T, DESCRIPTOR >.

Definition at line 73 of file operator.h.

  {
    return _mask->weight();
  }

---

The documentation for this class was generated from the following files:

• src/core/platform/gpu/cuda/operator.h
• src/core/platform/gpu/cuda/operator.hh