olb::IntegratePorousElementFieldsO< FIELDS >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > > Struct Template Reference

#include <integral.h>

Collaboration diagram for olb::IntegratePorousElementFieldsO< FIELDS >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >:

Public Member Functions
void	setup (ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &blockLattice)
void	apply (ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &blockLattice)

Detailed Description

template<typename... FIELDS>
template<typename T, typename DESCRIPTOR>
struct olb::IntegratePorousElementFieldsO< FIELDS >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >

Definition at line 33 of file integral.h.

Member Function Documentation

apply()

template<typename... FIELDS>

template<typename T , typename DESCRIPTOR >

void olb::IntegratePorousElementFieldsO< FIELDS >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >::apply

(

ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &

blockLattice

)

Definition at line 127 of file integral.hh.

{
  auto& tag = blockLattice.template getField<fields::fsi::REDUCED_ELEMENT_TAG>();
 
  auto begin = thrust::make_zip_iterator(gpu::cuda::thrust_tuple_cat(
                 gpu::cuda::make_thrust_tuple_of_device_data(tag),
                 gpu::cuda::make_thrust_tuple_of_device_data(
                   blockLattice.template getField<FIELDS>())...
               ));
  auto end = thrust::partition(thrust::device,
                               begin,
                               begin + blockLattice.getNcells(),
                               gpu::cuda::first_tuple_element_non_zero{});
  std::size_t nActiveCells = end - begin;
 
  if (nActiveCells > 0) {
    // Group per-cell field values by FSI tag
    thrust::sort_by_key(thrust::device,
                        tag[0].deviceData(),
                        tag[0].deviceData() + nActiveCells,
                        thrust::make_zip_iterator(gpu::cuda::thrust_tuple_cat(
                          gpu::cuda::make_thrust_tuple_of_device_data(
                            blockLattice.template getField<FIELDS>())...
                        )));
  }
 
  // Sum up grouped force values
  auto& parameters = blockLattice.template getData<OperatorParameters<IntegratePorousElementFieldsO>>().parameters;
  if (nActiveCells > 0) {
    auto reducedElementTag = parameters.template get<fields::array_of<fields::fsi::REDUCED_ELEMENT_TAG>>();
    const std::size_t nKeys = thrust::reduce_by_key(
      thrust::device,
      tag[0].deviceData(),
      tag[0].deviceData() + nActiveCells,
      thrust::make_zip_iterator(gpu::cuda::thrust_tuple_cat(
        gpu::cuda::make_thrust_tuple_of_device_data(
          blockLattice.template getField<FIELDS>())...
      )),
      reducedElementTag,
      thrust::make_zip_iterator(gpu::cuda::thrust_tuple_cat(
        gpu::cuda::make_thrust_tuple_of(parameters.template get<fields::array_of<FIELDS>>())...
      )),
      thrust::equal_to<T>{},
      gpu::cuda::sum_of_thrust_tuple{}
    ).first - reducedElementTag;
    parameters.template set<fields::fsi::REDUCED_ELEMENTS_COUNT>(nKeys);
  } else {
    parameters.template set<fields::fsi::REDUCED_ELEMENTS_COUNT>(0);
  }
 
  if (nActiveCells > 0) {
    thrust::fill(thrust::device,
                 tag[0].deviceData(),
                 tag[0].deviceData() + nActiveCells,
                 0);
  }
}

References olb::BlockStructureD< D >::getNcells(), olb::gpu::cuda::make_thrust_tuple_of(), olb::gpu::cuda::make_thrust_tuple_of_device_data(), and olb::gpu::cuda::thrust_tuple_cat().

Here is the call graph for this function:

setup()

template<typename... FIELDS>

template<typename T , typename DESCRIPTOR >

void olb::IntegratePorousElementFieldsO< FIELDS >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >::setup

(

ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > &

blockLattice

)

Definition at line 119 of file integral.hh.

{
  blockLattice.template getData<OperatorParameters<IntegratePorousElementFieldsO>>();
}

---

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

• src/fsi/platform/gpu/cuda/integral.h
• src/fsi/platform/gpu/cuda/integral.hh