Skip to content

Commit 275c909

Browse files
kunkunblueberrykunkunblueberry
committed
fix:some templates are not declared and force.cpp
1 parent 83804a3 commit 275c909

5 files changed

Lines changed: 97 additions & 15 deletions

File tree

‎source/module_hamilt_lcao/hamilt_lcaodft/operator_lcao/op_exx_lcao.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -240,7 +240,7 @@ OperatorEXX<OperatorLCAO<TK, TR>>::OperatorEXX(HS_Matrix_K<TK>* hsk_in,
240240
else if (this->add_hexx_type == Add_Hexx_Type::R)
241241
{
242242
// read in Hexx(R)
243-
const std::string restart_HR_path = PARAM.globalv.global_readin_dir + "HexxR" + std::to_string(PARAM.globalv.myrank);
243+
const std::string restart_HR_path = GlobalC::restart.folder + "HexxR" + std::to_string(PARAM.globalv.myrank);
244244
bool all_exist = true;
245245
for (int is = 0; is < PARAM.inp.nspin; ++is)
246246
{

‎source/module_hamilt_pw/hamilt_pwdft/forces.cpp‎

Lines changed: 56 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -16,7 +16,7 @@
1616
#include "module_hamilt_general/module_surchem/surchem.h"
1717
#include "module_hamilt_general/module_vdw/vdw.h"
1818
#include "kernels/force_op.h"
19-
19+
#include <type_traits>
2020
#ifdef _OPENMP
2121
#include <omp.h>
2222
#endif
@@ -579,7 +579,7 @@ void Forces<FPTYPE, Device>::cal_force_loc(const UnitCell& ucell,
579579
syncmem_var_h2d_op()(this->ctx, this->cpu_ctx, forcelc_d, forcelc.c, this->nat * 3);
580580
syncmem_var_h2d_op()(this->ctx, this->cpu_ctx, vloc_d, vloc.c, vloc.nr * vloc.nc);
581581

582-
hamilt::cal_force_loc_op<FPTYPE, Device>()(
582+
/* hamilt::cal_force_loc_op<FPTYPE, Device>()(
583583
this->nat,
584584
rho_basis->npw,
585585
ucell.tpiba * ucell.omega,
@@ -590,7 +590,34 @@ void Forces<FPTYPE, Device>::cal_force_loc(const UnitCell& ucell,
590590
aux_d,
591591
vloc_d,
592592
vloc.nc,
593-
forcelc_d);
593+
forcelc_d);*/
594+
if constexpr (std::is_same<Device, base_device::DEVICE_GPU>::value) {
595+
hamilt::cal_force_loc_sincos_op<FPTYPE, Device>()(
596+
this->ctx,
597+
this->nat,
598+
rho_basis->npw,
599+
ucell.ntype,
600+
gcar_d,
601+
tau_d,
602+
vloc_d,
603+
aux_d,
604+
static_cast<FPTYPE>(ucell.tpiba * ucell.omega),
605+
forcelc_d);
606+
} else {
607+
hamilt::cal_force_loc_op<FPTYPE, Device>()(
608+
this->nat,
609+
rho_basis->npw,
610+
ucell.tpiba * ucell.omega,
611+
iat2it_d,
612+
ig2gg_d,
613+
gcar_d,
614+
tau_d,
615+
aux_d,
616+
vloc_d,
617+
vloc.nc,
618+
forcelc_d);
619+
}
620+
594621
syncmem_var_d2h_op()(this->cpu_ctx, this->ctx, forcelc.c, forcelc_d, this->nat * 3);
595622

596623
delmem_int_op()(this->ctx,iat2it_d);
@@ -788,7 +815,7 @@ void Forces<FPTYPE, Device>::cal_force_ew(const UnitCell& ucell,
788815
syncmem_complex_h2d_op()(this->ctx, this->cpu_ctx, aux_d, aux.data(), rho_basis->npw);
789816
syncmem_var_h2d_op()(this->ctx, this->cpu_ctx, forceion_d, forceion.c, this->nat * 3);
790817

791-
hamilt::cal_force_ew_op<FPTYPE, Device>()(
818+
/* hamilt::cal_force_ew_op<FPTYPE, Device>()(
792819
this->nat,
793820
rho_basis->npw,
794821
rho_basis->ig_gge0,
@@ -798,7 +825,31 @@ void Forces<FPTYPE, Device>::cal_force_ew(const UnitCell& ucell,
798825
it_fact_d,
799826
aux_d,
800827
forceion_d);
801-
828+
*/
829+
if constexpr (std::is_same<Device, base_device::DEVICE_GPU>::value) {
830+
hamilt::cal_force_ew_sincos_op<FPTYPE, Device>()(
831+
this->ctx,
832+
this->nat,
833+
rho_basis->npw,
834+
rho_basis->ig_gge0,
835+
gcar_d,
836+
tau_d,
837+
it_fact_d,
838+
aux_d,
839+
forceion_d);
840+
} else {
841+
hamilt::cal_force_ew_op<FPTYPE, Device>()(
842+
this->nat,
843+
rho_basis->npw,
844+
rho_basis->ig_gge0,
845+
iat2it_d,
846+
gcar_d,
847+
tau_d,
848+
it_fact_d,
849+
aux_d,
850+
forceion_d);
851+
}
852+
802853
syncmem_var_d2h_op()(this->cpu_ctx, this->ctx, forceion.c, forceion_d, this->nat * 3);
803854
delmem_int_op()(this->ctx,iat2it_d);
804855
delmem_var_op()(this->ctx,gcar_d);

‎source/module_hamilt_pw/hamilt_pwdft/kernels/force_op.h‎

Lines changed: 30 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -179,6 +179,9 @@ struct cal_force_ew_op{
179179
FPTYPE* forceion
180180
) {};
181181
};
182+
183+
template <typename FPTYPE, typename Device> struct cal_force_loc_sincos_op;
184+
template <typename FPTYPE, typename Device> struct cal_force_ew_sincos_op;
182185
#if __CUDA || __UT_USE_CUDA || __ROCM || __UT_USE_ROCM
183186
template <typename FPTYPE>
184187
struct cal_vkb1_nl_op<FPTYPE, base_device::DEVICE_GPU>
@@ -335,6 +338,32 @@ struct cal_force_ew_op<FPTYPE, base_device::DEVICE_GPU>{
335338
FPTYPE* forceion
336339
);
337340
};
341+
template <typename FPTYPE>
342+
struct cal_force_loc_sincos_op<FPTYPE, base_device::DEVICE_GPU> {
343+
void operator()(const base_device::DEVICE_GPU* ctx,
344+
const int& nat,
345+
const int& npw,
346+
const int& ntype,
347+
const FPTYPE* gcar,
348+
const FPTYPE* tau,
349+
const FPTYPE* vloc_per_type,
350+
const std::complex<FPTYPE>* aux,
351+
const FPTYPE& scale_factor,
352+
FPTYPE* force);
353+
};
354+
355+
template <typename FPTYPE>
356+
struct cal_force_ew_sincos_op<FPTYPE, base_device::DEVICE_GPU> {
357+
void operator()(const base_device::DEVICE_GPU* ctx,
358+
const int& nat,
359+
const int& npw,
360+
const int& ig_gge0,
361+
const FPTYPE* gcar,
362+
const FPTYPE* tau,
363+
const FPTYPE* it_facts,
364+
const std::complex<FPTYPE>* aux,
365+
FPTYPE* force);
366+
};
338367
#endif // __CUDA || __UT_USE_CUDA || __ROCM || __UT_USE_ROCM
339368
} // namespace hamilt
340-
#endif // W_ABACUS_DEVELOP_ABACUS_DEVELOP_SOURCE_source_pw_HAMILT_PWDFT_KERNELS_FORCE_OP_H
369+
#endif // W_ABACUS_DEVELOP_ABACUS_DEVELOP_SOURCE_source_pw_HAMILT_PWDFT_KERNELS_FORCE_OP_H

‎source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/force_op.hip.cu‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,8 @@
1010

1111
namespace hamilt {
1212

13+
__device__ __forceinline__ void sincos_(float x, float* s, float* c) { sincosf(x, s, c); }
14+
__device__ __forceinline__ void sincos_(double x, double* s, double* c) { sincos(x, s, c); }
1315
template <typename FPTYPE>
1416
__global__ void cal_vkb1_nl(
1517
const int npwx,
@@ -658,7 +660,7 @@ __global__ void cal_force_loc_sincos_kernel(
658660

659661
// Use HIP intrinsic for sincos
660662
FPTYPE sinp, cosp;
661-
sincos(phase, &sinp, &cosp);
663+
sincos_(phase, &sinp, &cosp);
662664

663665
// Calculate force factor
664666
const FPTYPE vloc_factor = vloc_per_type[iat * npw + ig];
@@ -718,7 +720,7 @@ __global__ void cal_force_ew_sincos_kernel(
718720

719721
// Use HIP intrinsic for sincos
720722
FPTYPE sinp, cosp;
721-
sincos(phase, &sinp, &cosp);
723+
sincos_(phase, &sinp, &cosp);
722724

723725
// Calculate Ewald sum contribution (fixed sign error)
724726
const FPTYPE factor = it_fact * (-cosp * aux[ig].imag() + sinp * aux[ig].real());

‎source/module_hsolver/kernels/rocm/dngvd_op.hip.cu‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -128,8 +128,8 @@ void dngvd_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const ba
128128
hipsolverErrcheck(hipsolverDnChegvd_bufferSize(
129129
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
130130
nstart,
131-
reinterpret_cast<const float2 *>(_vcc), ldh,
132-
reinterpret_cast<const float2 *>(_scc), ldh,
131+
const_cast<float2*>(reinterpret_cast<const float2 *>(_vcc)), ldh,
132+
const_cast<float2*>(reinterpret_cast<const float2 *>(_scc)), ldh,
133133
_eigenvalue,
134134
&lwork));
135135

@@ -140,7 +140,7 @@ void dngvd_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const ba
140140
hipsolverErrcheck(hipsolverDnChegvd(
141141
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
142142
nstart,
143-
reinterpret_cast<float2 *>(_vcc), ldh,
143+
const_cast<float2*>(reinterpret_cast<float2 *>(_vcc)), ldh,
144144
const_cast<float2 *>(reinterpret_cast<const float2 *>(_scc)), ldh,
145145
_eigenvalue,
146146
work, lwork, devInfo));
@@ -206,8 +206,8 @@ void dngvd_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const b
206206
hipsolverErrcheck(hipsolverDnZhegvd_bufferSize(
207207
hipsolver_H, HIPSOLVER_EIG_TYPE_1, HIPSOLVER_EIG_MODE_VECTOR, uplo,
208208
nstart,
209-
reinterpret_cast<const double2 *>(_vcc), ldh,
210-
reinterpret_cast<const double2 *>(_scc), ldh,
209+
const_cast<double2*>(reinterpret_cast<const double2 *>(_vcc)), ldh,
210+
const_cast<double2*>(reinterpret_cast<const double2 *>(_scc)), ldh,
211211
_eigenvalue,
212212
&lwork));
213213

@@ -365,4 +365,4 @@ void dngvx_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DE
365365
}
366366
#endif // __LCAO
367367

368-
} // namespace hsolver
368+
} // namespace hsolver

0 commit comments

Comments
 (0)