Cuda (.cu) file for the Lattice Boltzmann implementation on GPUs. More...

#include <stdio.h>
#include <cuda.h>
#include <stdlib.h>
#include "lbgpu.h"
#include "config.h"

Include dependency graph for lbgpu.cu:

Defines
#define	GAUSSRANDOM
#define	BOUNCEBACK
Functions
device funktions called by kernel funktions
__device__ void	atomicadd (float *address, float value)
	atomic add function for sveral cuda architectures
__device__ void	random_01 (LB_randomnr_gpu *rn)
	randomgenerator which generates numbers [0,1]
__device__ void	gaussian_random (LB_randomnr_gpu *rn)
	gaussian random nummber generator for thermalisation
__device__ void	index_to_xyz (unsigned int index, unsigned int *xyz)
	tranformation from 1d array-index to xyz
__device__ void	calc_m_from_n (LB_nodes_gpu n_a, unsigned int index, float *mode)
	calculation of the modes from the velocitydensities (space-transform.)
__device__ void	relax_modes (float *mode, unsigned int index, LB_node_force_gpu node_f)
	lb_relax_modes, means collision update of the modes
__device__ void	thermalize_modes (float mode, unsigned int index, LB_randomnr_gpu rn)
	thermalization of the modes with gaussian random numbers
__device__ void	normalize_modes (float *mode)
	normalization of the modes need befor backtransformation into velocity space
__device__ void	calc_n_from_modes_push (LB_nodes_gpu n_b, float *mode, unsigned int index)
	backtransformation from modespace to desityspace and streaming with the push method using pbc
__device__ void	bounce_back_read (LB_nodes_gpu n_b, LB_nodes_gpu n_a, unsigned int index, float LB_boundary_velocity, float LB_boundary_force)
	Bounce back boundary conditions.
__device__ void	bounce_back_write (LB_nodes_gpu n_b, LB_nodes_gpu n_a, unsigned int index)
	bounce back read kernel needed to avoid raceconditions
__device__ void	apply_forces (unsigned int index, float *mode, LB_node_force_gpu node_f)
	add of (external) forces within the modespace, needed for particle-interaction
__device__ void	calc_values (LB_nodes_gpu n_a, float mode, LB_values_gpu d_v, unsigned int index, unsigned int singlenode)
	function used to calc physical values of every node
__device__ void	calc_mode (float *mode, LB_nodes_gpu n_a, unsigned int node_index)
Coupling part
__device__ void	calc_viscous_force (LB_nodes_gpu n_a, float delta, LB_particle_gpu particle_data, LB_particle_force_gpu particle_force, unsigned int part_index, LB_randomnr_gpu rn_part, float delta_j, unsigned int node_index)
	(Eq.
__device__ void	calc_node_force (float delta, float delta_j, unsigned int *node_index, LB_node_force_gpu node_f)
	calcutlation of the node force caused by the particles, with atomicadd due to avoiding race conditions (Eq.
Host functions to setup and call kernels
void	lb_init_GPU (LB_parameters_gpu *lbpar_gpu)
	initialization for the lb gpu fluid called from host
void	lb_reinit_GPU (LB_parameters_gpu *lbpar_gpu)
	reinitialization for the lb gpu fluid called from host
void	lb_realloc_particle_GPU (LB_parameters_gpu lbpar_gpu, LB_particle_gpu *host_data)
	setup and call particle reallocation from the host
void	lb_init_boundaries_GPU (int host_n_lb_boundaries, int number_of_boundnodes, int host_boundary_node_list, int host_boundary_index_list, float *host_LB_Boundary_velocity)
	setup and call boundaries from the host
void	lb_reinit_extern_nodeforce_GPU (LB_parameters_gpu *lbpar_gpu)
	setup and call extern single node force initialization from the host
void	lb_init_extern_nodeforces_GPU (int n_extern_nodeforces, LB_extern_nodeforce_gpu host_extern_nodeforces, LB_parameters_gpu lbpar_gpu)
	setup and call extern single node force initialization from the host
void	lb_particle_GPU (LB_particle_gpu *host_data)
	setup and call particle kernel from the host
void	lb_copy_forces_GPU (LB_particle_force_gpu *host_forces)
	setup and call kernel to copy particle forces to host
void	lb_get_values_GPU (LB_values_gpu *host_values)
	setup and call kernel for getting macroscopic fluid values of all nodes
void	lb_get_boundary_flags_GPU (unsigned int *host_bound_array)
	get all the boundary flags for all nodes
void	lb_print_node_GPU (int single_nodeindex, LB_values_gpu *host_print_values)
	setup and call kernel for getting macroscopic fluid values of a single node
void	lb_calc_fluid_mass_GPU (double *mass)
	setup and call kernel to calculate the total momentum of the hole fluid
void	lb_calc_fluid_momentum_GPU (double *host_mom)
	setup and call kernel to calculate the total momentum of the hole fluid
void	lb_calc_fluid_temperature_GPU (double *host_temp)
	setup and call kernel to calculate the temperature of the hole fluid
void	lb_save_checkpoint_GPU (float host_checkpoint_vd, unsigned int host_checkpoint_seed, unsigned int host_checkpoint_boundary, float host_checkpoint_force)
	setup and call kernel for getting macroscopic fluid values of all nodes
void	lb_load_checkpoint_GPU (float host_checkpoint_vd, unsigned int host_checkpoint_seed, unsigned int host_checkpoint_boundary, float host_checkpoint_force)
	setup and call kernel for setting macroscopic fluid values of all nodes
void	lb_get_boundary_flag_GPU (int single_nodeindex, unsigned int *host_flag)
	setup and call kernel to get the boundary flag of a single node
void	lb_set_node_velocity_GPU (int single_nodeindex, float *host_velocity)
	set the net velocity at a single node
void	reinit_parameters_GPU (LB_parameters_gpu *lbpar_gpu)
	reinit of params
void	lb_integrate_GPU ()
	integration kernel for the lb gpu fluid update called from host
void	lb_gpu_get_boundary_forces (double *forces)
void	lb_free_GPU ()
	free gpu memory kernel called from the host (not used anymore)
Variables
static LB_values_gpu *	device_values = NULL
	defining structures residing in global memory
static LB_nodes_gpu	nodes_a
	structs for velocity densities
static LB_nodes_gpu	nodes_b
static LB_particle_force_gpu *	particle_force = NULL
	struct for particle force
static LB_particle_gpu *	particle_data = NULL
	struct for particle position and veloctiy
static LB_node_force_gpu	node_f
	struct for node force
static LB_particle_seed_gpu *	part = NULL
	struct for storing particle rn seed
static LB_extern_nodeforce_gpu *	extern_nodeforces = NULL
static float *	LB_boundary_force = NULL
static float *	LB_boundary_velocity = NULL
static int *	boundary_node_list
	pointer for bound index array
static int *	boundary_index_list
static __device__ __constant__ int	n_lb_boundaries_gpu = 0
static size_t	size_of_boundindex
static int *	gpu_check = NULL
	pointers for additional cuda check flag
static int *	h_gpu_check = NULL
static unsigned int	intflag = 1
static LB_nodes_gpu *	current_nodes = NULL
static size_t	size_of_values
	defining size values for allocating global memory
static size_t	size_of_forces
static size_t	size_of_positions
static size_t	size_of_seed
static size_t	size_of_extern_nodeforces
static __device__ __constant__ LB_parameters_gpu	para
	parameters residing in constant memory
static const float	c_sound_sq = 1.f/3.f
cudaStream_t	stream [1]
	cudasteams for parallel computing on cpu and gpu
cudaError_t	err
cudaError_t	_err
int	initflag = 0
int	partinitflag = 0
System setup and Kernel funktions
#define	cuda_safe_mem(a) _cuda_safe_mem((a), __FILE__, __LINE__)
#define	KERNELCALL(_f, _a, _b, _params)
__global__ void	calc_n_equilibrium (LB_nodes_gpu n_a, int *gpu_check)
	kernel to calculate local populations from hydrodynamic fields given by the tcl values.
__global__ void	set_u_equilibrium (LB_nodes_gpu n_a, int single_nodeindex, float *velocity)
	kernel to calculate local populations from hydrodynamic fields from given flow field velocities.
__global__ void	init_particle_force (LB_particle_force_gpu particle_force, LB_particle_seed_gpu part)
	kernel for the initalisation of the particle force array
__global__ void	reset_particle_force (LB_particle_force_gpu *particle_force)
	kernel for the initalisation of the partikel force array
__global__ void	reinit_node_force (LB_node_force_gpu node_f)
	(re-)initialization of the node force / set up of external force in lb units
__global__ void	init_boundaries (int boundary_node_list, int boundary_index_list, int number_of_boundnodes, LB_nodes_gpu n_a, LB_nodes_gpu n_b)
	set the boundary flag for all boundary nodes
__global__ void	reset_boundaries (LB_nodes_gpu n_a, LB_nodes_gpu n_b)
	reset the boundary flag of every node
__global__ void	integrate (LB_nodes_gpu n_a, LB_nodes_gpu n_b, LB_values_gpu *d_v, LB_node_force_gpu node_f)
	integrationstep of the lb-fluid-solver
__global__ void	calc_fluid_particle_ia (LB_nodes_gpu n_a, LB_particle_gpu particle_data, LB_particle_force_gpu particle_force, LB_node_force_gpu node_f, LB_particle_seed_gpu *part)
	part interaction kernel
__global__ void	bb_read (LB_nodes_gpu n_a, LB_nodes_gpu n_b, float LB_boundary_velocity, float LB_boundary_force)
	Bounce back boundary read kernel.
__global__ void	bb_write (LB_nodes_gpu n_a, LB_nodes_gpu n_b)
	Bounce back boundary write kernel.
__global__ void	values (LB_nodes_gpu n_a, LB_values_gpu *d_v)
	get physical values of the nodes (density, velocity, ...)
__global__ void	lb_get_boundaries (LB_nodes_gpu n_a, unsigned int *device_bound_array)
	get boundary flags
__global__ void	init_extern_nodeforces (int n_extern_nodeforces, LB_extern_nodeforce_gpu *extern_nodeforces, LB_node_force_gpu node_f)
	set extern force on single nodes kernel
__global__ void	lb_print_node (int single_nodeindex, LB_values_gpu *d_p_v, LB_nodes_gpu n_a)
	print single node values kernel
__global__ void	calc_mass (LB_nodes_gpu n_a, float *sum)
	calculate mass of the hole fluid kernel
__global__ void	momentum (LB_nodes_gpu n_a, float *sum, LB_node_force_gpu node_f)
	calculate momentum of the hole fluid kernel
__global__ void	temperature (LB_nodes_gpu n_a, float *cpu_jsquared)
	calculate temperature of the fluid kernel
__global__ void	lb_get_boundary_flag (int single_nodeindex, unsigned int *device_flag, LB_nodes_gpu n_a)
	print single node boundary flag
void	_cuda_safe_mem (cudaError_t err, char *file, unsigned int line)
	erroroutput for memory allocation and memory copy

Detailed Description

Cuda (.cu) file for the Lattice Boltzmann implementation on GPUs.

Header file for lbgpu.h.

Definition in file lbgpu.cu.

Define Documentation

#define BOUNCEBACK

Value:

shift = para.agrid*para.agrid*para.agrid*para.agrid*para.rho*2.*3.*weight*para.tau*(v[0]*c[0] + v[1]*c[1] + v[2]*c[2]); \
  pop_to_bounce_back = n_b.vd[population*para.number_of_nodes + index ]; \
  to_index_x = (x+c[0]+para.dim_x)%para.dim_x; \
  to_index_y = (y+c[1]+para.dim_y)%para.dim_y; \
  to_index_z = (z+c[2]+para.dim_z)%para.dim_z; \
  to_index = to_index_x + para.dim_x*to_index_y + para.dim_x*para.dim_y*to_index_z; \
  if (n_b.boundary[to_index] == 0) \
  { \
    boundary_force[0] += (2*pop_to_bounce_back+shift)*c[0]/para.tau/para.tau/para.agrid; \
    boundary_force[1] += (2*pop_to_bounce_back+shift)*c[1]/para.tau/para.tau/para.agrid; \
    boundary_force[2] += (2*pop_to_bounce_back+shift)*c[2]/para.tau/para.tau/para.agrid; \
    n_b.vd[inverse*para.number_of_nodes + to_index ] = pop_to_bounce_back + shift; \
  }

Referenced by bounce_back_read().

#define cuda_safe_mem ( a ) _cuda_safe_mem((a), __FILE__, __LINE__)

Definition at line 1413 of file lbgpu.cu.

Referenced by lb_calc_fluid_mass_GPU(), lb_calc_fluid_momentum_GPU(), lb_calc_fluid_temperature_GPU(), lb_get_boundary_flag_GPU(), lb_get_boundary_flags_GPU(), lb_gpu_get_boundary_forces(), lb_init_boundaries_GPU(), lb_init_extern_nodeforces_GPU(), lb_init_GPU(), lb_integrate_GPU(), lb_print_node_GPU(), lb_realloc_particle_GPU(), lb_reinit_extern_nodeforce_GPU(), lb_reinit_GPU(), lb_set_node_velocity_GPU(), and reinit_parameters_GPU().

#define GAUSSRANDOM

Definition at line 37 of file lbgpu.cu.

#define KERNELCALL	(	_f,
		_a,
		_b,
		_params
	)

Value:

_f<<<_a, _b, 0, stream[0]>>>_params; \
_err=cudaGetLastError(); \
if (_err!=cudaSuccess){ \
  printf("CUDA error: %s\n", cudaGetErrorString(_err)); \
  fprintf(stderr, "error calling %s with dim %d %d %d #thpb %d in %s:%u\n", #_f, _a.x, _a.y, _a.z, _b, __FILE__, __LINE__); \
  exit(EXIT_FAILURE); \
}

Definition at line 1415 of file lbgpu.cu.

Referenced by lb_calc_fluid_mass_GPU(), lb_calc_fluid_momentum_GPU(), lb_calc_fluid_temperature_GPU(), lb_copy_forces_GPU(), lb_get_boundary_flag_GPU(), lb_get_boundary_flags_GPU(), lb_get_values_GPU(), lb_init_boundaries_GPU(), lb_init_extern_nodeforces_GPU(), lb_init_GPU(), lb_integrate_GPU(), lb_particle_GPU(), lb_print_node_GPU(), lb_realloc_particle_GPU(), lb_reinit_extern_nodeforce_GPU(), lb_reinit_GPU(), and lb_set_node_velocity_GPU().

Function Documentation

void _cuda_safe_mem	(	cudaError_t	err,
		char *	file,
		unsigned int	line
	)

erroroutput for memory allocation and memory copy

Parameters:

err	cuda error code
*file	.cu file were the error took place
line	line of the file were the error took place

Definition at line 1399 of file lbgpu.cu.

References _err.

__device__ void apply_forces	(	unsigned int	index,
		float *	mode,
		LB_node_force_gpu	node_f
	)

add of (external) forces within the modespace, needed for particle-interaction

Parameters:

index	node index / thread index (Input)
mode	Pointer to the local register values mode (Input/Output)
node_f	Pointer to local node force (Input)

hydrodynamic momentum density is redefined when forces present

update momentum modes

update stress modes

Definition at line 602 of file lbgpu.cu.

References LB_parameters_gpu::agrid, LB_parameters_gpu::ext_force, LB_parameters_gpu::external_force, LB_node_force_gpu::force, LB_parameters_gpu::gamma_bulk, LB_parameters_gpu::gamma_shear, LB_parameters_gpu::number_of_nodes, para, LB_parameters_gpu::rho, and LB_parameters_gpu::tau.

Referenced by integrate().

__device__ void atomicadd	(	float *	address,
		float	value
	)		`[inline]`

atomic add function for sveral cuda architectures

Definition at line 99 of file lbgpu.cu.

Referenced by bounce_back_read(), calc_mass(), calc_node_force(), momentum(), and temperature().

__global__ void bb_read	(	LB_nodes_gpu	n_a,
		LB_nodes_gpu	n_b,
		float *	LB_boundary_velocity,
		float *	LB_boundary_force
	)

Bounce back boundary read kernel.

Parameters:

n_a	Pointer to local node residing in array a (Input)
n_b	Pointer to local node residing in array b (Input)
LB_boundary_velocity	The constant velocity at the boundary, set by the user (Input)
LB_boundary_force	The force on the boundary nodes (Output)

Definition at line 1239 of file lbgpu.cu.

References bounce_back_read(), LB_parameters_gpu::number_of_nodes, and para.

Referenced by lb_integrate_GPU().

__global__ void bb_write	(	LB_nodes_gpu	n_a,
		LB_nodes_gpu	n_b
	)

Bounce back boundary write kernel.

Parameters:

n_a	Pointer to local node residing in array a (Input)
n_b	Pointer to local node residing in array b (Input)

Definition at line 1252 of file lbgpu.cu.

References bounce_back_write(), LB_parameters_gpu::number_of_nodes, and para.

__device__ void bounce_back_read	(	LB_nodes_gpu	n_b,
		LB_nodes_gpu	n_a,
		unsigned int	index,
		float *	LB_boundary_velocity,
		float *	LB_boundary_force
	)

Bounce back boundary conditions.

The populations that have propagated into a boundary node are bounced back to the node they came from. This results in no slip boundary conditions.

[cf. Ladd and Verberg, J. Stat. Phys. 104(5/6):1191-1251, 2001]

Parameters:

index	node index / thread index (Input)
n_b	Pointer to local node residing in array b (Input)
n_a	Pointer to local node residing in array a (Output) (temp stored in buffer a)
LB_boundary_velocity	The constant velocity at the boundary, set by the user (Input)
LB_boundary_force	The force on the boundary nodes (Output)

store vd temporary in second lattice to avoid race conditions

Definition at line 453 of file lbgpu.cu.

References atomicadd(), BOUNCEBACK, LB_nodes_gpu::boundary, and index_to_xyz().

Referenced by bb_read().

__device__ void bounce_back_write	(	LB_nodes_gpu	n_b,
		LB_nodes_gpu	n_a,
		unsigned int	index
	)

bounce back read kernel needed to avoid raceconditions

Parameters:

index	node index / thread index (Input)
n_b	Pointer to local node residing in array b (Input)
n_a	Pointer to local node residing in array a (Output) (temp stored in buffer a)

stream vd from boundary node back to origin node

Definition at line 566 of file lbgpu.cu.

References LB_nodes_gpu::boundary, LB_parameters_gpu::dim_x, LB_parameters_gpu::dim_y, LB_parameters_gpu::dim_z, index_to_xyz(), LB_parameters_gpu::number_of_nodes, para, and LB_nodes_gpu::vd.

Referenced by bb_write().

__global__ void calc_fluid_particle_ia	(	LB_nodes_gpu	n_a,
		LB_particle_gpu *	particle_data,
		LB_particle_force_gpu *	particle_force,
		LB_node_force_gpu	node_f,
		LB_particle_seed_gpu *	part
	)

part interaction kernel

Parameters:

n_a	Pointer to local node residing in array a (Input)
*particle_data	Pointer to the particle position and velocity (Input)
*particle_force	Pointer to the particle force (Input)
*part	Pointer to the rn array of the particles (Input)
node_f	Pointer to local node force (Input)

calc of the force which act on the particle

calc of the force which acts back to the fluid node

Definition at line 1214 of file lbgpu.cu.

References calc_node_force(), calc_viscous_force(), LB_parameters_gpu::number_of_particles, para, LB_randomnr_gpu::seed, and LB_particle_seed_gpu::seed.

Referenced by lb_particle_GPU().

__device__ void calc_m_from_n	(	LB_nodes_gpu	n_a,
		unsigned int	index,
		float *	mode
	)

calculation of the modes from the velocitydensities (space-transform.)

Parameters:

n_a	Pointer to local node residing in array a (Input)
index	node index / thread index (Input)
mode	Pointer to the local register values mode (Output)

Definition at line 171 of file lbgpu.cu.

References LB_parameters_gpu::number_of_nodes, para, and LB_nodes_gpu::vd.

Referenced by integrate(), lb_print_node(), and values().

__global__ void calc_mass	(	LB_nodes_gpu	n_a,
		float *	sum
	)

calculate mass of the hole fluid kernel

Parameters:

*sum	Pointer to result storage value (Output)
n_a	Pointer to local node residing in array a (Input)

Definition at line 1326 of file lbgpu.cu.

References LB_parameters_gpu::agrid, atomicadd(), calc_mode(), LB_parameters_gpu::number_of_nodes, para, and LB_parameters_gpu::rho.

Referenced by lb_calc_fluid_mass_GPU().

__device__ void calc_mode	(	float *	mode,
		LB_nodes_gpu	n_a,
		unsigned int	node_index
	)

Parameters:

node_index	node index around (8) particle (Input)
*mode	Pointer to the local register values mode (Output)
n_a	Pointer to local node residing in array a(Input)

mass mode

momentum modes

Definition at line 731 of file lbgpu.cu.

References LB_parameters_gpu::number_of_nodes, para, and LB_nodes_gpu::vd.

Referenced by calc_mass(), calc_viscous_force(), momentum(), set_u_equilibrium(), and temperature().

__global__ void calc_n_equilibrium	(	LB_nodes_gpu	n_a,
		int *	gpu_check
	)

kernel to calculate local populations from hydrodynamic fields given by the tcl values.

The mapping is given in terms of the equilibrium distribution.

Eq. (2.15) Ladd, J. Fluid Mech. 271, 295-309 (1994) Eq. (4) in Berk Usta, Ladd and Butler, JCP 122, 094902 (2005)

Parameters:

n_a	Pointer to the lattice site (Input).
*gpu_check	additional check if gpu kernel are executed(Input).

default values for fields in lattice units

reduce the pressure tensor to the part needed here

update the q=0 sublattice

update the q=1 sublattice

update the q=2 sublattice

set different seed for randomgen on every node

Definition at line 909 of file lbgpu.cu.

References LB_parameters_gpu::agrid, c_sound_sq, LB_parameters_gpu::number_of_nodes, para, LB_parameters_gpu::rho, LB_nodes_gpu::seed, LB_nodes_gpu::vd, and LB_parameters_gpu::your_seed.

Referenced by lb_init_GPU(), and lb_reinit_GPU().

__device__ void calc_n_from_modes_push	(	LB_nodes_gpu	n_b,
		float *	mode,
		unsigned int	index
	)

backtransformation from modespace to desityspace and streaming with the push method using pbc

Parameters:

index	node index / thread index (Input)
mode	Pointer to the local register values mode (Input)
*n_b	Pointer to local node residing in array b (Output)

Definition at line 411 of file lbgpu.cu.

References LB_parameters_gpu::dim_x, LB_parameters_gpu::dim_y, LB_parameters_gpu::dim_z, index_to_xyz(), LB_parameters_gpu::number_of_nodes, para, and LB_nodes_gpu::vd.

Referenced by integrate().

__device__ void calc_node_force	(	float *	delta,
		float *	delta_j,
		unsigned int *	node_index,
		LB_node_force_gpu	node_f
	)

calcutlation of the node force caused by the particles, with atomicadd due to avoiding race conditions (Eq.

(14) Ahlrichs and Duenweg, JCP 111(17):8225 (1999))

Parameters:

*delta	Pointer for the weighting of particle position (Input)
*delta_j	Pointer for the weighting of particle momentum (Input)
node_index	node index around (8) particle (Input)
node_f	Pointer to the node force (Output).

Definition at line 861 of file lbgpu.cu.

References atomicadd(), LB_node_force_gpu::force, LB_parameters_gpu::number_of_nodes, and para.

Referenced by calc_fluid_particle_ia().

__device__ void calc_values	(	LB_nodes_gpu	n_a,
		float *	mode,
		LB_values_gpu *	d_v,
		unsigned int	index,
		unsigned int	singlenode
	)

function used to calc physical values of every node

Parameters:

index	node index / thread index (Input)
mode	Pointer to the local register values mode (Input)
n_a	Pointer to local node residing in array a for boundary flag(Input)
*d_v	Pointer to local device values (Input/Output)
singlenode	Flag, if there is only one node This function is a clone of lb_calc_local_fields and additionally performs unit conversions.

Definition at line 660 of file lbgpu.cu.

References LB_parameters_gpu::agrid, LB_parameters_gpu::gamma_bulk, LB_parameters_gpu::gamma_shear, i, para, LB_values_gpu::pi, LB_parameters_gpu::rho, LB_values_gpu::rho, LB_parameters_gpu::tau, and LB_values_gpu::v.

Referenced by lb_print_node(), and values().

__device__ void calc_viscous_force	(	LB_nodes_gpu	n_a,
		float *	delta,
		LB_particle_gpu *	particle_data,
		LB_particle_force_gpu *	particle_force,
		unsigned int	part_index,
		LB_randomnr_gpu *	rn_part,
		float *	delta_j,
		unsigned int *	node_index
	)

(Eq.

(12) Ahlrichs and Duenweg, JCP 111(17):8225 (1999))

Parameters:

n_a	Pointer to local node residing in array a (Input)
*delta	Pointer for the weighting of particle position (Output)
*delta_j	Pointer for the weighting of particle momentum (Output)
*particle_data	Pointer to the particle position and velocity (Input)
*particle_force	Pointer to the particle force (Input)
part_index	particle id / thread id (Input)
*rn_part	Pointer to randomnumber array of the particle
node_index	node index around (8) particle (Output)

see ahlrichs + duenweg page 8227 equ (10) and (11)

further value used for interpolation of fluid velocity at part pos near boundaries

calculate viscous force take care to rescale velocities with time_step and transform to MD units (Eq. (9) Ahlrichs and Duenweg, JCP 111(17):8225 (1999))

add stochastik force of zero mean (Ahlrichs, Duennweg equ. 15)

delta_j for transform momentum transfer to lattice units which is done in calc_node_force (Eq. (12) Ahlrichs and Duenweg, JCP 111(17):8225 (1999))

Definition at line 765 of file lbgpu.cu.

References LB_parameters_gpu::agrid, calc_mode(), LB_parameters_gpu::dim_x, LB_parameters_gpu::dim_y, LB_parameters_gpu::dim_z, LB_particle_force_gpu::f, LB_parameters_gpu::friction, gaussian_random(), i, LB_parameters_gpu::lb_coupl_pref, LB_parameters_gpu::lb_coupl_pref2, LB_particle_gpu::mu_E, my_left, LB_particle_gpu::p, para, random_01(), LB_randomnr_gpu::randomnr, LB_parameters_gpu::rho, LB_parameters_gpu::tau, LB_parameters_gpu::time_step, and LB_particle_gpu::v.

Referenced by calc_fluid_particle_ia().

__device__ void gaussian_random ( LB_randomnr_gpu * rn )

gaussian random nummber generator for thermalisation

Parameters:

*rn	Pointer to randomnumber array of the local node node or particle

On every second call two gaussian random numbers are calculated via the Box-Muller transformation.

draw two uniform random numbers in the unit circle

perform Box-Muller transformation

Definition at line 132 of file lbgpu.cu.

References random_01(), and LB_randomnr_gpu::randomnr.

Referenced by calc_particle_lattice_ia(), calc_viscous_force(), lb_thermalize_modes(), partial_momentum_update(), simple_momentum_update(), and thermalize_modes().

__device__ void index_to_xyz	(	unsigned int	index,
		unsigned int *	xyz
	)

tranformation from 1d array-index to xyz

Parameters:

index	node index / thread index (Input)
xyz	Pointer to calculated xyz array (Output)

Definition at line 157 of file lbgpu.cu.

References LB_parameters_gpu::dim_x, LB_parameters_gpu::dim_y, and para.

Referenced by bounce_back_read(), bounce_back_write(), and calc_n_from_modes_push().

__global__ void init_boundaries	(	int *	boundary_node_list,
		int *	boundary_index_list,
		int	number_of_boundnodes,
		LB_nodes_gpu	n_a,
		LB_nodes_gpu	n_b
	)

set the boundary flag for all boundary nodes

Parameters:

boundary_node_list	The indices of the boundary nodes
boundary_index_list	The flag representing the corresponding boundary
number_of_boundnodes	The number of boundary nodes
n_a	Pointer to local node residing in array a (Input)
n_b	Pointer to local node residing in array b (Input)

Definition at line 1145 of file lbgpu.cu.

References LB_nodes_gpu::boundary.

Referenced by lb_init_boundaries_GPU().

__global__ void init_extern_nodeforces	(	int	n_extern_nodeforces,
		LB_extern_nodeforce_gpu *	extern_nodeforces,
		LB_node_force_gpu	node_f
	)

set extern force on single nodes kernel

Parameters:

n_extern_nodeforces	number of nodes (Input)
*extern_nodeforces	Pointer to extern node force array (Input)
node_f	node force struct (Output)

Definition at line 1295 of file lbgpu.cu.

References LB_parameters_gpu::agrid, LB_node_force_gpu::force, LB_extern_nodeforce_gpu::force, LB_extern_nodeforce_gpu::index, LB_parameters_gpu::number_of_nodes, para, and LB_parameters_gpu::tau.

Referenced by lb_init_extern_nodeforces_GPU().

__global__ void init_particle_force	(	LB_particle_force_gpu *	particle_force,
		LB_particle_seed_gpu *	part
	)

kernel for the initalisation of the particle force array

Parameters:

*particle_force	Pointer to local particle force (Output)
*part	Pointer to the particle rn seed storearray (Output)

Definition at line 1083 of file lbgpu.cu.

References LB_particle_force_gpu::f, LB_parameters_gpu::number_of_particles, para, LB_particle_seed_gpu::seed, and LB_parameters_gpu::your_seed.

Referenced by lb_init_GPU(), and lb_realloc_particle_GPU().

__global__ void integrate	(	LB_nodes_gpu	n_a,
		LB_nodes_gpu	n_b,
		LB_values_gpu *	d_v,
		LB_node_force_gpu	node_f
	)

integrationstep of the lb-fluid-solver

Parameters:

n_a	Pointer to local node residing in array a (Input)
n_b	Pointer to local node residing in array b (Input)
*d_v	Pointer to local device values (Input)
node_f	Pointer to local node force (Input)

every node is connected to a thread via the index

the 19 moments (modes) are only temporary register values

storing the seed into a register value

calc_m_from_n

lb_relax_modes

lb_thermalize_modes

if external force is used apply node force

lb_calc_n_from_modes_push

calc of velocity densities and streaming with pbc

rewriting the seed back to the global memory

Definition at line 1174 of file lbgpu.cu.

References apply_forces(), calc_m_from_n(), calc_n_from_modes_push(), LB_parameters_gpu::fluct, normalize_modes(), LB_parameters_gpu::number_of_nodes, LB_parameters_gpu::number_of_particles, para, relax_modes(), LB_nodes_gpu::seed, LB_randomnr_gpu::seed, and thermalize_modes().

Referenced by lb_integrate_GPU().

void lb_calc_fluid_mass_GPU ( double * mass )

setup and call kernel to calculate the total momentum of the hole fluid

Parameters:

*mass value of the mass calcutated on the GPU

values for the kernel call

Definition at line 1744 of file lbgpu.cu.

References calc_mass(), cuda_safe_mem, KERNELCALL, lbpar_gpu, and LB_parameters_gpu::number_of_nodes.

Referenced by tclcommand_analyze_fluid_mass_gpu().

void lb_calc_fluid_momentum_GPU ( double * host_mom )

setup and call kernel to calculate the total momentum of the hole fluid

Parameters:

host_mom value of the momentum calcutated on the GPU

values for the kernel call

Definition at line 1768 of file lbgpu.cu.

References LB_parameters_gpu::agrid, cuda_safe_mem, KERNELCALL, lbpar_gpu, momentum(), LB_parameters_gpu::number_of_nodes, and LB_parameters_gpu::tau.

Referenced by tclcommand_analyze_fluid_parse_momentum_gpu().

void lb_calc_fluid_temperature_GPU ( double * host_temp )

setup and call kernel to calculate the temperature of the hole fluid

Parameters:

host_temp value of the temperatur calcutated on the GPU

values for the kernel call

Definition at line 1793 of file lbgpu.cu.

References LB_parameters_gpu::agrid, cuda_safe_mem, LB_parameters_gpu::dim_x, LB_parameters_gpu::dim_y, LB_parameters_gpu::dim_z, KERNELCALL, lbpar_gpu, LB_parameters_gpu::number_of_nodes, LB_parameters_gpu::rho, LB_parameters_gpu::tau, and temperature.

Referenced by tclcommand_analyze_fluid_parse_temperature_gpu().

void lb_copy_forces_GPU ( LB_particle_force_gpu * host_forces )

setup and call kernel to copy particle forces to host

Parameters:

*host_forces contains the particle force computed on the GPU

Copy result from device memory to host memory

values for the particle kernel

reset part forces with zero

Definition at line 1671 of file lbgpu.cu.

References KERNELCALL, lbpar_gpu, LB_parameters_gpu::number_of_particles, reset_particle_force(), and size_of_forces.

Referenced by lb_send_forces_gpu().

void lb_free_GPU ( )

free gpu memory kernel called from the host (not used anymore)

Definition at line 1935 of file lbgpu.cu.

References para, and stream.

__global__ void lb_get_boundaries	(	LB_nodes_gpu	n_a,
		unsigned int *	device_bound_array
	)

get boundary flags

Parameters:

n_a	Pointer to local node residing in array a (Input)
device_bound_array	Pointer to local device values (Input)

Definition at line 1281 of file lbgpu.cu.

References LB_nodes_gpu::boundary, LB_parameters_gpu::number_of_nodes, and para.

Referenced by lb_get_boundary_flags_GPU().

__global__ void lb_get_boundary_flag	(	int	single_nodeindex,
		unsigned int *	device_flag,
		LB_nodes_gpu	n_a
	)

print single node boundary flag

Parameters:

single_nodeindex	index of the node (Input)
*device_flag	Pointer to result storage array (Input)
n_a	Pointer to local node residing in array a (Input)

Definition at line 1386 of file lbgpu.cu.

References LB_nodes_gpu::boundary.

Referenced by lb_get_boundary_flag_GPU().

void lb_get_boundary_flag_GPU	(	int	single_nodeindex,
		unsigned int *	host_flag
	)

setup and call kernel to get the boundary flag of a single node

Parameters:

single_nodeindex	number of the node to get the flag for
host_flag	her goes the value of the boundary flag

Definition at line 1839 of file lbgpu.cu.

References cuda_safe_mem, KERNELCALL, and lb_get_boundary_flag().

Referenced by lb_lbnode_get_boundary().

void lb_get_boundary_flags_GPU ( unsigned int * host_bound_array )

get all the boundary flags for all nodes

Parameters:

host_bound_array here go the values of the boundary flag

values for the kernel call

Definition at line 1707 of file lbgpu.cu.

References cuda_safe_mem, KERNELCALL, lb_get_boundaries(), lbpar_gpu, and LB_parameters_gpu::number_of_nodes.

Referenced by lb_lbfluid_print_boundary(), and lb_lbfluid_print_vtk_boundary().

void lb_get_values_GPU ( LB_values_gpu * host_values )

setup and call kernel for getting macroscopic fluid values of all nodes

Parameters:

*host_values struct to save the gpu values

values for the kernel call

Definition at line 1691 of file lbgpu.cu.

References KERNELCALL, lbpar_gpu, LB_parameters_gpu::number_of_nodes, size_of_values, and values().

Referenced by lb_lbfluid_print_velocity(), and lb_lbfluid_print_vtk_velocity().

void lb_gpu_get_boundary_forces ( double * forces )

Definition at line 1923 of file lbgpu.cu.

References cuda_safe_mem, free, i, LB_boundary_force, malloc, and n_lb_boundaries.

Referenced by lbboundary_get_force().

void lb_init_boundaries_GPU	(	int	host_n_lb_boundaries,
		int	number_of_boundnodes,
		int *	host_boundary_node_list,
		int *	host_boundary_index_list,
		float *	host_LB_Boundary_velocity
	)

setup and call boundaries from the host

Parameters:

host_n_lb_boundaries	number of LB boundaries
number_of_boundnodes	number of boundnodes
host_boundary_node_list	The indices of the boundary nodes
host_boundary_index_list	The flag representing the corresponding boundary
host_LB_Boundary_velocity	The constant velocity at the boundary, set by the user (Input)

values for the kernel call

Definition at line 1573 of file lbgpu.cu.

References boundary_index_list, boundary_node_list, cuda_safe_mem, init_boundaries(), KERNELCALL, LB_boundary_force, LB_boundary_velocity, lbpar_gpu, n_lb_boundaries, n_lb_boundaries_gpu, LB_parameters_gpu::number_of_nodes, reset_boundaries(), and size_of_boundindex.

Referenced by lb_init_boundaries().

void lb_init_extern_nodeforces_GPU	(	int	n_extern_nodeforces,
		LB_extern_nodeforce_gpu *	host_extern_nodeforces,
		LB_parameters_gpu *	lbpar_gpu
	)

setup and call extern single node force initialization from the host

Parameters:

n_extern_nodeforces	number of nodes on which the external force has to be applied
*host_extern_nodeforces	Pointer to the host extern node forces
*lbpar_gpu	Pointer to host parameter struct

Definition at line 1635 of file lbgpu.cu.

References cuda_safe_mem, LB_parameters_gpu::external_force, init_extern_nodeforces(), KERNELCALL, para, and size_of_extern_nodeforces.

Referenced by lb_lbnode_set_extforce_GPU().

void lb_init_GPU ( LB_parameters_gpu * lbpar_gpu )

initialization for the lb gpu fluid called from host

Parameters:

*lbpar_gpu Pointer to parameters to setup the lb field

Allocate structs in device memory

write parameters in const memory

check flag if lb gpu init works

values for the kernel call

values for the particle kernel

calc of veloctiydensities from given parameters and initialize the Node_Force array with zero

init part forces with zero

Definition at line 1433 of file lbgpu.cu.

References LB_nodes_gpu::boundary, calc_n_equilibrium(), cuda_safe_mem, errexit(), LB_node_force_gpu::force, gpu_check, h_gpu_check, init_particle_force(), initflag, intflag, KERNELCALL, malloc, nodes_a, LB_parameters_gpu::number_of_nodes, LB_parameters_gpu::number_of_particles, para, reinit_node_force(), reset_boundaries(), LB_nodes_gpu::seed, size_of_forces, size_of_positions, size_of_seed, size_of_values, stream, and LB_nodes_gpu::vd.

Referenced by lb_init_gpu().

void lb_integrate_GPU ( )

integration kernel for the lb gpu fluid update called from host

values for the kernel call

call of fluid step

Definition at line 1883 of file lbgpu.cu.

References bb_read(), cuda_safe_mem, integrate(), intflag, KERNELCALL, LB_boundary_force, LB_boundary_velocity, lbpar_gpu, n_lb_boundaries, nodes_a, nodes_b, and LB_parameters_gpu::number_of_nodes.

Referenced by lattice_boltzmann_update_gpu().

void lb_load_checkpoint_GPU	(	float *	host_checkpoint_vd,
		unsigned int *	host_checkpoint_seed,
		unsigned int *	host_checkpoint_boundary,
		float *	host_checkpoint_force
	)

setup and call kernel for setting macroscopic fluid values of all nodes

Parameters:

*host_values struct to set stored values

Definition at line 1825 of file lbgpu.cu.

References LB_nodes_gpu::boundary, LB_node_force_gpu::force, intflag, lbpar_gpu, LB_parameters_gpu::number_of_nodes, LB_nodes_gpu::seed, and LB_nodes_gpu::vd.

Referenced by lb_lbfluid_load_checkpoint().

void lb_particle_GPU ( LB_particle_gpu * host_data )

setup and call particle kernel from the host

Parameters:

**host_data Pointer to the host particle positions and velocities

get espresso md particle values

call of the particle kernel

values for the particle kernel

Definition at line 1655 of file lbgpu.cu.

References calc_fluid_particle_ia(), KERNELCALL, lbpar_gpu, LB_parameters_gpu::number_of_particles, size_of_positions, and stream.

Referenced by lb_calc_particle_lattice_ia_gpu().

__global__ void lb_print_node	(	int	single_nodeindex,
		LB_values_gpu *	d_p_v,
		LB_nodes_gpu	n_a
	)

print single node values kernel

Parameters:

single_nodeindex	index of the node (Input)
*d_p_v	Pointer to result storage array (Input)
n_a	Pointer to local node residing in array a (Input)

Definition at line 1311 of file lbgpu.cu.

References calc_m_from_n(), and calc_values().

Referenced by lb_print_node_GPU().

void lb_print_node_GPU	(	int	single_nodeindex,
		LB_values_gpu *	host_print_values
	)

setup and call kernel for getting macroscopic fluid values of a single node

Definition at line 1726 of file lbgpu.cu.

References cuda_safe_mem, KERNELCALL, and lb_print_node().

Referenced by lb_lbnode_get_pi_neq(), lb_lbnode_get_rho(), and lb_lbnode_get_u().

void lb_realloc_particle_GPU	(	LB_parameters_gpu *	lbpar_gpu,
		LB_particle_gpu **	host_data
	)

setup and call particle reallocation from the host

Parameters:

*lbpar_gpu	Pointer to parameters to setup the lb field
**host_data	Pointer to host information data

Allocate struct for particle positions

pinned memory mode - use special function to get OS-pinned memory

values for the particle kernel

Definition at line 1531 of file lbgpu.cu.

References cuda_safe_mem, init_particle_force(), KERNELCALL, LB_parameters_gpu::number_of_particles, para, partinitflag, size_of_forces, size_of_positions, and size_of_seed.

Referenced by lb_realloc_particles_gpu().

void lb_reinit_extern_nodeforce_GPU ( LB_parameters_gpu * lbpar_gpu )

setup and call extern single node force initialization from the host

Parameters:

*lbpar_gpu Pointer to host parameter struct

values for the kernel call

Definition at line 1617 of file lbgpu.cu.

References cuda_safe_mem, KERNELCALL, LB_parameters_gpu::number_of_nodes, para, and reinit_node_force().

Referenced by lb_lbfluid_set_ext_force().

void lb_reinit_GPU ( LB_parameters_gpu * lbpar_gpu )

reinitialization for the lb gpu fluid called from host

Parameters:

*lbpar_gpu Pointer to parameters to setup the lb field

write parameters in const memory

values for the kernel call

calc of veloctiydensities from given parameters and initialize the Node_Force array with zero

Definition at line 1512 of file lbgpu.cu.

References calc_n_equilibrium(), cuda_safe_mem, gpu_check, KERNELCALL, LB_parameters_gpu::number_of_nodes, and para.

Referenced by lb_reinit_fluid_gpu().

void lb_save_checkpoint_GPU	(	float *	host_checkpoint_vd,
		unsigned int *	host_checkpoint_seed,
		unsigned int *	host_checkpoint_boundary,
		float *	host_checkpoint_force
	)

setup and call kernel for getting macroscopic fluid values of all nodes

Parameters:

*host_values struct to save the gpu values

Definition at line 1814 of file lbgpu.cu.

References LB_nodes_gpu::boundary, LB_node_force_gpu::force, lbpar_gpu, LB_parameters_gpu::number_of_nodes, LB_nodes_gpu::seed, and LB_nodes_gpu::vd.

Referenced by lb_lbfluid_save_checkpoint().

void lb_set_node_velocity_GPU	(	int	single_nodeindex,
		float *	host_velocity
	)

set the net velocity at a single node

Parameters:

single_nodeindex	the node to set the velocity for
host_velocity	the velocity to set

Definition at line 1859 of file lbgpu.cu.

References cuda_safe_mem, KERNELCALL, and set_u_equilibrium().

Referenced by lb_lbnode_set_u().

__global__ void momentum	(	LB_nodes_gpu	n_a,
		float *	sum,
		LB_node_force_gpu	node_f
	)

calculate momentum of the hole fluid kernel

Parameters:

node_f	node force struct (Input)
*sum	Pointer to result storage value (Output)
n_a	Pointer to local node residing in array a (Input)

Definition at line 1345 of file lbgpu.cu.

References atomicadd(), LB_nodes_gpu::boundary, calc_mode(), LB_node_force_gpu::force, LB_parameters_gpu::number_of_nodes, and para.

Referenced by lb_calc_fluid_momentum(), lb_calc_fluid_momentum_GPU(), predict_momentum_particles(), and tclcommand_analyze_parse_and_print_momentum().

__device__ void normalize_modes ( float * mode )

normalization of the modes need befor backtransformation into velocity space

Parameters:

mode	Pointer to the local register values mode (Input/Output)

normalization factors enter in the back transformation

Definition at line 381 of file lbgpu.cu.

Referenced by integrate().

__device__ void random_01 ( LB_randomnr_gpu * rn )

randomgenerator which generates numbers [0,1]

Parameters:

*rn	Pointer to randomnumber array of the local node or particle

Definition at line 116 of file lbgpu.cu.

References LB_randomnr_gpu::randomnr, and LB_randomnr_gpu::seed.

Referenced by calc_viscous_force(), gaussian_random(), and thermalize_modes().

__global__ void reinit_node_force ( LB_node_force_gpu node_f )

(re-)initialization of the node force / set up of external force in lb units

Parameters:

node_f Pointer to local node force (Input)

Definition at line 1114 of file lbgpu.cu.

References LB_parameters_gpu::agrid, LB_parameters_gpu::ext_force, LB_parameters_gpu::external_force, LB_node_force_gpu::force, LB_parameters_gpu::number_of_nodes, para, and LB_parameters_gpu::tau.

Referenced by lb_init_GPU(), and lb_reinit_extern_nodeforce_GPU().

void reinit_parameters_GPU ( LB_parameters_gpu * lbpar_gpu )

reinit of params

Parameters:

*lbpar_gpu struct containing the paramters of the fluid

write parameters in const memory

Definition at line 1877 of file lbgpu.cu.

References cuda_safe_mem, and para.

Referenced by lb_reinit_parameters_gpu().

__device__ void relax_modes	(	float *	mode,
		unsigned int	index,
		LB_node_force_gpu	node_f
	)

lb_relax_modes, means collision update of the modes

Parameters:

index	node index / thread index (Input)
mode	Pointer to the local register values mode (Input/Output)
node_f	Pointer to local node force (Input)

re-construct the real density remember that the populations are stored as differences to their equilibrium value

if forces are present, the momentum density is redefined to inlcude one half-step of the force action. See the Chapman-Enskog expansion in [Ladd & Verberg].

equilibrium part of the stress modes (eq13 schiller)

relax the stress modes (eq14 schiller)

relax the ghost modes (project them out)

ghost modes have no equilibrium part due to orthogonality

Definition at line 245 of file lbgpu.cu.

References LB_parameters_gpu::agrid, LB_node_force_gpu::force, LB_parameters_gpu::gamma_bulk, LB_parameters_gpu::gamma_even, LB_parameters_gpu::gamma_odd, LB_parameters_gpu::gamma_shear, LB_parameters_gpu::number_of_nodes, para, and LB_parameters_gpu::rho.

Referenced by integrate().