Implementations of Nvidia CUDA specifics. More...

Namespaces
namespace	concepts

namespace	device
	Basic wrappers of common CUDA functions.

namespace	kernel
	CUDA kernels to execute collisions and post processors.

Classes
struct	AnyDeviceFieldArrayD
	Type-erased pointer to FieldArrayD device data. More...

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

class	Column
	Plain column for CUDA GPU targets. More...

class	ConcreteDynamics
	Implementation of gpu::cuda::Dynamics for concrete DYNAMICS. More...

class	CyclicColumn
	Virtual memory based cyclic column for usage in ColumnVector. More...

class	DataOnlyCell
	Device-side implementation of the data-only Cell concept for collision steps. More...

class	DeviceBlockLattice
	Device-side view of a block lattice. More...

class	DeviceContext
	Structure for passing pointers to on-device data into CUDA kernels. More...

struct	DynamicDispatchCollision
	Last node in a MaskedDynamics chain in kernel::call_operators. More...

struct	DYNAMICS
	On-device field mirroring BlockDynamicsMap. More...

struct	Dynamics
	Virtual interface for device-side dynamically-dispatched dynamics access. More...

struct	FieldArrayPointer
	Host-side version of gpu::cuda::AnyDeviceFieldArrayD. More...

class	FieldPtr
	Pointer to row of a D-dimensional field. More...

struct	first_tuple_element_non_zero

class	ListedCollision
	List-based application of DYNAMICS::collide for use in kernel::call_list_operators. More...

struct	ListedPostProcessor
	List-based application of OPERATOR::apply. More...

class	ListedPostProcessorWithParameters
	List-based application of OPERATOR::apply with parameters. More...

class	MaskedCollision
	Masked application of DYNAMICS::collide for use in kernel::call_operators. More...

class	MaskedPostProcessor
	Masked application of OPERATOR::apply. More...

struct	maximum_and_plus
	Function object for simulateneously computing maximum and sum in a single thrust::reduce. More...

struct	pair
	Plain pair type with single-value constructor for use in gpu::cuda::maximum_and_plus. More...

class	ParticleCouplingWithParameters
	Unrestricted application of COUPLING::apply with parameters. More...

struct	SharedMemory

struct	SharedMemory< double >

struct	sum_of_thrust_tuple

struct	UnmaskedCoupling
	Unrestricted application of COUPLING::apply. More...

class	UnmaskedCouplingWithParameters
	Unrestricted application of COUPLING::apply with parameters. More...

Functions
template<typename FIELD , typename CONTEXT >
void	gather_field (CONTEXT &lattice, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Blocking gather of FIELD at given indices into buffer.

template<typename FIELD , typename CONTEXT >
void	async_gather_field (cudaStream_t stream, CONTEXT &lattice, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Non-blocking gather of FIELD at given indices into buffer.

template<typename FIELD , typename SOURCE , typename TARGET >
void	async_copy_field (cudaStream_t stream, SOURCE &sourceLattice, TARGET &targetLattice, const thrust::device_vector< CellID > &sourceIndices, const thrust::device_vector< CellID > &targetIndices)
	Non-blocking copy of FIELD at given indices from sourceLattice to targetLattice.

void	gather_any_fields (thrust::device_vector< AnyDeviceFieldArrayD > &fields, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Blocking gather of fields at given indices into buffer.

void	async_gather_any_fields (cudaStream_t stream, thrust::device_vector< AnyDeviceFieldArrayD > &fields, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Non-blocking gather of fields at given indices into buffer.

void	async_copy_any_fields (cudaStream_t stream, thrust::device_vector< AnyDeviceFieldArrayD > &sourceFields, thrust::device_vector< AnyDeviceFieldArrayD > &targetFields, const thrust::device_vector< CellID > &sourceIndices, const thrust::device_vector< CellID > &targetIndices)
	Non-blocking copy of fields at given indices from sourceIndices to targetIndices.

template<typename FIELD , typename CONTEXT >
void	scatter_field (CONTEXT &lattice, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Blocking scatter of FIELD data in buffer to given indices.

template<typename FIELD , typename CONTEXT >
void	async_scatter_field (cudaStream_t stream, CONTEXT &lattice, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Non-blocking scatter of FIELD data in buffer to given indices.

void	scatter_any_fields (thrust::device_vector< AnyDeviceFieldArrayD > &fields, const thrust::device_vector< CellID > &indices, std::uint8_t *buffer)
	Blocking scatter of fields data in buffer to given indices.

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.

template<typename T , typename DESCRIPTOR , typename... DYNAMICS>
std::function< void(ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &)>	getFusedCollisionO ()
	Helper for constructing fused collision operators.

template<typename CONTEXT , typename... ARGS>
void	call_operators (CONTEXT &lattice, bool *subdomain, ARGS &&... args)
	Apply masked collision operators to lattice.

template<typename CONTEXT , typename... ARGS>
void	async_call_operators (cudaStream_t stream, CONTEXT &lattice, bool *subdomain, ARGS &&... args)
	Apply masked collision operators to lattice (async)

template<typename CONTEXT , typename... ARGS>
void	call_operators_with_statistics (CONTEXT &lattice, bool *subdomain, ARGS &&... args)
	Apply masked collision operators to lattice while tracking statistics.

template<typename CONTEXT , typename... ARGS>
void	async_call_operators_with_statistics (cudaStream_t stream, CONTEXT &lattice, bool *subdomain, ARGS &&... args)
	Apply masked collision operators to lattice while tracking statistics (async)

template<typename CONTEXT , typename... ARGS>
void	call_list_operators (CONTEXT &lattice, const gpu::cuda::Column< CellID > &cells, ARGS &&... args)
	Apply operators to listed cell indices.

template<typename CONTEXT , typename... ARGS>
void	async_call_list_operators (cudaStream_t stream, CONTEXT &lattice, const gpu::cuda::Column< CellID > &cells, ARGS &&... args)
	Apply operators to listed cell indices (async version)

template<typename CONTEXT , typename... ARGS>
void	async_call_list_operators_with_statistics (cudaStream_t stream, CONTEXT &lattice, const gpu::cuda::Column< CellID > &cells, ARGS &&... args)
	Apply ListedCollision with statistics (async version)

template<typename CONTEXT , typename... ARGS>
void	call_coupling_operators (CONTEXT &lattices, bool *subdomain, ARGS &&... args)
	Apply coupling on subdomain.

template<typename CONTEXT , typename... ARGS>
void	call_particle_coupling_operators (CONTEXT &lattices, ARGS &&... args)
	Apply coupling on particles and lattices.

template<typename F , std::size_t... INDICES>
auto	make_thrust_tuple_f (F &&f, std::index_sequence< INDICES... >)

template<unsigned D, typename F >
auto	make_thrust_tuple_f (F &&f)

template<typename A , std::size_t... AIs>
auto	thrust_tuple_cat (A &a, std::index_sequence< AIs... >)

template<typename A , typename B , std::size_t... AIs, std::size_t... BIs>
auto	thrust_tuple_cat (A &a, B &b, std::index_sequence< AIs... >, std::index_sequence< BIs... >)

template<typename A , typename B , typename C , std::size_t... AIs, std::size_t... BIs, std::size_t... CIs>
auto	thrust_tuple_cat (A &a, B &b, C &c, std::index_sequence< AIs... >, std::index_sequence< BIs... >, std::index_sequence< CIs... >)

template<typename... ARGS>
auto	thrust_tuple_cat (ARGS &&... args)

template<typename T , unsigned D>
auto	make_thrust_tuple_of (Vector< T, D > &f)

template<typename T >
auto	make_thrust_tuple_of (T &f)

template<typename VECTOR >
auto	make_thrust_tuple_of_device_data (VECTOR &f)

template<typename COARSE , typename FINE , typename CONTEXT , typename PARAMETERS , typename OPERATOR >
void	call_refinement_coupling_operator (COARSE &cLattice, FINE &fLattice, CONTEXT &data, PARAMETERS &parameters, meta::id< OPERATOR >)

Variables
template<typename CONTEXT , typename TYPE >
__constant__ std::size_t	field_type_index
	Mapping of TYPE in CONTEXT to runtime-fixed index.

Detailed Description

Implementations of Nvidia CUDA specifics.

Function Documentation

◆ async_call_list_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::async_call_list_operators	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		const gpu::cuda::Column< CellID > &	cells,
		ARGS &&...	args )

Apply operators to listed cell indices (async version)

Definition at line 483 of file operator.hh.

                                               {
  const auto block_size = 32;
  const auto block_count = (cells.size() + block_size - 1) / block_size;
  kernel::call_list_operators<<<block_count,block_size,0,stream>>>(
    lattice,
    cells.deviceData(), cells.size(),
    std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_list_operators(), olb::gpu::cuda::device::check(), olb::gpu::cuda::Column< T >::deviceData(), and olb::gpu::cuda::Column< T >::size().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_call_list_operators_with_statistics()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::async_call_list_operators_with_statistics	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		const gpu::cuda::Column< CellID > &	cells,
		ARGS &&...	args )

Apply ListedCollision with statistics (async version)

Definition at line 498 of file operator.hh.

                                                               {
  const auto block_size = 32;
  const auto block_count = (cells.size() + block_size - 1) / block_size;
  kernel::call_list_operators_with_statistics<<<block_count,block_size,0,stream>>>(
    lattice,
    cells.deviceData(), cells.size(),
    std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_list_operators_with_statistics(), olb::gpu::cuda::device::check(), olb::gpu::cuda::Column< T >::deviceData(), and olb::gpu::cuda::Column< T >::size().

Here is the call graph for this function:

◆ async_call_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::async_call_operators	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		bool *	subdomain,
		ARGS &&...	args )

Apply masked collision operators to lattice (async)

Definition at line 433 of file operator.hh.

                                                                                                  {
  const auto block_size = 32;
  const auto block_count = (lattice.getNcells() + block_size - 1) / block_size;
  kernel::call_operators<CONTEXT,ARGS...><<<block_count,block_size,0,stream>>>(
    lattice, subdomain, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_operators(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

◆ async_call_operators_with_statistics()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::async_call_operators_with_statistics	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		bool *	subdomain,
		ARGS &&...	args )

Apply masked collision operators to lattice while tracking statistics (async)

Definition at line 456 of file operator.hh.

                                                                                                                  {
  const auto block_size = 32;
  const auto block_count = (lattice.getNcells() + block_size - 1) / block_size;
  kernel::call_operators_with_statistics<CONTEXT,ARGS...><<<block_count,block_size,0,stream>>>(
    lattice, subdomain, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_operators_with_statistics(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

◆ async_copy_any_fields()

void olb::gpu::cuda::async_copy_any_fields	(	cudaStream_t	stream,
		thrust::device_vector< AnyDeviceFieldArrayD > &	sourceFields,
		thrust::device_vector< AnyDeviceFieldArrayD > &	targetFields,
		const thrust::device_vector< CellID > &	sourceIndices,
		const thrust::device_vector< CellID > &	targetIndices )

Non-blocking copy of fields at given indices from sourceIndices to targetIndices.

Definition at line 276 of file communicator.hh.

                                                                             {
  const auto block_size = 32;
  const auto block_count = (sourceIndices.size() + block_size - 1) / block_size;
  kernel::copy_any_fields<<<block_count,block_size,0,stream>>>(
    sourceFields.data().get(),  targetFields.data().get(),  sourceFields.size(),
    sourceIndices.data().get(), targetIndices.data().get(), sourceIndices.size());
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::copy_any_fields().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_copy_field()

template<typename FIELD , typename SOURCE , typename TARGET >

void olb::gpu::cuda::async_copy_field	(	cudaStream_t	stream,
		SOURCE &	sourceLattice,
		TARGET &	targetLattice,
		const thrust::device_vector< CellID > &	sourceIndices,
		const thrust::device_vector< CellID > &	targetIndices )

Non-blocking copy of FIELD at given indices from sourceLattice to targetLattice.

Definition at line 231 of file communicator.hh.

                                                                        {
  const auto block_size = 32;
  const auto block_count = (sourceIndices.size() + block_size - 1) / block_size;
  kernel::copy_field<SOURCE,TARGET,FIELD><<<block_count,block_size,0,stream>>>(
    sourceLattice,
    targetLattice,
    sourceIndices.data().get(),
    targetIndices.data().get(),
    sourceIndices.size());
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::copy_field().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_gather_any_fields()

void olb::gpu::cuda::async_gather_any_fields	(	cudaStream_t	stream,
		thrust::device_vector< AnyDeviceFieldArrayD > &	fields,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Non-blocking gather of fields at given indices into buffer.

Definition at line 262 of file communicator.hh.

                                                 {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::gather_any_fields<<<block_count,block_size,0,stream>>>(
    fields.data().get(), fields.size(),
    indices.data().get(), indices.size(),
    buffer);
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::gather_any_fields().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_gather_field()

template<typename FIELD , typename CONTEXT >

void olb::gpu::cuda::async_gather_field	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Non-blocking gather of FIELD at given indices into buffer.

Definition at line 216 of file communicator.hh.

                                            {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::gather_field<CONTEXT,FIELD><<<block_count,block_size,0,stream>>>(
    lattice,
    indices.data().get(), indices.size(),
    reinterpret_cast<typename FIELD::template value_type<typename CONTEXT::value_t>*>(buffer));
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::gather_field().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_scatter_any_fields()

void olb::gpu::cuda::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.

Definition at line 330 of file communicator.hh.

                                                  {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::scatter_any_fields<<<block_count,block_size,0,stream>>>(
    fields.data().get(), fields.size(),
    indices.data().get(), indices.size(),
    buffer);
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::scatter_any_fields().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ async_scatter_field()

template<typename FIELD , typename CONTEXT >

void olb::gpu::cuda::async_scatter_field	(	cudaStream_t	stream,
		CONTEXT &	lattice,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Non-blocking scatter of FIELD data in buffer to given indices.

Definition at line 303 of file communicator.hh.

                                             {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::scatter_field<CONTEXT,FIELD><<<block_count,block_size,0,stream>>>(
    lattice,
    indices.data().get(), indices.size(),
    reinterpret_cast<typename FIELD::template value_type<typename CONTEXT::value_t>*>(buffer));
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::scatter_field().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ call_coupling_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::call_coupling_operators	(	CONTEXT &	lattices,
		bool *	subdomain,
		ARGS &&...	args )

Apply coupling on subdomain.

Definition at line 513 of file operator.hh.

                                                                                 {
  const auto nCells = lattices.template get<0>().getNcells();
  const auto block_size = 32;
  const auto block_count = (nCells + block_size - 1) / block_size;
  kernel::call_coupling_operators<CONTEXT,ARGS...><<<block_count,block_size>>>(
    lattices, subdomain, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_coupling_operators(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ call_list_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::call_list_operators	(	CONTEXT &	lattice,
		const gpu::cuda::Column< CellID > &	cells,
		ARGS &&...	args )

Apply operators to listed cell indices.

Used to call post processors in ConcreteBlockO with OperatorScope::PerCell

Definition at line 469 of file operator.hh.

                                         {
  const auto block_size = 32;
  const auto block_count = (cells.size() + block_size - 1) / block_size;
  kernel::call_list_operators<CONTEXT,ARGS...><<<block_count, block_size>>>(
    lattice,
    cells.deviceData(), cells.size(),
    std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_list_operators(), olb::gpu::cuda::device::check(), olb::gpu::cuda::Column< T >::deviceData(), and olb::gpu::cuda::Column< T >::size().

Here is the call graph for this function:

◆ call_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::call_operators	(	CONTEXT &	lattice,
		bool *	subdomain,
		ARGS &&...	args )

Apply masked collision operators to lattice.

ARGS are instances of MaskedCollision or DynamicDispatchCollision

Definition at line 423 of file operator.hh.

                                                                       {
  const auto block_size = 32;
  const auto block_count = (lattice.getNcells() + block_size - 1) / block_size;
  kernel::call_operators<CONTEXT,ARGS...><<<block_count,block_size>>>(
    lattice, subdomain, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_operators(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ call_operators_with_statistics()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::call_operators_with_statistics	(	CONTEXT &	lattice,
		bool *	subdomain,
		ARGS &&...	args )

Apply masked collision operators to lattice while tracking statistics.

ARGS are instances of MaskedCollision or DynamicDispatchCollision

Definition at line 446 of file operator.hh.

                                                                                       {
  const auto block_size = 32;
  const auto block_count = (lattice.getNcells() + block_size - 1) / block_size;
  kernel::call_operators_with_statistics<CONTEXT,ARGS...><<<block_count,block_size>>>(
    lattice, subdomain, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_operators_with_statistics(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ call_particle_coupling_operators()

template<typename CONTEXT , typename... ARGS>

void olb::gpu::cuda::call_particle_coupling_operators	(	CONTEXT &	lattices,
		ARGS &&...	args )

Apply coupling on particles and lattices.

Definition at line 524 of file operator.hh.

                                                                         {
  const auto nParticles = lattices.get(meta::id<names::Points>{}).getNcells();
  const auto block_size = 32;
  const auto block_count = (nParticles + block_size - 1) / block_size;
  kernel::call_particle_coupling_operators<CONTEXT,ARGS...><<<block_count,block_size>>>(
    lattices, nParticles, std::forward<decltype(args)>(args)...);
  device::check();
}

References olb::gpu::cuda::kernel::call_particle_coupling_operators(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ call_refinement_coupling_operator()

template<typename COARSE , typename FINE , typename CONTEXT , typename PARAMETERS , typename OPERATOR >

void olb::gpu::cuda::call_refinement_coupling_operator	(	COARSE &	cLattice,
		FINE &	fLattice,
		CONTEXT &	data,
		PARAMETERS &	parameters,
		meta::id< OPERATOR >	)

Definition at line 77 of file operator.hh.

                                                         {
  const std::size_t n = data.getNcells();
  if (n > 0) {
    const auto block_size = 32;
    const auto block_count = (n + block_size - 1) / block_size;
    kernel::call_refinement_coupling_operator<COARSE,FINE,CONTEXT,PARAMETERS,OPERATOR>
      <<<block_count,block_size>>>(cLattice, fLattice, data, parameters, n);
    device::check();
  }
}

References olb::gpu::cuda::kernel::call_refinement_coupling_operator(), and olb::gpu::cuda::device::check().

Here is the call graph for this function:

◆ gather_any_fields()

void olb::gpu::cuda::gather_any_fields	(	thrust::device_vector< AnyDeviceFieldArrayD > &	fields,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Blocking gather of fields at given indices into buffer.

Definition at line 249 of file communicator.hh.

                                           {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::gather_any_fields<<<block_count,block_size>>>(
    fields.data().get(), fields.size(),
    indices.data().get(), indices.size(),
    buffer);
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::gather_any_fields().

Here is the call graph for this function:

◆ gather_field()

template<typename FIELD , typename CONTEXT >

void olb::gpu::cuda::gather_field	(	CONTEXT &	lattice,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Blocking gather of FIELD at given indices into buffer.

Definition at line 204 of file communicator.hh.

                                                                                                    {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::gather_field<CONTEXT,FIELD><<<block_count,block_size>>>(
    lattice,
    indices.data().get(), indices.size(),
    reinterpret_cast<typename FIELD::template value_type<typename CONTEXT::value_t>*>(buffer));
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::gather_field().

Here is the call graph for this function:

◆ getFusedCollisionO()

template<typename T , typename DESCRIPTOR , typename... DYNAMICS>

std::function< void(ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &)> olb::gpu::cuda::getFusedCollisionO ( )

Helper for constructing fused collision operators.

This is a convenient way for potentially improving performance by injecting application knowledge. E.g. if the lattice contains primarily BGK and BounceBack dynamics this can be declared using:

superLattice.forBlocksOnPlatform<Platform::GPU_CUDA>([](auto& block) {
  block.setCollisionO(
    gpu::cuda::getFusedCollisionO<T,DESCRIPTOR,
                                  BGKdynamics<T,DESCRIPTOR>,
                                  BounceBack<T,DESCRIPTOR>,
                                  BounceBackVelocity<T,DESCRIPTOR>>());
});

Definition at line 280 of file operator.hh.

                     {
  return [](ConcreteBlockLattice<T,DESCRIPTOR,Platform::GPU_CUDA>& block) {
    bool* subdomain = block.template getData<CollisionSubdomainMask>().deviceData();
    DeviceContext<T,DESCRIPTOR> lattice(block);
    if (block.statisticsEnabled()) {
      call_operators_with_statistics(
        lattice,
        subdomain,
        MaskedCollision<T,DESCRIPTOR,DYNAMICS>{
          block.template getData<OperatorParameters<DYNAMICS>>().parameters,
          block.template getData<DynamicsMask<DYNAMICS>>().deviceData()
        }...,
        DynamicDispatchCollision{});
    } else {
      call_operators(
        lattice,
        subdomain,
        MaskedCollision<T,DESCRIPTOR,DYNAMICS>{
          block.template getData<OperatorParameters<DYNAMICS>>().parameters,
          block.template getData<DynamicsMask<DYNAMICS>>().deviceData()
        }...,
        DynamicDispatchCollision{});
    }
  };
}

References call_operators(), and call_operators_with_statistics().

Here is the call graph for this function:

◆ make_thrust_tuple_f() [1/2]

template<unsigned D, typename F >

auto olb::gpu::cuda::make_thrust_tuple_f ( F && f )

Definition at line 46 of file integral.hh.

                                {
  return make_thrust_tuple_f<F>(std::forward<F&&>(f), std::make_index_sequence<D>{});
}

References make_thrust_tuple_f().

Here is the call graph for this function:

◆ make_thrust_tuple_f() [2/2]

template<typename F , std::size_t... INDICES>

auto olb::gpu::cuda::make_thrust_tuple_f	(	F &&	f,
		std::index_sequence< INDICES... >	)

Definition at line 41 of file integral.hh.

                                                               {
  return thrust::make_tuple(f(INDICES)...);
}

Here is the caller graph for this function:

◆ make_thrust_tuple_of() [1/2]

template<typename T >

auto olb::gpu::cuda::make_thrust_tuple_of ( T & f )

Definition at line 82 of file integral.hh.

                                {
  return thrust::make_tuple(f);
}

◆ make_thrust_tuple_of() [2/2]

template<typename T , unsigned D>

auto olb::gpu::cuda::make_thrust_tuple_of ( Vector< T, D > & f )

Definition at line 75 of file integral.hh.

                                          {
  return gpu::cuda::make_thrust_tuple_f<f.d>([&](unsigned iDim) {
    return f[iDim];
  });
}

References olb::GenericVector< T, D, IMPL >::d, and make_thrust_tuple_f().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ make_thrust_tuple_of_device_data()

template<typename VECTOR >

auto olb::gpu::cuda::make_thrust_tuple_of_device_data ( VECTOR & f )

Definition at line 87 of file integral.hh.

                                                 {
  return gpu::cuda::make_thrust_tuple_f<f.d>([&](unsigned iDim) {
    return f[iDim].deviceData();
  });
}

References make_thrust_tuple_f().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ scatter_any_fields()

void olb::gpu::cuda::scatter_any_fields	(	thrust::device_vector< AnyDeviceFieldArrayD > &	fields,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Blocking scatter of fields data in buffer to given indices.

Definition at line 317 of file communicator.hh.

                                            {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::scatter_any_fields<<<block_count,block_size>>>(
    fields.data().get(), fields.size(),
    indices.data().get(), indices.size(),
    buffer);
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::scatter_any_fields().

Here is the call graph for this function:

◆ scatter_field()

template<typename FIELD , typename CONTEXT >

void olb::gpu::cuda::scatter_field	(	CONTEXT &	lattice,
		const thrust::device_vector< CellID > &	indices,
		std::uint8_t *	buffer )

Blocking scatter of FIELD data in buffer to given indices.

Definition at line 291 of file communicator.hh.

                                                                                                     {
  const auto block_size = 32;
  const auto block_count = (indices.size() + block_size - 1) / block_size;
  kernel::scatter_field<CONTEXT,FIELD><<<block_count,block_size>>>(
    lattice,
    indices.data().get(), indices.size(),
    reinterpret_cast<typename FIELD::template value_type<typename CONTEXT::value_t>*>(buffer));
  device::check();
}

References olb::gpu::cuda::device::check(), and olb::gpu::cuda::kernel::scatter_field().

Here is the call graph for this function:

◆ thrust_tuple_cat() [1/4]

template<typename A , typename B , typename C , std::size_t... AIs, std::size_t... BIs, std::size_t... CIs>

auto olb::gpu::cuda::thrust_tuple_cat	(	A &	a,
		B &	b,
		C &	c,
		std::index_sequence< AIs... >	,
		std::index_sequence< BIs... >	,
		std::index_sequence< CIs... >	)

Definition at line 62 of file integral.hh.

                                                                                                                         {
  return thrust::make_tuple(thrust::get<AIs>(a)...,
                            thrust::get<BIs>(b)...,
                            thrust::get<CIs>(c)...);
}

◆ thrust_tuple_cat() [2/4]

template<typename A , typename B , std::size_t... AIs, std::size_t... BIs>

auto olb::gpu::cuda::thrust_tuple_cat	(	A &	a,
		B &	b,
		std::index_sequence< AIs... >	,
		std::index_sequence< BIs... >	)

Definition at line 56 of file integral.hh.

                                                                                        {
  return thrust::make_tuple(thrust::get<AIs>(a)...,
                            thrust::get<BIs>(b)...);
}

◆ thrust_tuple_cat() [3/4]

template<typename A , std::size_t... AIs>

auto olb::gpu::cuda::thrust_tuple_cat	(	A &	a,
		std::index_sequence< AIs... >	)

Definition at line 51 of file integral.hh.

                                                       {
  return thrust::make_tuple(thrust::get<AIs>(a)...);
}

Here is the caller graph for this function:

◆ thrust_tuple_cat() [4/4]

template<typename... ARGS>

auto olb::gpu::cuda::thrust_tuple_cat ( ARGS &&... args )

Definition at line 69 of file integral.hh.

                                      {
  return thrust_tuple_cat(args...,
                          std::make_index_sequence<thrust::tuple_size<ARGS>::value>{}...);
}

References thrust_tuple_cat().

Here is the call graph for this function:

Variable Documentation

◆ field_type_index

template<typename CONTEXT , typename TYPE >

__constant__ std::size_t olb::gpu::cuda::field_type_index

Mapping of TYPE in CONTEXT to runtime-fixed index.

Used for dynamic access to field arrays

Definition at line 40 of file registry.hh.

Namespaces

Classes

Functions

Variables

Detailed Description

Function Documentation

◆ async_call_list_operators()

◆ async_call_list_operators_with_statistics()

◆ async_call_operators()

◆ async_call_operators_with_statistics()

◆ async_copy_any_fields()

◆ async_copy_field()

◆ async_gather_any_fields()

◆ async_gather_field()

◆ async_scatter_any_fields()

◆ async_scatter_field()

◆ call_coupling_operators()

◆ call_list_operators()

◆ call_operators()

◆ call_operators_with_statistics()

◆ call_particle_coupling_operators()

◆ call_refinement_coupling_operator()

◆ gather_any_fields()

◆ gather_field()

◆ getFusedCollisionO()

◆ make_thrust_tuple_f() [1/2]

◆ make_thrust_tuple_f() [2/2]

◆ make_thrust_tuple_of() [1/2]

◆ make_thrust_tuple_of() [2/2]

◆ make_thrust_tuple_of_device_data()

◆ scatter_any_fields()

◆ scatter_field()

◆ thrust_tuple_cat() [1/4]

◆ thrust_tuple_cat() [2/4]

◆ thrust_tuple_cat() [3/4]

◆ thrust_tuple_cat() [4/4]

Variable Documentation

◆ field_type_index