Skip to content

Commit

Permalink
Merge pull request #486 from toxa81/develop
Browse files Browse the repository at this point in the history
bugfix in computing derivative of augmentation operator on GPU
  • Loading branch information
toxa81 authored Mar 25, 2020
2 parents 6df5b0b + d0efeb2 commit 2b2720b
Show file tree
Hide file tree
Showing 4 changed files with 12 additions and 31 deletions.
34 changes: 7 additions & 27 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -88,8 +88,6 @@ RUN wget -O - https://apt.kitware.com/keys/kitware-archive-latest.asc 2>/dev/nul
RUN apt-add-repository 'deb https://apt.kitware.com/ubuntu/ bionic main'
RUN apt-get update
RUN apt-get install -y cmake
RUN apt-get install -y kitware-archive-keyring
RUN apt-key --keyring /etc/apt/trusted.gpg del C1F34CDD40CD72DA

WORKDIR /root
ENTRYPOINT ["bash", "-l"]
Expand Down Expand Up @@ -335,29 +333,11 @@ Have you got any questions, feel free to contact us:

## Acknowledgements
The development of SIRIUS library would not be possible without support of the following organizations:

|List of supporting organizations |
|:------------------------------------------------:|
| |
| Swiss Federal Institute of Technology in Zürich | |
| ![ethz](doc/images/logo_ethz.png) |
| https://www.ethz.ch/ |
| |
| Swiss National Supercomputing Centre |
| ![cscs](doc/images/logo_cscs.png) |
| https://www.cscs.ch/ |
| |
| Platform for Advanced Scientific Computing |
| ![pasc](doc/images/logo_pasc.png) |
| https://www.pasc-ch.org/ |
| |
| MAX (MAterials design at the eXascale) <br> European Centre of Excellence |
| ![pasc](doc/images/logo_max.png) |
| http://www.max-centre.eu/ |
| |
| Partnership for Advanced Computing in Europe |
| ![pasc](doc/images/logo_prace.png) |
| https://prace-ri.eu/ |


| Logo | Name | URL |
|:----:|:----:|:---:|
|![ethz](doc/images/logo_ethz.png) | Swiss Federal Institute of Technology in Zürich | https://www.ethz.ch/ |
|![cscs](doc/images/logo_cscs.png) | Swiss National Supercomputing Centre | https://www.cscs.ch/ |
|![pasc](doc/images/logo_pasc.png) | Platform for Advanced Scientific Computing | https://www.pasc-ch.org/ |
|![pasc](doc/images/logo_max.png) | MAX (MAterials design at the eXascale) <br> European Centre of Excellence | http://www.max-centre.eu/ |
|![pasc](doc/images/logo_prace.png) | Partnership for Advanced Computing in Europe | https://prace-ri.eu/ |

5 changes: 2 additions & 3 deletions src/density/augmentation_operator.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -277,7 +277,7 @@ void Augmentation_operator_gvec_deriv::prepare(Atom_type const& atom_type__,
{
PROFILE("sirius::Augmentation_operator_gvec_deriv::prepare");

int lmax_beta = atom_type__.indexr().lmax();
int lmax_beta = atom_type__.lmax_beta();

/* number of beta- radial functions */
int nbrf = atom_type__.mt_radial_basis_size();
Expand Down Expand Up @@ -394,7 +394,6 @@ void Augmentation_operator_gvec_deriv::generate_pw_coeffs(Atom_type const& atom_

switch (atom_type__.parameters().processing_unit()) {
case device_t::CPU: {
auto gc = gaunt_coefs_->get_full_set_L3();
#pragma omp parallel for schedule(static)
for (int igloc = 0; igloc < gvec_count; igloc++) {
/* index of the G-vector shell */
Expand Down Expand Up @@ -428,7 +427,7 @@ void Augmentation_operator_gvec_deriv::generate_pw_coeffs(Atom_type const& atom_
aug_op_pw_coeffs_deriv_gpu(gvec_count, gvec_shell_.at(memory_t::device), gvec_cart_.at(memory_t::device),
idx_.at(memory_t::device), static_cast<int>(idx_.size(1)),
gc.at(memory_t::device), static_cast<int>(gc.size(0)), static_cast<int>(gc.size(1)),
rlm_g_.at(memory_t::device), rlm_dg_.at(memory_t::device), lmmax,
rlm_g_.at(memory_t::device), rlm_dg_.at(memory_t::device), static_cast<int>(rlm_g_.size(0)),
ri_values_.at(memory_t::device), ri_dg_values_.at(memory_t::device), static_cast<int>(ri_values_.size(0)),
static_cast<int>(ri_values_.size(1)), q_pw_.at(memory_t::device), static_cast<int>(q_pw_.size(0)),
fourpi, nu__, lmax_q);
Expand Down
2 changes: 1 addition & 1 deletion src/gpu/augmentation_operator.cu
Original file line number Diff line number Diff line change
Expand Up @@ -105,7 +105,7 @@ __global__ void aug_op_pw_coeffs_deriv_gpu_kernel(int ngvec__, int const* gvec_s
z.x += d * phase.x;
z.y -= d * phase.y;
/* i^l */
phase = accCmul(phase, make_accDoubleComplex(0, -1));
phase = accCmul(phase, make_accDoubleComplex(0, 1));
}
q_pw__[array2D_offset(idx12, 2 * igloc, ld5__)] = z.x * fourpi__;
q_pw__[array2D_offset(idx12, 2 * igloc + 1, ld5__)] = z.y * fourpi__;
Expand Down
2 changes: 2 additions & 0 deletions src/unit_cell/atom_type.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -1485,7 +1485,9 @@ inline void Atom_type::print_info() const
std::printf("number of aw basis functions : %i\n", indexb().size_aw());
std::printf("number of lo basis functions : %i\n", indexb().size_lo());
if (!parameters_.full_potential()) {
std::printf("lmax of beta-projectors : %i\n", this->lmax_beta());
std::printf("number of ps wavefunctions : %i\n", this->indexr_wfs().size());
std::printf("charge augmentation : %s\n", utils::boolstr(this->augment()).c_str());
}
std::printf("Hubbard correction : %s\n", utils::boolstr(this->hubbard_correction()).c_str());
if (parameters_.hubbard_correction() && this->hubbard_correction_) {
Expand Down

0 comments on commit 2b2720b

Please sign in to comment.