Skip to content

Fix/gpu elecstate pw op spin stride #7776

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
chengleizheng wants to merge 4 commits into
base: develop
Choose a base branch
from
Open

Fix/gpu elecstate pw op spin stride #7776

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
bc21c53
Fix(GPU): correct spin stride of rho in elecstate_pw_op kernel
Aug 4, 2026
9654edb
Merge branch 'deepmodeling:develop' into fix/gpu-elecstate-pw-op-spin…
chengleizheng Aug 4, 2026
96ef4cd
Fix(Test): add rho_stride argument to elecstate_op unit test calls
Aug 5, 2026
07bc74a
rename rho_stride to nrxx_dense in elecstate_pw_op
Aug 5, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 3 additions & 2 deletions source/source_estate/elecstate_pw.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -227,6 +227,7 @@ void ElecStatePW<T, Device>::rhoBandK(const psi::Psi<T, Device>& psi)
PARAM.globalv.domag,
PARAM.globalv.domag_z,
this->basis->nrxx,
this->charge->nrxx,
w1,
this->rho,
this->wfcr,
Expand All @@ -249,7 +250,7 @@ void ElecStatePW<T, Device>::rhoBandK(const psi::Psi<T, Device>& psi)
if (w1 != 0.0)
{
// replaced by denghui at 20221110
elecstate_pw_op()(this->ctx, current_spin, this->basis->nrxx, w1, this->rho, this->wfcr);
elecstate_pw_op()(this->ctx, current_spin, this->basis->nrxx, this->charge->nrxx, w1, this->rho, this->wfcr);
}

// kinetic energy density
Expand All @@ -272,7 +273,7 @@ void ElecStatePW<T, Device>::rhoBandK(const psi::Psi<T, Device>& psi)

this->basis->recip_to_real(this->ctx, this->wfcr, this->wfcr, ik);

elecstate_pw_op()(this->ctx, current_spin, this->charge->nrxx, w1, this->kin_r, this->wfcr);
elecstate_pw_op()(this->ctx, current_spin, this->charge->nrxx, this->charge->nrxx, w1, this->kin_r, this->wfcr);
}
}
}
Expand Down
2 changes: 1 addition & 1 deletion source/source_estate/elecstate_pw_cal_tau.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ void ElecStatePW<T, Device>::cal_tau(const psi::Psi<T, Device>& psi)

this->basis->recip_to_real(this->ctx, this->wfcr, this->wfcr, ik);

elecstate_pw_op()(this->ctx, current_spin, this->charge->nrxx, w1, this->kin_r, this->wfcr);
elecstate_pw_op()(this->ctx, current_spin, this->charge->nrxx, this->charge->nrxx, w1, this->kin_r, this->wfcr);
}
}
}
Expand Down
38 changes: 22 additions & 16 deletions source/source_estate/kernels/cuda/elecstate_op.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,11 +14,12 @@ __global__ void elecstate_pw(
const int nrxx,
const FPTYPE w1,
FPTYPE* rho,
const thrust::complex<FPTYPE>* wfcr)
const thrust::complex<FPTYPE>* wfcr,
const int nrxx_dense)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if(idx >= nrxx) {return;}
rho[spin * nrxx + idx] += w1 * norm(wfcr[idx]);
rho[spin * nrxx_dense + idx] += w1 * norm(wfcr[idx]);
}

template<typename FPTYPE>
Expand All @@ -29,46 +30,49 @@ __global__ void elecstate_pw(
const FPTYPE w1,
FPTYPE* rho,
const thrust::complex<FPTYPE>* wfcr,
const thrust::complex<FPTYPE>* wfcr_another_spin)
const thrust::complex<FPTYPE>* wfcr_another_spin,
const int nrxx_dense)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if(idx >= nrxx) {return;}
rho[0 * nrxx + idx] += w1 * (norm(wfcr[idx]) + norm(wfcr_another_spin[idx]));
rho[0 * nrxx_dense + idx] += w1 * (norm(wfcr[idx]) + norm(wfcr_another_spin[idx]));

if (DOMAG) {
rho[1 * nrxx + idx] += w1 * 2.0
rho[1 * nrxx_dense + idx] += w1 * 2.0
* (wfcr[idx].real() * wfcr_another_spin[idx].real()
+ wfcr[idx].imag() * wfcr_another_spin[idx].imag());
rho[2 * nrxx + idx] += w1 * 2.0
rho[2 * nrxx_dense + idx] += w1 * 2.0
* (wfcr[idx].real() * wfcr_another_spin[idx].imag()
- wfcr_another_spin[idx].real() * wfcr[idx].imag());
rho[3 * nrxx + idx] += w1 * (norm(wfcr[idx]) - norm(wfcr_another_spin[idx]));
rho[3 * nrxx_dense + idx] += w1 * (norm(wfcr[idx]) - norm(wfcr_another_spin[idx]));
}
else if(DOMAG_Z) {
rho[1 * nrxx + idx] = 0;
rho[2 * nrxx + idx] = 0;
rho[3 * nrxx + idx] += w1 * (norm(wfcr[idx]) - norm(wfcr_another_spin[idx]));
rho[1 * nrxx_dense + idx] = 0;
rho[2 * nrxx_dense + idx] = 0;
rho[3 * nrxx_dense + idx] += w1 * (norm(wfcr[idx]) - norm(wfcr_another_spin[idx]));
}
else {
rho[0 * nrxx + idx] = 0;
rho[1 * nrxx + idx] = 0;
rho[2 * nrxx + idx] = 0;
rho[3 * nrxx + idx] = 0;
rho[0 * nrxx_dense + idx] = 0;
rho[1 * nrxx_dense + idx] = 0;
rho[2 * nrxx_dense + idx] = 0;
rho[3 * nrxx_dense + idx] = 0;
}
}

template <typename FPTYPE>
void elecstate_pw_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* ctx,
const int& spin,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr)
{
const int block = (nrxx + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK;
elecstate_pw<FPTYPE><<<block, THREADS_PER_BLOCK>>>(
spin, nrxx, w1, rho[0],
reinterpret_cast<const thrust::complex<FPTYPE>*>(wfcr)
reinterpret_cast<const thrust::complex<FPTYPE>*>(wfcr),
nrxx_dense
);

CHECK_CUDA_SYNC();
Expand All @@ -79,6 +83,7 @@ void elecstate_pw_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_dev
const bool& DOMAG,
const bool& DOMAG_Z,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr,
Expand All @@ -88,7 +93,8 @@ void elecstate_pw_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_dev
elecstate_pw<FPTYPE><<<block, THREADS_PER_BLOCK>>>(
DOMAG, DOMAG_Z, nrxx, w1, rho[0],
reinterpret_cast<const thrust::complex<FPTYPE>*>(wfcr),
reinterpret_cast<const thrust::complex<FPTYPE>*>(wfcr_another_spin)
reinterpret_cast<const thrust::complex<FPTYPE>*>(wfcr_another_spin),
nrxx_dense
);

CHECK_CUDA_SYNC();
Expand Down
2 changes: 2 additions & 0 deletions source/source_estate/kernels/elecstate_op.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@ struct elecstate_pw_op<FPTYPE, base_device::DEVICE_CPU>
void operator()(const base_device::DEVICE_CPU* /*ctx*/,
const int& spin,
const int& nrxx,
const int& /*nrxx_dense*/,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr)
Expand All @@ -29,6 +30,7 @@ struct elecstate_pw_op<FPTYPE, base_device::DEVICE_CPU>
const bool& DOMAG,
const bool& DOMAG_Z,
const int& nrxx,
const int& /*nrxx_dense*/,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr,
Expand Down
10 changes: 8 additions & 2 deletions source/source_estate/kernels/elecstate_op.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,8 @@ struct elecstate_pw_op {
/// Input Parameters
/// @param ctx - which device this function runs on
/// @param spin - current spin
/// @param nrxx - number of planewaves
/// @param nrxx - number of real-space grid points on this process (basis/wfc grid)
/// @param nrxx_dense - nrxx of the dense (charge) grid, used as the stride between spin components of rho
/// @param weight - input constant
/// @param wfcr - input array, psi in real space
///
Expand All @@ -24,6 +25,7 @@ struct elecstate_pw_op {
const Device* ctx,
const int& spin,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& weight,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr);
Expand All @@ -34,7 +36,8 @@ struct elecstate_pw_op {
/// @param ctx - which device this function runs on
/// @param DOMAG - PARAM.globalv.domag
/// @param DOMAG_Z - PARAM.globalv.domag_z
/// @param nrxx - number of planewaves
/// @param nrxx - number of real-space grid points on this process (basis/wfc grid)
/// @param nrxx_dense - nrxx of the dense (charge) grid, used as the stride between spin components of rho
/// @param weight - input constant
/// @param wfcr - input array, psi in real space
/// @param wfcr_another_spin - input array, psi in real space
Expand All @@ -46,6 +49,7 @@ struct elecstate_pw_op {
const bool& DOMAG,
const bool& DOMAG_Z,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& weight,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr,
Expand All @@ -59,6 +63,7 @@ struct elecstate_pw_op<FPTYPE, base_device::DEVICE_GPU>
void operator()(const base_device::DEVICE_GPU* ctx,
const int& spin,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr);
Expand All @@ -67,6 +72,7 @@ struct elecstate_pw_op<FPTYPE, base_device::DEVICE_GPU>
const bool& DOMAG,
const bool& DOMAG_Z,
const int& nrxx,
const int& nrxx_dense,
const FPTYPE& w1,
FPTYPE** rho,
const std::complex<FPTYPE>* wfcr,
Expand Down
26 changes: 15 additions & 11 deletions source/source_estate/kernels/test/elecstate_op_test.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,10 +63,11 @@ TEST_F(TestModuleElecstateMultiDevice, elecstate_pw_op_cpu)
double ** rho = new double* [1];
rho[0] = rho_data.data();
elecstate_cpu_op()(
this->cpu_ctx,
this->cpu_ctx,
this->spin, this->nrxx,
this->w1,
rho,
this->nrxx,
this->w1,
rho,
this->wfcr.data());

// check the result
Expand All @@ -85,12 +86,13 @@ TEST_F(TestModuleElecstateMultiDevice, elecstate_pw_spin_op_cpu)
rho[2] = rho_data.data() + this->nrxx * 2;
rho[3] = rho_data.data() + this->nrxx * 3;
elecstate_cpu_op()(
this->cpu_ctx,
this->cpu_ctx,
this->DOMAG,
this->DOMAG_Z,
this->nrxx,
this->w1,
rho,
this->nrxx,
this->w1,
rho,
this->wfcr_2.data(),
this->wfcr_another_spin_2.data());

Expand All @@ -114,10 +116,11 @@ TEST_F(TestModuleElecstateMultiDevice, elecstate_pw_op_gpu)
double ** rho = new double* [1];
rho[0] = d_rho_data;
elecstate_gpu_op()(
this->gpu_ctx,
this->gpu_ctx,
this->spin, this->nrxx,
this->nrxx,
this->w1,
rho,
rho,
d_wfcr);

syncmem_var_d2h_op()(rho_data.data(), d_rho_data, rho_data.size());
Expand Down Expand Up @@ -149,12 +152,13 @@ TEST_F(TestModuleElecstateMultiDevice, elecstate_pw_spin_op_gpu)
rho[3] = d_rho_data_2 + this->nrxx * 3;

elecstate_gpu_op()(
this->gpu_ctx,
this->gpu_ctx,
this->DOMAG,
this->DOMAG_Z,
this->nrxx,
this->w1,
rho,
this->nrxx,
this->w1,
rho,
d_wfcr_2,
d_wfcr_another_spin_2);

Expand Down
3 changes: 2 additions & 1 deletion source/source_io/module_wf/read_wf2rho_pw.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,7 @@ void ModuleIO::read_wf2rho_pw(
PARAM.globalv.domag,
PARAM.globalv.domag_z,
nrxx,
nrxx,
w1,
chg.rho,
rho_tmp.data(),
Expand All @@ -158,7 +159,7 @@ void ModuleIO::read_wf2rho_pw(
if (w1 != 0.0)
{
base_device::DEVICE_CPU* ctx = nullptr;
elecstate::elecstate_pw_op<double, base_device::DEVICE_CPU>()(ctx, is, nrxx,
elecstate::elecstate_pw_op<double, base_device::DEVICE_CPU>()(ctx, is, nrxx, nrxx,
w1, chg.rho, rho_tmp.data());
}
}
Expand Down
2 changes: 1 addition & 1 deletion source/source_pw/module_stodft/sto_iter.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -645,7 +645,7 @@ void Stochastic_Iter<T, Device>::cal_storho(const UnitCell& ucell,
{
wfc_basis->recip_to_real(this->ctx, tmpout, porter, ik);
const auto w1 = static_cast<Real>(this->pkv->wk[ik]);
elecstate::elecstate_pw_op<Real, Device>()(this->ctx, current_spin, nrxx, w1, pes->rho, porter);
elecstate::elecstate_pw_op<Real, Device>()(this->ctx, current_spin, nrxx, pes->charge->nrxx, w1, pes->rho, porter);
// for (int ir = 0; ir < nrxx; ++ir)
// {
// pes->charge->rho[0][ir] += norm(porter[ir]) * this->pkv->wk[ik];
Expand Down
Loading