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/copy.h>
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/transform.h>
#include <array>
#include <vector>

Include dependency graph for FieldAccessorsSinglePrecisionCUDA.cu:

Go to the source code of this file.

Classes
struct	algo_rescale
	Rescale values in a device vector by some constant. More...

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`.

static dim3	calculate_dim_grid (uint const threads_x, uint const blocks_per_grid_y, uint const threads_per_block)

__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)

__global__ void	walberla::lbm::accessor::Vector::kernel_set_from_list (gpu::FieldAccessor< float > vec, int const RESTRICT const indices, float const RESTRICT const values, uint length)

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.

void	walberla::lbm::accessor::Vector::set_from_list (gpu::GPUField< float > const *field, thrust::device_vector< int > const &indices, thrust::device_vector< float > const &values, uint gl)

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_rho (gpu::FieldAccessor< float > pdf, float const RESTRICT const pos, float RESTRICT const rho_out, float const density, uint n_pos, uint gl)

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

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

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

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

void	walberla::lbm::accessor::Interpolation::add_force (gpu::GPUField< float > const *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, float const density)

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

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

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

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

void	walberla::lbm::accessor::Density::set (gpu::GPUField< float > *pdf_field, std::vector< float > const &values, float const density, 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, float const density)

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, float const density, 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, float const density, CellInterval const &ci)

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

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

__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, float const density, Cell const &cell)

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

Matrix3< float >	walberla::lbm::accessor::PressureTensor::reduce (gpu::GPUField< float > const *pdf_field, float const density)