Lattice field accessors. More...

#include <core/DataTypes.h>
#include <core/cell/Cell.h>
#include <core/cell/CellInterval.h>
#include <core/math/Matrix3.h>
#include <core/math/Vector3.h>
#include <field/iterators/IteratorMacros.h>
#include <gpu/FieldAccessor.h>
#include <gpu/FieldIndexing.h>
#include <gpu/GPUField.h>
#include <gpu/Kernel.h>
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
#include <array>
#include <vector>

Include dependency graph for FieldAccessorsSinglePrecisionCUDA.cu:

Go to the source code of this file.

Namespaces
namespace	walberla
	\file PackInfoPdfDoublePrecision.cpp \author pystencils

namespace	walberla::lbm

namespace	walberla::lbm::accessor

namespace	walberla::lbm::accessor::Population

namespace	walberla::lbm::accessor::Vector

namespace	walberla::lbm::accessor::Interpolation

namespace	walberla::lbm::accessor::Equilibrium

namespace	walberla::lbm::accessor::Density

namespace	walberla::lbm::accessor::Velocity

namespace	walberla::lbm::accessor::Force

namespace	walberla::lbm::accessor::MomentumDensity

namespace	walberla::lbm::accessor::PressureTensor

Macros
#define	RESTRICT

Functions
static __forceinline__ __device__ uint	getLinearIndex (uint3 blockIdx, uint3 threadIdx, uint3 gridDim, uint3 blockDim, uint fOffset)
	Get linear index of flattened data with original layout `fzyx`.

__global__ void	walberla::lbm::accessor::Population::kernel_get (gpu::FieldAccessor< float > pdf, float *RESTRICT pop)

__global__ void	walberla::lbm::accessor::Population::kernel_set (gpu::FieldAccessor< float > pdf, float const *RESTRICT pop)

__global__ void	walberla::lbm::accessor::Population::kernel_broadcast (gpu::FieldAccessor< float > pdf, float const *RESTRICT pop)

__global__ void	walberla::lbm::accessor::Population::kernel_set_vel (gpu::FieldAccessor< float > pdf, gpu::FieldAccessor< float > velocity, gpu::FieldAccessor< float > force, float const *RESTRICT pop)

std::array< float, 19u >	walberla::lbm::accessor::Population::get (gpu::GPUField< float > const *pdf_field, Cell const &cell)
	Get populations from a single cell.

void	walberla::lbm::accessor::Population::set (gpu::GPUField< float > *pdf_field, std::array< float, 19u > const &pop, Cell const &cell)
	Set populations on a single cell.

void	walberla::lbm::accessor::Population::set (gpu::GPUField< float > pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > const *force_field, std::array< float, 19u > const &pop, Cell const &cell)
	Set populations and recalculate velocities on a single cell.

void	walberla::lbm::accessor::Population::initialize (gpu::GPUField< float > *pdf_field, std::array< float, 19u > const &pop)
	Initialize all cells with the same value.

std::vector< float >	walberla::lbm::accessor::Population::get (gpu::GPUField< float > const *pdf_field, CellInterval const &ci)
	Get populations from a cell interval.

void	walberla::lbm::accessor::Population::set (gpu::GPUField< float > *pdf_field, std::vector< float > const &values, CellInterval const &ci)
	Set populations on a cell interval.

void	walberla::lbm::accessor::Population::set (gpu::GPUField< float > pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > const *force_field, std::vector< float > const &values, CellInterval const &ci)
	Set populations and recalculate velocities on a cell interval.

__global__ void	walberla::lbm::accessor::Vector::kernel_get (gpu::FieldAccessor< float > vec, float *u_out)

__global__ void	walberla::lbm::accessor::Vector::kernel_set (gpu::FieldAccessor< float > vec, float const *RESTRICT u_in)

__global__ void	walberla::lbm::accessor::Vector::kernel_broadcast (gpu::FieldAccessor< float > vec, float const *RESTRICT u_in)

__global__ void	walberla::lbm::accessor::Vector::kernel_add (gpu::FieldAccessor< float > vec, float const *RESTRICT u_in)

__global__ void	walberla::lbm::accessor::Vector::kernel_broadcast_add (gpu::FieldAccessor< float > vec, float const *RESTRICT u_in)

Vector3< float >	walberla::lbm::accessor::Vector::get (gpu::GPUField< float > const *field, Cell const &cell)
	Get value from a single cell.

void	walberla::lbm::accessor::Vector::set (gpu::GPUField< float > *field, Vector3< float > const &vec, Cell const &cell)
	Set value on a single cell.

void	walberla::lbm::accessor::Vector::add (gpu::GPUField< float > *field, Vector3< float > const &vec, Cell const &cell)
	Add value to a single cell.

void	walberla::lbm::accessor::Vector::initialize (gpu::GPUField< float > *field, Vector3< float > const &vec)
	Initialize all cells with the same value.

void	walberla::lbm::accessor::Vector::add_to_all (gpu::GPUField< float > *field, Vector3< float > const &vec)
	Add value to all cells.

std::vector< float >	walberla::lbm::accessor::Vector::get (gpu::GPUField< float > const *vec_field, CellInterval const &ci)
	Get values from a cell interval.

void	walberla::lbm::accessor::Vector::set (gpu::GPUField< float > *vec_field, std::vector< float > const &values, CellInterval const &ci)
	Set values on a cell interval.

static __forceinline__ __device__ void	walberla::lbm::accessor::Interpolation::calculate_weights (float const RESTRICT const pos, int RESTRICT const corner, float *RESTRICT const weights, uint gl)
	Calculate interpolation weights.

__global__ void	walberla::lbm::accessor::Interpolation::kernel_get (gpu::FieldAccessor< float > vec, float const RESTRICT const pos, float RESTRICT const vel, uint n_pos, uint gl)

__global__ void	walberla::lbm::accessor::Interpolation::kernel_set (gpu::FieldAccessor< float > vec, float const RESTRICT const pos, float const RESTRICT const forces, uint n_pos, uint gl)

static dim3	walberla::lbm::accessor::Interpolation::calculate_dim_grid (uint const threads_x, uint const blocks_per_grid_y, uint const threads_per_block)

std::vector< float >	walberla::lbm::accessor::Interpolation::get (gpu::GPUField< float > const *vec_field, std::vector< float > const &pos, uint gl)

void	walberla::lbm::accessor::Interpolation::set (gpu::GPUField< float > const *vec_field, std::vector< float > const &pos, std::vector< float > const &forces, uint gl)

__device__ void	walberla::lbm::accessor::Equilibrium::kernel_set_device (gpu::FieldAccessor< float > pdf, float const *RESTRICT const u, float rho)

__global__ void	walberla::lbm::accessor::Density::kernel_get (gpu::FieldAccessor< float > pdf, float *RESTRICT rho_out)

__global__ void	walberla::lbm::accessor::Density::kernel_set (gpu::FieldAccessor< float > pdf, float const *RESTRICT rho_in)

float	walberla::lbm::accessor::Density::get (gpu::GPUField< float > const *pdf_field, Cell const &cell)

void	walberla::lbm::accessor::Density::set (gpu::GPUField< float > *pdf_field, const float rho, Cell const &cell)

std::vector< float >	walberla::lbm::accessor::Density::get (gpu::GPUField< float > const *pdf_field, CellInterval const &ci)

void	walberla::lbm::accessor::Density::set (gpu::GPUField< float > *pdf_field, std::vector< float > const &values, CellInterval const &ci)

__global__ void	walberla::lbm::accessor::Velocity::kernel_get (gpu::FieldAccessor< float > pdf, gpu::FieldAccessor< float > force, float *RESTRICT u_out)

__global__ void	walberla::lbm::accessor::Velocity::kernel_set (gpu::FieldAccessor< float > pdf, gpu::FieldAccessor< float > velocity, gpu::FieldAccessor< float > force, float const *RESTRICT u_in)

Vector3< float >	walberla::lbm::accessor::Velocity::get (gpu::GPUField< float > const pdf_field, gpu::GPUField< float > const force_field, Cell const &cell)

std::vector< float >	walberla::lbm::accessor::Velocity::get (gpu::GPUField< float > const pdf_field, gpu::GPUField< float > const force_field, CellInterval const &ci)

void	walberla::lbm::accessor::Velocity::set (gpu::GPUField< float > pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > const *force_field, Vector3< float > const &u, Cell const &cell)

void	walberla::lbm::accessor::Velocity::set (gpu::GPUField< float > pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > const *force_field, std::vector< float > const &values, CellInterval const &ci)

__global__ void	walberla::lbm::accessor::Force::kernel_set (gpu::FieldAccessor< float > pdf, gpu::FieldAccessor< float > velocity, gpu::FieldAccessor< float > force, float const *RESTRICT f_in)

void	walberla::lbm::accessor::Force::set (gpu::GPUField< float > const pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > *force_field, Vector3< float > const &u, Cell const &cell)

void	walberla::lbm::accessor::Force::set (gpu::GPUField< float > const pdf_field, gpu::GPUField< float > velocity_field, gpu::GPUField< float > *force_field, std::vector< float > const &values, CellInterval const &ci)

__global__ void	walberla::lbm::accessor::MomentumDensity::kernel_sum (gpu::FieldAccessor< float > pdf, gpu::FieldAccessor< float > force, float *RESTRICT out)

Vector3< float >	walberla::lbm::accessor::MomentumDensity::reduce (gpu::GPUField< float > const pdf_field, gpu::GPUField< float > const force_field)

__global__ void	walberla::lbm::accessor::PressureTensor::kernel_get (gpu::FieldAccessor< float > pdf, float *RESTRICT p_out)

Matrix3< float >	walberla::lbm::accessor::PressureTensor::get (gpu::GPUField< float > const *pdf_field, Cell const &cell)

std::vector< float >	walberla::lbm::accessor::PressureTensor::get (gpu::GPUField< float > const *pdf_field, CellInterval const &ci)