OpenLB 1.7
Loading...
Searching...
No Matches
Public Member Functions | List of all members
olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE > Class Template Reference

Private implementation of heterogeneous copy task between CPU_* source and GPU_CUDA target. More...

#include <communicator.hh>

+ Inheritance diagram for olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >:
+ Collaboration diagram for olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >:

Public Member Functions

 HeterogeneousCopyTaskDataForGpuTarget (const std::vector< std::type_index > &fields, const std::vector< CellID > &targetCells, ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &target, const std::vector< CellID > &sourceCells, ConcreteBlockLattice< T, DESCRIPTOR, SOURCE > &source)
 
void copy () override
 
void wait () override
 
- Public Member Functions inherited from olb::ConcreteHeterogeneousCopyTask
virtual ~ConcreteHeterogeneousCopyTask ()
 

Detailed Description

template<typename T, typename DESCRIPTOR, Platform SOURCE>
class olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >

Private implementation of heterogeneous copy task between CPU_* source and GPU_CUDA target.

Definition at line 590 of file communicator.hh.

Constructor & Destructor Documentation

◆ HeterogeneousCopyTaskDataForGpuTarget()

template<typename T , typename DESCRIPTOR , Platform SOURCE>
olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >::HeterogeneousCopyTaskDataForGpuTarget ( const std::vector< std::type_index > & fields,
const std::vector< CellID > & targetCells,
ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > & target,
const std::vector< CellID > & sourceCells,
ConcreteBlockLattice< T, DESCRIPTOR, SOURCE > & source )
inline

Definition at line 608 of file communicator.hh.

611 :
612 _targetFields(target.getDataRegistry().deviceFieldArrays(fields)),
613 _onlyPopulationField(fields.size() == 1 && fields[0] == typeid(descriptors::POPULATION)),
614 _targetCells(targetCells),
615 _sourceCells(sourceCells),
616 _target(target),
617 _source(source, fields),
618 _stream(std::make_unique<gpu::cuda::device::Stream>(cudaStreamNonBlocking)),
619 _buffer(_source.size(_sourceCells))
620 { }
auto & getDataRegistry()
Return reference to Data's FieldTypeRegistry.

Member Function Documentation

◆ copy()

template<typename T , typename DESCRIPTOR , Platform SOURCE>
void olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >::copy ( )
inlineoverridevirtual

Implements olb::ConcreteHeterogeneousCopyTask.

Definition at line 622 of file communicator.hh.

622 {
623 _source.serialize(_sourceCells, _buffer.data());
625
626 if (_onlyPopulationField) {
627 gpu::cuda::DeviceContext<T,DESCRIPTOR> lattice(_target);
628 gpu::cuda::async_scatter_field<descriptors::POPULATION>(_stream->get(), lattice, _targetCells, _buffer.deviceData());
629 } else {
630 gpu::cuda::async_scatter_any_fields(_stream->get(), _targetFields, _targetCells, _buffer.deviceData());
631 }
632 }
const T * deviceData() const
Definition column.hh:146
const T * data() const
Definition column.hh:134
void setProcessingContext(ProcessingContext)
Definition column.hh:158
void async_scatter_any_fields(cudaStream_t stream, thrust::device_vector< AnyDeviceFieldArrayD > &fields, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
Non-blocking scatter of fields data in buffer to given indices.
@ Simulation
Data available on host for e.g. functor evaluation.

References olb::gpu::cuda::async_scatter_any_fields(), olb::gpu::cuda::Column< T >::data(), olb::gpu::cuda::Column< T >::deviceData(), olb::MultiConcreteCommunicatable< COMMUNICATEE >::serialize(), olb::gpu::cuda::Column< T >::setProcessingContext(), and olb::Simulation.

+ Here is the call graph for this function:

◆ wait()

template<typename T , typename DESCRIPTOR , Platform SOURCE>
void olb::HeterogeneousCopyTaskDataForGpuTarget< T, DESCRIPTOR, SOURCE >::wait ( )
inlineoverridevirtual

Implements olb::ConcreteHeterogeneousCopyTask.

Definition at line 634 of file communicator.hh.

634 {
635 _stream->synchronize();
636 }

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