Data Structures
class	context_info

class	gpu_vector

class	grid_info

struct	kernel_buffers
	Device buffers used by the collocate and integrate kernels. Which fields are non-null depends on which kernel is launched. More...

struct	kernel_params
	Parameters of the collocate kernel. More...

struct	ldiffs_value
	Differences in angular momentum. More...

struct	orbital
	Orbital angular momentum. More...

class	smem_parameters

struct	smem_task
	data needed for calculating the coefficients, forces, and stress More...

struct	smem_task_reduced
	data needed for collocate and integrate kernels More...

struct	task_info
	Internal representation of a task. More...

Functions
template<typename T , bool IS_FUNC_AB>
__device__ __inline__ void	block_to_cab (const kernel_params &params, const smem_task< T > &task, T *cab)
	Decontracts the subblock, going from spherical to cartesian harmonics.

template<typename T , bool IS_FUNC_AB>
__global__	__launch_bounds__ (64) void calculate_coefficients(const kernel_params dev_)

template<typename T , typename T3 , bool distributed__, bool orthogonal_>
__global__	__launch_bounds__ (64) void collocate_kernel(const kernel_params dev_)

constexpr size_t	align_up_elems (size_t n_elems, size_t elem_alignment)

size_t	round_up16 (size_t n)

template<typename T >
__device__ __inline__ T	block_reduce_64 (T *table, const T val_, const int tid)

template<typename T , typename T3 , bool COMPUTE_TAU>
__global__	__launch_bounds__ (64) void compute_hab_forces(const kernel_params dev_)

template<typename T , typename T3 , bool distributed__, bool orthogonal_, int lbatch = 20>
__global__	__launch_bounds__ (64) void integrate_kernel(const kernel_params dev_)

__device__ __inline__ double	fac (const int i)
	Factorial function, e.g. fac(5) = 5! = 120.

__host__ __device__ __inline__ int	ncoset (const int l)
	Number of Cartesian orbitals up to given angular momentum quantum.

__host__ __device__ __inline__ int	coset (int lx, int ly, int lz)
	Maps three angular momentum components to a single zero based index.

__device__ __inline__ orbital	up (const int i, const orbital &a)
	Increase i'th component of given orbital angular momentum.

__inline__ __device__ orbital	down (const int i, const orbital &a)
	Decrease i'th component of given orbital angular momentum.

__inline__ __device__ int	idx (const orbital a)
	Return coset index of given orbital angular momentum.

__device__ __inline__ double	power (const double x, const int expo)

template<typename T = double>
__device__ __inline__ void	prep_term (const orbital a, const orbital b, const T value, const int n, T *cab)
	Adds given value to matrix element cab[idx(b)][idx(a)].

static void	init_constant_memory ()
	Initializes the device's constant memory.

__inline__ __device__ unsigned int	block_index ()

__inline__ __device__ unsigned int	thread_global_index ()

__inline__ __device__ double3	compute_coordinates (const double *__restrict__ dh_, const double x, const double y, const double z)

__inline__ __device__ float3	compute_coordinates (const float *__restrict__ dh_, const float x, const float y, const float z)

template<typename T >
__inline__ __device__ void	compute_alpha (const smem_task< T > &task, T *__restrict__ alpha)
	Computes the polynomial expansion coefficients: (x-a)lxa (x-b)lxb -> sum_{ls} alpha(ls,lxa,lxb,1)(x-p)*ls.

__host__ __inline__ __device__ void	convert_to_lattice_coordinates (const double dh_inv_, const double3 __restrict__ const rp, double3 *__restrict__ rp_c)

__host__ __inline__ __device__ void	convert_from_lattice_coordinates_to_cartesian (const double __restrict__ dh_, const double3 __restrict__ const rp, double3 *__restrict__ rp_c)

__host__ __inline__ __device__ void	convert_to_lattice_coordinates (const float dh_inv_, const float3 __restrict__ const rp, float3 *__restrict__ rp_c)

__host__ __inline__ __device__ void	convert_from_lattice_coordinates_to_cartesian (const float __restrict__ dh_, const float3 __restrict__ const rp, float3 *__restrict__ rp_c)

template<typename T , typename T3 , bool orthorhombic_>
__inline__ T	compute_cube_properties (const T radius, const T const __restrict__ dh_, const T const __restrict__ dh_inv_, const T3 __restrict__ rp, T3 __restrict__ roffset, int3 __restrict__ cubecenter, int3 __restrict__ lb_cube, int3 *__restrict__ cube_size)

__inline__ __device__ void	compute_window_size (const int3 grid_size, const int border_mask, const int3 &border_width, int3 &window_size, int3 &window_shift)

__inline__ __device__ int	wrap_grid_index (const int idx, const int full_size)
	Wraps a cube-local coordinate into the periodic full grid along one axis. Single point of entry shared by collocate_kernel and integrate_kernel.

template<typename T , typename T3 , bool distributed__>
__device__ __inline__ void	setup_task_cube_center (const kernel_params &dev_, smem_task_reduced< T, T3 > &task)
	Shifts task.cube_center into the local grid's coordinate system and, for distributed grids, derives the border window. Single point of entry shared by collocate_kernel and integrate_kernel so the two cannot drift apart on which axis maps to which.

template<typename T >
__device__ static __inline__ void	cab_to_cxyz (const smem_task< T > &task, const T __restrict__ alpha, const T __restrict__ cab, T *__restrict__ cxyz)
	Transforms coefficients C_ab into C_xyz.

template<typename T >
__device__ static __inline__ void	cxyz_to_cab (const smem_task< T > &task, const T __restrict__ alpha, const T __restrict__ cxyz, T *__restrict__ cab)
	Transforms coefficients C_xyz into C_ab.

template<typename T , typename T3 >
__device__ __inline__ void	fill_smem_task_reduced (const kernel_params &dev, const int task_id, smem_task_reduced< T, T3 > &task)
	Copies a task from global to shared memory.

template<typename T >
__device__ __inline__ void	fill_smem_task_coef (const kernel_params &dev, const int task_id, smem_task< T > &task)
	Copies a task from global to shared memory and does precomputations.

template<typename T >
__inline__ __device__ T *	allocate_workspace (const kernel_params &dev_)

template<typename T , typename T3 >
__device__ __inline__ T	calculate_ymix_ymax_boundaries (smem_task_reduced< T, T3 > &task, const int z, int &ymin, int &ymax)

template<typename T , typename T3 >
__device__ __inline__ void	calculate_xmin_xmax_boundaries (smem_task_reduced< T, T3 > &task, const int y, const T kremain, int &xmin, int &xmax)

template<typename T >
__device__ __inline__ void	prepare_pab_AB (const orbital a, const orbital b, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_AB, ie. identity transformation.

template<typename T >
__device__ __inline__ void	prepare_pab_DADB (const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_DADB.

template<typename T >
__device__ void	prepare_pab_ADBmDAB (const int idir, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_ADBmDAB_{X,Y,Z}.

template<typename T >
__device__ void	prepare_pab_ARDBmDARB (const int idir, const int ir, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_ARDBmDARB_{X,Y,Z}{X,Y,Z}.

template<typename T >
__device__ void	prepare_pab_DABpADB (const int idir, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_DABpADB_{X,Y,Z}.

template<typename T >
__device__ __inline__ void	prepare_pab_Di (const int ider, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_{DX,DY,DZ}.

template<typename T >
__device__ __inline__ void	oneterm_dijdij (const int idir, const T func_a, const orbital a, const orbital b, const T zetb, const int n, T *cab)
	Helper for grid_prepare_pab_DiDj.

template<typename T >
__device__ __inline__ void	prepare_pab_DiDj (const int ider1, const int ider2, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_{DXDY,DYDZ,DZDX}.

template<typename T >
__device__ void	oneterm_diidii (const int idir, const T func_a, const orbital a, const orbital b, const T zetb, const int n, T *cab)
	Helper for grid_prepare_pab_Di2.

template<typename T >
__device__ __inline__ void	prepare_pab_Di2 (const int ider, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Implementation of function GRID_FUNC_{DXDX,DYDY,DZDZ}.

template<typename T >
__device__ __inline__ void	prepare_pab (const enum grid_func func, const orbital a, const orbital b, const T zeta, const T zetb, const T pab_val, const int n, T *cab)
	Transforms a given element of the density matrix according to func.

ldiffs_value	prepare_get_ldiffs (const enum grid_func func)
	Returns difference in angular momentum range for given func.

template<typename T >
__device__ __inline__ T	get_term (const orbital &a, const orbital &b, const int n, const T *cab)
	Returns matrix element cab[idx(b)][idx(a)].

template<typename T >
__device__ __inline__ T	get_force_a_normal (const orbital &a, const orbital &b, const int i, const T zeta, const int n, const T *cab)
	Returns i'th component of force on atom a for compute_tau=false.

template<bool compute_tau, typename T >
__device__ __inline__ double	get_force_a (const orbital &a, const orbital &b, const int i, const T zeta, const T zetb, const int n, const T *cab)
	Returns i'th component of force on atom a.

template<typename T >
__device__ __inline__ double	get_force_b_normal (const orbital &a, const orbital &b, const int i, const T zetb, const T rab[3], const int n, const T *cab)
	Returns i'th component of force on atom b for compute_tau=false.

template<bool compute_tau, typename T >
__device__ __inline__ T	get_force_b (const orbital &a, const orbital &b, const int i, const T zeta, const T zetb, const T rab[3], const int n, const T *cab)
	Returns i'th component of force on atom b.

template<typename T >
__device__ __inline__ double	get_virial_a_normal (const orbital &a, const orbital &b, const int i, const int j, const T zeta, const int n, const T *cab)
	Returns element i,j of virial on atom a for compute_tau=false.

template<bool compute_tau, typename T >
__device__ __inline__ T	get_virial_a (const orbital &a, const orbital &b, const int i, const int j, const T zeta, const T zetb, const int n, const T *cab)
	Returns element i,j of virial on atom a.

template<typename T >
__device__ __inline__ double	get_virial_b_normal (const orbital &a, const orbital &b, const int i, const int j, const T zetb, const T rab[3], const int n, const T *cab)
	Returns element i,j of virial on atom b for compute_tau=false.

template<bool compute_tau, typename T >
__device__ __inline__ double	get_virial_b (const orbital &a, const orbital &b, const int i, const int j, const T zeta, const T zetb, const T rab[3], const int n, const T *cab)
	Returns element i,j of virial on atom b.

template<bool compute_tau, typename T >
__device__ __inline__ T	get_hab (const orbital &a, const orbital &b, const T zeta, const T zetb, const int n, const T *cab)
	Returns element i,j of hab matrix.

ldiffs_value	process_get_ldiffs (bool calculate_forces, bool calculate_virial, bool compute_tau)
	Returns difference in angular momentum range for given flags.

Variables
__constant__ orbital	coset_inv [1330]

__constant__ int	binomial_coef [19][19]

Function Documentation

◆ block_to_cab()

template<typename T , bool IS_FUNC_AB>

__device__ __inline__ void rocm_backend::block_to_cab	(	const kernel_params &	params,
		const smem_task< T > &	task,
		T *	cab
	)

Decontracts the subblock, going from spherical to cartesian harmonics.

Definition at line 35 of file grid_gpu_collocate.cu.

Here is the call graph for this function:

◆ __launch_bounds__() [1/4]

template<typename T , bool IS_FUNC_AB>

__global__ rocm_backend::__launch_bounds__ ( 64 ) const

Definition at line 81 of file grid_gpu_collocate.cu.

Here is the call graph for this function:

◆ __launch_bounds__() [2/4]

template<typename T , typename T3 , bool distributed__, bool orthogonal_>

__global__ rocm_backend::__launch_bounds__ ( 64 ) const

Definition at line 150 of file grid_gpu_collocate.cu.

Here is the call graph for this function:

◆ align_up_elems()

constexpr size_t rocm_backend::align_up_elems	(	size_t	n_elems,
		size_t	elem_alignment
	)

constexpr

Definition at line 39 of file grid_gpu_context.cu.

Here is the caller graph for this function:

◆ round_up16()

size_t rocm_backend::round_up16 ( size_t n )

inline

Definition at line 38 of file grid_gpu_context.h.

Here is the caller graph for this function:

◆ block_reduce_64()

template<typename T >

__device__ __inline__ T rocm_backend::block_reduce_64	(	T *	table,
		const T	val_,
		const int	tid
	)

Definition at line 37 of file grid_gpu_integrate.cu.

◆ __launch_bounds__() [3/4]

template<typename T , typename T3 , bool COMPUTE_TAU>

__global__ rocm_backend::__launch_bounds__ ( 64 ) const

Definition at line 128 of file grid_gpu_integrate.cu.

Here is the call graph for this function:

◆ __launch_bounds__() [4/4]

template<typename T , typename T3 , bool distributed__, bool orthogonal_, int lbatch = 20>

__global__ rocm_backend::__launch_bounds__ ( 64 ) const

Cuda kernel for calcualting the coefficients of a potential (density, etc...) for a given pair of gaussian basis sets

We compute the discretized version of the following integral

$$\int _\infty ^\infty V_{ijk} P^\alpha_iP^beta_j P ^ \gamma _ k Exp(- \eta |r_{ijk} - r_c|^2)$$

where in practice the summation is over a finite domain. The discrete form has this shape

$$ \sum_{ijk < dmoain} V_{ijk} (x_i - x_c)^\alpha (y_j - y_c)^\beta (z_k - z_c) ^ \gamma Exp(- \eta |r_{ijk} - r_c|^2) $$

where $0 \le \alpha + \beta + \gamma \le lmax$

It is formely the same operation than collocate (from a discrete point of view) but the implementation differ because of technical reasons.

So most of the code is the same except the core of the routine that is specialized to the integration.

Definition at line 360 of file grid_gpu_integrate.cu.

Here is the call graph for this function:

◆ fac()

__device__ __inline__ double rocm_backend::fac ( const int i )

Factorial function, e.g. fac(5) = 5! = 120.

Author: Ole Schuett

Definition at line 143 of file grid_gpu_internal_header.h.

◆ ncoset()

__host__ __device__ __inline__ int rocm_backend::ncoset ( const int l )

Number of Cartesian orbitals up to given angular momentum quantum.

Author: Ole Schuett

Definition at line 168 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ coset()

__host__ __device__ __inline__ int rocm_backend::coset	(	int	lx,
		int	ly,
		int	lz
	)

Maps three angular momentum components to a single zero based index.

Definition at line 180 of file grid_gpu_internal_header.h.

◆ up()

__device__ __inline__ orbital rocm_backend::up	(	const int	i,
		const orbital &	a
	)

Increase i'th component of given orbital angular momentum.

Definition at line 192 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ down()

__inline__ __device__ orbital rocm_backend::down	(	const int	i,
		const orbital &	a
	)

Decrease i'th component of given orbital angular momentum.

Definition at line 201 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ idx()

__inline__ __device__ int rocm_backend::idx ( const orbital a )

Return coset index of given orbital angular momentum.

Definition at line 210 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ power()

__device__ __inline__ double rocm_backend::power	(	const double	x,
		const int	expo
	)

Definition at line 214 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ prep_term()

template<typename T = double>

__device__ __inline__ void rocm_backend::prep_term	(	const orbital	a,
		const orbital	b,
		const T	value,
		const int	n,
		T *	cab
	)

Adds given value to matrix element cab[idx(b)][idx(a)].

Definition at line 225 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

◆ init_constant_memory()

static void rocm_backend::init_constant_memory ( )

inlinestatic

Initializes the device's constant memory.

Author: Ole Schuett

Definition at line 234 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ block_index()

__inline__ __device__ unsigned int rocm_backend::block_index ( )

Definition at line 287 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ thread_global_index()

__inline__ __device__ unsigned int rocm_backend::thread_global_index ( )

Definition at line 292 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ compute_coordinates() [1/2]

__inline__ __device__ double3 rocm_backend::compute_coordinates	(	const double *__restrict__	dh_,
		const double	x,
		const double	y,
		const double	z
	)

Definition at line 297 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ compute_coordinates() [2/2]

__inline__ __device__ float3 rocm_backend::compute_coordinates	(	const float *__restrict__	dh_,
		const float	x,
		const float	y,
		const float	z
	)

Definition at line 310 of file grid_gpu_internal_header.h.

◆ compute_alpha()

template<typename T >

__inline__ __device__ void rocm_backend::compute_alpha	(	const smem_task< T > &	task,
		T *__restrict__	alpha
	)

Computes the polynomial expansion coefficients: (x-a)**lxa (x-b)**lxb -> sum_{ls} alpha(ls,lxa,lxb,1)*(x-p)**ls.

Definition at line 329 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ convert_to_lattice_coordinates() [1/2]

__host__ __inline__ __device__ void rocm_backend::convert_to_lattice_coordinates	(	const double *	dh_inv_,
		const double3 *__restrict__ const	rp,
		double3 *__restrict__	rp_c
	)

Definition at line 364 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ convert_from_lattice_coordinates_to_cartesian() [1/2]

__host__ __inline__ __device__ void rocm_backend::convert_from_lattice_coordinates_to_cartesian	(	const double *__restrict__	dh_,
		const double3 *__restrict__ const	rp,
		double3 *__restrict__	rp_c
	)

Definition at line 373 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ convert_to_lattice_coordinates() [2/2]

__host__ __inline__ __device__ void rocm_backend::convert_to_lattice_coordinates	(	const float *	dh_inv_,
		const float3 *__restrict__ const	rp,
		float3 *__restrict__	rp_c
	)

Definition at line 382 of file grid_gpu_internal_header.h.

◆ convert_from_lattice_coordinates_to_cartesian() [2/2]

__host__ __inline__ __device__ void rocm_backend::convert_from_lattice_coordinates_to_cartesian	(	const float *__restrict__	dh_,
		const float3 *__restrict__ const	rp,
		float3 *__restrict__	rp_c
	)

Definition at line 391 of file grid_gpu_internal_header.h.

◆ compute_cube_properties()

template<typename T , typename T3 , bool orthorhombic_>

__inline__ T rocm_backend::compute_cube_properties	(	const T	radius,
		const T *const __restrict__	dh_,
		const T *const __restrict__	dh_inv_,
		const T3 *__restrict__	rp,
		T3 *__restrict__	roffset,
		int3 *__restrict__	cubecenter,
		int3 *__restrict__	lb_cube,
		int3 *__restrict__	cube_size
	)

Definition at line 400 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

◆ compute_window_size()

__inline__ __device__ void rocm_backend::compute_window_size	(	const int3	grid_size,
		const int	border_mask,
		const int3 &	border_width,
		int3 &	window_size,
		int3 &	window_shift
	)

Definition at line 538 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ wrap_grid_index()

__inline__ __device__ int rocm_backend::wrap_grid_index	(	const int	idx,
		const int	full_size
	)

Wraps a cube-local coordinate into the periodic full grid along one axis. Single point of entry shared by collocate_kernel and integrate_kernel.

Definition at line 570 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ setup_task_cube_center()

template<typename T , typename T3 , bool distributed__>

__device__ __inline__ void rocm_backend::setup_task_cube_center	(	const kernel_params &	dev_,
		smem_task_reduced< T, T3 > &	task
	)

Shifts task.cube_center into the local grid's coordinate system and, for distributed grids, derives the border window. Single point of entry shared by collocate_kernel and integrate_kernel so the two cannot drift apart on which axis maps to which.

Definition at line 585 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

◆ cab_to_cxyz()

template<typename T >

__device__ static __inline__ void rocm_backend::cab_to_cxyz	(	const smem_task< T > &	task,
		const T *__restrict__	alpha,
		const T *__restrict__	cab,
		T *__restrict__	cxyz
	)

static

Transforms coefficients C_ab into C_xyz.

Definition at line 606 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ cxyz_to_cab()

template<typename T >

__device__ static __inline__ void rocm_backend::cxyz_to_cab	(	const smem_task< T > &	task,
		const T *__restrict__	alpha,
		const T *__restrict__	cxyz,
		T *__restrict__	cab
	)

static

Transforms coefficients C_xyz into C_ab.

Definition at line 656 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ fill_smem_task_reduced()

template<typename T , typename T3 >

__device__ __inline__ void rocm_backend::fill_smem_task_reduced	(	const kernel_params &	dev,
		const int	task_id,
		smem_task_reduced< T, T3 > &	task
	)

Copies a task from global to shared memory.

Definition at line 708 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ fill_smem_task_coef()

template<typename T >

__device__ __inline__ void rocm_backend::fill_smem_task_coef	(	const kernel_params &	dev,
		const int	task_id,
		smem_task< T > &	task
	)

Copies a task from global to shared memory and does precomputations.

Definition at line 762 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ allocate_workspace()

template<typename T >

__inline__ __device__ T * rocm_backend::allocate_workspace ( const kernel_params & dev_ )

Definition at line 899 of file grid_gpu_internal_header.h.

Here is the call graph for this function:

◆ calculate_ymix_ymax_boundaries()

template<typename T , typename T3 >

__device__ __inline__ T rocm_backend::calculate_ymix_ymax_boundaries	(	smem_task_reduced< T, T3 > &	task,
		const int	z,
		int &	ymin,
		int &	ymax
	)

Definition at line 906 of file grid_gpu_internal_header.h.

Here is the caller graph for this function:

◆ calculate_xmin_xmax_boundaries()

template<typename T , typename T3 >

__device__ __inline__ void rocm_backend::calculate_xmin_xmax_boundaries	(	smem_task_reduced< T, T3 > &	task,
		const int	y,
		const T	kremain,
		int &	xmin,
		int &	xmax
	)

Definition at line 922 of file grid_gpu_internal_header.h.

◆ prepare_pab_AB()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab_AB	(	const orbital	a,
		const orbital	b,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_AB, ie. identity transformation.

Definition at line 28 of file grid_gpu_prepare_pab.h.

◆ prepare_pab_DADB()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab_DADB	(	const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_DADB.

Definition at line 41 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_ADBmDAB()

template<typename T >

__device__ void rocm_backend::prepare_pab_ADBmDAB	(	const int	idir,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_ADBmDAB_{X,Y,Z}.

Definition at line 62 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_ARDBmDARB()

template<typename T >

__device__ void rocm_backend::prepare_pab_ARDBmDARB	(	const int	idir,
		const int	ir,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_ARDBmDARB_{X,Y,Z}{X,Y,Z}.

Definition at line 83 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_DABpADB()

template<typename T >

__device__ void rocm_backend::prepare_pab_DABpADB	(	const int	idir,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_DABpADB_{X,Y,Z}.

Author: Ole Schuett

Definition at line 108 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_Di()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab_Di	(	const int	ider,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_{DX,DY,DZ}.

Author: Ole Schuett

Definition at line 131 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ oneterm_dijdij()

template<typename T >

__device__ __inline__ void rocm_backend::oneterm_dijdij	(	const int	idir,
		const T	func_a,
		const orbital	a,
		const orbital	b,
		const T	zetb,
		const int	n,
		T *	cab
	)

Helper for grid_prepare_pab_DiDj.

Definition at line 155 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_DiDj()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab_DiDj	(	const int	ider1,
		const int	ider2,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_{DXDY,DYDZ,DZDX}.

Author: Ole Schuett

Definition at line 185 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ oneterm_diidii()

template<typename T >

__device__ void rocm_backend::oneterm_diidii	(	const int	idir,
		const T	func_a,
		const orbital	a,
		const orbital	b,
		const T	zetb,
		const int	n,
		T *	cab
	)

Helper for grid_prepare_pab_Di2.

Author: Ole Schuett

Definition at line 212 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab_Di2()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab_Di2	(	const int	ider,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Implementation of function GRID_FUNC_{DXDX,DYDY,DZDZ}.

Definition at line 227 of file grid_gpu_prepare_pab.h.

Here is the call graph for this function:

◆ prepare_pab()

template<typename T >

__device__ __inline__ void rocm_backend::prepare_pab	(	const enum grid_func	func,
		const orbital	a,
		const orbital	b,
		const T	zeta,
		const T	zetb,
		const T	pab_val,
		const int	n,
		T *	cab
	)

Transforms a given element of the density matrix according to func.

Parameters

func	Transformation function to apply, one of GRID_FUNC_*.
{a,b}	Orbital angular momenta.
zet_{a,b}	Gaussian exponents.
pab_val	Input matrix element of pab.
n	Leading dimensions of output matrix cab.
cab	Output matrix.

Definition at line 255 of file grid_gpu_prepare_pab.h.

Here is the caller graph for this function:

◆ prepare_get_ldiffs()

ldiffs_value rocm_backend::prepare_get_ldiffs ( const enum grid_func func )

inline

Returns difference in angular momentum range for given func.

Author: Ole Schuett

Definition at line 347 of file grid_gpu_prepare_pab.h.

Here is the caller graph for this function:

◆ get_term()

template<typename T >

__device__ __inline__ T rocm_backend::get_term	(	const orbital &	a,
		const orbital &	b,
		const int	n,
		const T *	cab
	)

Returns matrix element cab[idx(b)][idx(a)].

Author: Ole Schuett

Definition at line 27 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ get_force_a_normal()

template<typename T >

__device__ __inline__ T rocm_backend::get_force_a_normal	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const T	zeta,
		const int	n,
		const T *	cab
	)

Returns i'th component of force on atom a for compute_tau=false.

Author: Ole Schuett

Definition at line 37 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ get_force_a()

template<bool compute_tau, typename T >

__device__ __inline__ double rocm_backend::get_force_a	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const T	zeta,
		const T	zetb,
		const int	n,
		const T *	cab
	)

Returns i'th component of force on atom a.

Author: Ole Schuett

Definition at line 51 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

◆ get_force_b_normal()

template<typename T >

__device__ __inline__ double rocm_backend::get_force_b_normal	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const T	zetb,
		const T	rab[3],
		const int	n,
		const T *	cab
	)

Returns i'th component of force on atom b for compute_tau=false.

Author: Ole Schuett

Definition at line 81 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ get_force_b()

template<bool compute_tau, typename T >

__device__ __inline__ T rocm_backend::get_force_b	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const T	zeta,
		const T	zetb,
		const T	rab[3],
		const int	n,
		const T *	cab
	)

Returns i'th component of force on atom b.

Author: Ole Schuett

Definition at line 94 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

◆ get_virial_a_normal()

template<typename T >

__device__ __inline__ double rocm_backend::get_virial_a_normal	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const int	j,
		const T	zeta,
		const int	n,
		const T *	cab
	)

Returns element i,j of virial on atom a for compute_tau=false.

Author: Ole Schuett

Definition at line 125 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ get_virial_a()

template<bool compute_tau, typename T >

__device__ __inline__ T rocm_backend::get_virial_a	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const int	j,
		const T	zeta,
		const T	zetb,
		const int	n,
		const T *	cab
	)

Returns element i,j of virial on atom a.

Author: Ole Schuett

Definition at line 136 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

◆ get_virial_b_normal()

template<typename T >

__device__ __inline__ double rocm_backend::get_virial_b_normal	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const int	j,
		const T	zetb,
		const T	rab[3],
		const int	n,
		const T *	cab
	)

Returns element i,j of virial on atom b for compute_tau=false.

Author: Ole Schuett

Definition at line 168 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

Here is the caller graph for this function:

◆ get_virial_b()

template<bool compute_tau, typename T >

__device__ __inline__ double rocm_backend::get_virial_b	(	const orbital &	a,
		const orbital &	b,
		const int	i,
		const int	j,
		const T	zeta,
		const T	zetb,
		const T	rab[3],
		const int	n,
		const T *	cab
	)

Returns element i,j of virial on atom b.

Author: Ole Schuett

Definition at line 186 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

◆ get_hab()

template<bool compute_tau, typename T >

__device__ __inline__ T rocm_backend::get_hab	(	const orbital &	a,
		const orbital &	b,
		const T	zeta,
		const T	zetb,
		const int	n,
		const T *	cab
	)

Returns element i,j of hab matrix.

Author: Ole Schuett

Definition at line 217 of file grid_gpu_process_vab.h.

Here is the call graph for this function:

◆ process_get_ldiffs()

ldiffs_value rocm_backend::process_get_ldiffs	(	bool	calculate_forces,
		bool	calculate_virial,
		bool	compute_tau
	)

inline

Returns difference in angular momentum range for given flags.

Definition at line 241 of file grid_gpu_process_vab.h.

Here is the caller graph for this function:

Variable Documentation

◆ coset_inv

__constant__ orbital rocm_backend::coset_inv[1330]

Definition at line 62 of file grid_gpu_internal_header.h.

◆ binomial_coef

__constant__ int rocm_backend::binomial_coef[19][19]

Definition at line 63 of file grid_gpu_internal_header.h.

Data Structures

Functions

Variables

Function Documentation

◆ block_to_cab()

◆ __launch_bounds__() [1/4]

◆ __launch_bounds__() [2/4]

◆ align_up_elems()

◆ round_up16()

◆ block_reduce_64()

◆ __launch_bounds__() [3/4]

◆ __launch_bounds__() [4/4]

◆ fac()

◆ ncoset()

◆ coset()

◆ up()

◆ down()

◆ idx()

◆ power()

◆ prep_term()

◆ init_constant_memory()

◆ block_index()

◆ thread_global_index()

◆ compute_coordinates() [1/2]

◆ compute_coordinates() [2/2]

◆ compute_alpha()

◆ convert_to_lattice_coordinates() [1/2]

◆ convert_from_lattice_coordinates_to_cartesian() [1/2]

◆ convert_to_lattice_coordinates() [2/2]

◆ convert_from_lattice_coordinates_to_cartesian() [2/2]

◆ compute_cube_properties()

◆ compute_window_size()

◆ wrap_grid_index()

◆ setup_task_cube_center()

◆ cab_to_cxyz()

◆ cxyz_to_cab()

◆ fill_smem_task_reduced()

◆ fill_smem_task_coef()

◆ allocate_workspace()

◆ calculate_ymix_ymax_boundaries()

◆ calculate_xmin_xmax_boundaries()

◆ prepare_pab_AB()

◆ prepare_pab_DADB()

◆ prepare_pab_ADBmDAB()

◆ prepare_pab_ARDBmDARB()

◆ prepare_pab_DABpADB()

◆ prepare_pab_Di()

◆ oneterm_dijdij()

◆ prepare_pab_DiDj()

◆ oneterm_diidii()

◆ prepare_pab_Di2()

◆ prepare_pab()

◆ prepare_get_ldiffs()

◆ get_term()

◆ get_force_a_normal()

◆ get_force_a()

◆ get_force_b_normal()

◆ get_force_b()

◆ get_virial_a_normal()

◆ get_virial_a()

◆ get_virial_b_normal()

◆ get_virial_b()

◆ get_hab()

◆ process_get_ldiffs()

Variable Documentation

◆ coset_inv

◆ binomial_coef