![]() |
(git:b77b4be)
|
Data Structures | |
class | context_info |
class | gpu_vector |
class | grid_info |
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 | |
__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__ 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 int *const grid_size, const int border_mask, const int *border_width, int3 *const window_size, int3 *const window_shift) |
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 > | |
__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] |
__device__ __inline__ double rocm_backend::fac | ( | const int | i | ) |
Factorial function, e.g. fac(5) = 5! = 120.
Definition at line 139 of file grid_hip_internal_header.h.
__host__ __device__ __inline__ int rocm_backend::ncoset | ( | const int | l | ) |
Number of Cartesian orbitals up to given angular momentum quantum.
Definition at line 164 of file grid_hip_internal_header.h.
__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 176 of file grid_hip_internal_header.h.
Increase i'th component of given orbital angular momentum.
Definition at line 188 of file grid_hip_internal_header.h.
Decrease i'th component of given orbital angular momentum.
Definition at line 197 of file grid_hip_internal_header.h.
__inline__ __device__ int rocm_backend::idx | ( | const orbital | a | ) |
Return coset index of given orbital angular momentum.
Definition at line 206 of file grid_hip_internal_header.h.
__device__ __inline__ double rocm_backend::power | ( | const double | x, |
const int | expo | ||
) |
Definition at line 210 of file grid_hip_internal_header.h.
__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 221 of file grid_hip_internal_header.h.
|
inlinestatic |
Initializes the device's constant memory.
Definition at line 230 of file grid_hip_internal_header.h.
__inline__ __device__ double3 rocm_backend::compute_coordinates | ( | const double *__restrict__ | dh_, |
const double | x, | ||
const double | y, | ||
const double | z | ||
) |
Definition at line 288 of file grid_hip_internal_header.h.
__inline__ __device__ float3 rocm_backend::compute_coordinates | ( | const float *__restrict__ | dh_, |
const float | x, | ||
const float | y, | ||
const float | z | ||
) |
Definition at line 301 of file grid_hip_internal_header.h.
__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 320 of file grid_hip_internal_header.h.
__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 356 of file grid_hip_internal_header.h.
__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 365 of file grid_hip_internal_header.h.
__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 374 of file grid_hip_internal_header.h.
__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 383 of file grid_hip_internal_header.h.
__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 392 of file grid_hip_internal_header.h.
__inline__ __device__ void rocm_backend::compute_window_size | ( | const int *const | grid_size, |
const int | border_mask, | ||
const int * | border_width, | ||
int3 *const | window_size, | ||
int3 *const | window_shift | ||
) |
Definition at line 526 of file grid_hip_internal_header.h.
|
static |
Transforms coefficients C_ab into C_xyz.
Definition at line 558 of file grid_hip_internal_header.h.
|
static |
Transforms coefficients C_xyz into C_ab.
Definition at line 609 of file grid_hip_internal_header.h.
__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 661 of file grid_hip_internal_header.h.
__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 701 of file grid_hip_internal_header.h.
__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 25 of file grid_hip_prepare_pab.h.
__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 38 of file grid_hip_prepare_pab.h.
__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 59 of file grid_hip_prepare_pab.h.
__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 80 of file grid_hip_prepare_pab.h.
__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}.
Definition at line 105 of file grid_hip_prepare_pab.h.
__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}.
Definition at line 128 of file grid_hip_prepare_pab.h.
__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 152 of file grid_hip_prepare_pab.h.
__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}.
Definition at line 182 of file grid_hip_prepare_pab.h.
__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.
Definition at line 209 of file grid_hip_prepare_pab.h.
__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 224 of file grid_hip_prepare_pab.h.
__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.
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 252 of file grid_hip_prepare_pab.h.
|
inline |
Returns difference in angular momentum range for given func.
Definition at line 344 of file grid_hip_prepare_pab.h.
__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)].
Definition at line 27 of file grid_hip_process_vab.h.
__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.
Definition at line 37 of file grid_hip_process_vab.h.
__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.
Definition at line 51 of file grid_hip_process_vab.h.
__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.
Definition at line 81 of file grid_hip_process_vab.h.
__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.
Definition at line 94 of file grid_hip_process_vab.h.
__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.
Definition at line 125 of file grid_hip_process_vab.h.
__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.
Definition at line 136 of file grid_hip_process_vab.h.
__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.
Definition at line 168 of file grid_hip_process_vab.h.
__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.
Definition at line 186 of file grid_hip_process_vab.h.
__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.
Definition at line 217 of file grid_hip_process_vab.h.
|
inline |
Returns difference in angular momentum range for given flags.
Definition at line 241 of file grid_hip_process_vab.h.
__constant__ orbital rocm_backend::coset_inv[1330] |
Definition at line 63 of file grid_hip_internal_header.h.
__constant__ int rocm_backend::binomial_coef[19][19] |
Definition at line 64 of file grid_hip_internal_header.h.