olb::BlockLatticeFieldReductionO< FIELD, REDUCTION_OP, CONDITION >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > > Struct Template Reference

#include <fieldReduction.h>

Collaboration diagram for olb::BlockLatticeFieldReductionO< FIELD, REDUCTION_OP, CONDITION >::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 FIELD, typename REDUCTION_OP, typename CONDITION>
template<typename T, typename DESCRIPTOR>
struct olb::BlockLatticeFieldReductionO< FIELD, REDUCTION_OP, CONDITION >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >

Definition at line 33 of file fieldReduction.h.

Member Function Documentation

apply()

template<typename FIELD , typename REDUCTION_OP , typename CONDITION >

template<typename T , typename DESCRIPTOR >

void olb::BlockLatticeFieldReductionO< FIELD, REDUCTION_OP, CONDITION >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >::apply

(

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

blockLattice

)

Definition at line 186 of file fieldReduction.hh.

{
  auto&       parameters = blockLattice.template getData<OperatorParameters<BlockLatticeFieldReductionO>>().parameters;
  const auto& blockField = blockLattice.template getField<FIELD>();
 
  FieldD<T, DESCRIPTOR, fields::array_of<FIELD>> elementField =
      parameters.template get<fields::array_of<FIELD>>(); //not auto. because of considering 1 dimensional field
 
  /*note
   -blockField[iD].deviceData(); is return  row pointer T* in device.
   -thrust::device_pointer_cast(elementField[iD])[0] is that thrust::device_pointer_cast is casting to thrust::device_ptr<T> from T* in device.
   -thrust::device_pointer_cast(
          blockField[iD].deviceData()); is that thrust::device_pointer_cast is casting to thrust::device_ptr<T> from T* in device.
   -thrust::raw_pointer_cast(d_partial.data());, with thrust::device_vector<T> d_partial(block_count, (T)0.0);, is that thrust::raw_pointer_cast is casting to T* from thrust::raw_pointer_cast in device.
 
  */
 
  for (unsigned iD = 0; iD < blockField.d; iD++) {
    cudaPointerAttributes attributes;
    cudaError_t           error = cudaPointerGetAttributes(&attributes, elementField[iD]);
    if (error == cudaSuccess) {
      if (attributes.devicePointer != nullptr) {
      }
      else if (attributes.hostPointer != nullptr) {
        std::cout << "elementField is on the host" << std::endl;
      }
    }
    else {
      std::cerr << "Error in cudaPointerGetAttributes: " << cudaGetErrorString(error) << std::endl;
    }
    auto devPtr = thrust::device_pointer_cast(
        blockField[iD].deviceData()); // thrust::device_pointer_cast is casting to thrust::device_ptr<T> from T*.
    if (devPtr == nullptr) {
      std::cerr << "Invalid device pointer at index " << iD << std::endl;
      return;
    }
    reductionFunctionDevice<T, DESCRIPTOR, REDUCTION_OP, CONDITION>(devPtr, blockLattice, elementField[iD]);
  }
}

References olb::reductionFunctionDevice().

Here is the call graph for this function:

setup()

template<typename FIELD , typename REDUCTION_OP , typename CONDITION >

template<typename T , typename DESCRIPTOR >

void olb::BlockLatticeFieldReductionO< FIELD, REDUCTION_OP, CONDITION >::type< ConcreteBlockLattice< T, DESCRIPTOR, Platform::GPU_CUDA > >::setup

(

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

blockLattice

)

Definition at line 178 of file fieldReduction.hh.

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

---

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

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