86 complex(real64),
allocatable :: phase(:, :)
89 complex(real64),
public,
allocatable :: phase_corr(:,:)
92 complex(real64),
allocatable :: phase_spiral(:,:)
95 type(accel_mem_t) :: buff_phase
96 type(accel_mem_t) :: buff_phase_spiral
97 type(accel_mem_t),
public :: buff_phase_corr
98 integer :: buff_phase_qn_start
99 real(real64),
public,
pointer :: spin(:,:,:) => null()
125 class(phase_t),
intent(inout) :: phase
126 type(grid_t),
intent(in) :: gr
127 type(distributed_t),
intent(in) :: kpt
128 type(kpoints_t),
intent(in) :: kpoints
129 type(states_elec_dim_t),
intent(in) :: d
130 type(space_t),
intent(in) :: space
132 integer :: ip, ik, sp
133 integer(int64) :: ip_inner_global
134 real(real64) :: kpoint(space%dim), x_global(space%dim)
141 phase%buff_phase_qn_start = kpt%start
143 if(kpoints%gamma_only())
then
148 safe_allocate(phase%phase(1:gr%np_part, kpt%start:kpt%end))
149 safe_allocate(phase%phase_corr(gr%np+1:gr%np_part, kpt%start:kpt%end))
151 do ik = kpt%start, kpt%end
153 do ip = gr%np + 1, gr%np_part
154 phase%phase_corr(ip, ik) =
m_one
160 if (gr%der%boundaries%spiralBC)
then
162 if (gr%parallel_in_domains) sp = gr%np + gr%pv%np_ghost
168 safe_allocate(phase%phase_spiral(1:gr%np_part-sp, 1:2))
171 do ip = sp + 1, gr%np_part
175 phase%phase_spiral(ip-sp, 1) = &
176 exp(
m_zi * sum((gr%x(ip, 1:space%dim)-x_global(1:space%dim)) * gr%der%boundaries%spiral_q(1:space%dim)))
177 phase%phase_spiral(ip-sp, 2) = &
178 exp(-
m_zi * sum((gr%x(ip, 1:space%dim)-x_global(1:space%dim)) * gr%der%boundaries%spiral_q(1:space%dim)))
183 call accel_write_buffer(phase%buff_phase_spiral, (gr%np_part-sp)*2, phase%phase_spiral)
191 if (gr%parallel_in_domains) sp = gr%np + gr%pv%np_ghost
193 do ik = kpt%start, kpt%end
194 kpoint(1:space%dim) = kpoints%get_point(d%get_kpoint_index(ik))
197 do ip = 1, gr%np_part
198 phase%phase(ip, ik) =
exp(-
m_zi * sum(gr%x(ip, 1:space%dim) * kpoint(1:space%dim)))
204 do ip = sp + 1, gr%np_part
210 phase%phase_corr(ip, ik) = phase%phase(ip, ik)* &
211 exp(
m_zi * sum(x_global(1:space%dim) * kpoint(1:space%dim)))
221 call accel_write_buffer(phase%buff_phase_corr, (gr%np_part - gr%np)*kpt%nlocal, phase%phase_corr)
229 subroutine phase_update_phases(phase, mesh, kpt, kpoints, d, space, uniform_vector_potential)
230 class(
phase_t),
intent(inout) :: phase
231 class(
mesh_t),
intent(in) :: mesh
235 type(
space_t),
intent(in) :: space
236 real(real64),
allocatable,
intent(in) :: uniform_vector_potential(:)
238 integer :: ik, ip, sp
239 integer(int64) :: ip_inner_global
240 real(real64) :: kpoint(space%dim)
241 real(real64),
allocatable :: x_global(:,:)
245 if (
allocated(uniform_vector_potential))
then
247 if (.not.
allocated(phase%phase))
then
248 safe_allocate(phase%phase(1:mesh%np_part, kpt%start:kpt%end))
251 mesh%np_part*kpt%nlocal)
255 if (.not.
allocated(phase%phase_corr))
then
256 safe_allocate(phase%phase_corr(mesh%np+1:mesh%np_part, kpt%start:kpt%end))
259 (mesh%np_part - mesh%np)*kpt%nlocal)
263 kpoint(1:space%dim) =
m_zero
267 if (mesh%parallel_in_domains) sp = mesh%np + mesh%pv%np_ghost
269 safe_allocate(x_global(1:space%dim,(sp + 1):mesh%np_part))
272 do ip = sp + 1, mesh%np_part
277 x_global(:,ip) =
mesh_x_global(mesh, ip_inner_global) - mesh%x(ip, 1:space%dim)
280 do ik = kpt%start, kpt%end
281 kpoint(1:space%dim) = kpoints%get_point(d%get_kpoint_index(ik))
283 kpoint(1:space%dim) = kpoint(1:space%dim) + uniform_vector_potential(1:space%dim)
286 do ip = 1, mesh%np_part
287 phase%phase(ip, ik) =
exp(cmplx(
m_zero, -sum(mesh%x(ip, 1:space%dim)*kpoint(1:space%dim)), real64))
292 do ip = sp + 1, mesh%np_part
293 phase%phase_corr(ip, ik) =
exp(cmplx(
m_zero, sum(x_global(1:space%dim, ip) * kpoint(1:space%dim)), real64))
298 safe_deallocate_a(x_global)
303 phase%phase_corr, async=.
true.)
313 class(
phase_t),
intent(inout) :: phase
326 safe_deallocate_a(phase%phase)
327 safe_deallocate_a(phase%phase_corr)
328 safe_deallocate_a(phase%phase_spiral)
337 class(
phase_t),
intent(in) :: phase
338 class(
mesh_t),
intent(in) :: mesh
340 logical,
optional,
intent(in) :: async
343 logical :: phase_correction
348 phase_correction = phase%is_allocated()
352 if (phase_correction)
then
353 call phase%apply_to(mesh, mesh%np, .false., psib, async=async)
363 class(
phase_t),
intent(in) :: phase
364 class(
mesh_t),
intent(in) :: mesh
366 logical,
optional,
intent(in) :: async
368 logical :: phase_correction
373 phase_correction = phase%is_allocated()
377 if (phase_correction)
then
378 call phase%apply_to(mesh, mesh%np, .
true., psib, async=async)
388 class(
phase_t),
intent(in) :: this
389 class(
mesh_t),
intent(in) :: mesh
390 integer,
intent(in) :: np
391 logical,
intent(in) :: conjugate
392 type(
wfs_elec_t),
target,
intent(inout) :: psib
393 type(
wfs_elec_t),
optional,
target,
intent(in) :: src
394 logical,
optional,
intent(in) :: async
396 integer :: ip, ii, sp
398 complex(real64) :: phase
399 integer(int64) :: wgsize, dim2, dim3
407 assert(np <= mesh%np_part)
411 if (
present(src)) src_ => src
413 assert(src_%has_phase .eqv. conjugate)
414 assert(src_%ik == psib%ik)
418 sp = min(np, mesh%np)
419 if (np > mesh%np .and. mesh%parallel_in_domains) sp = mesh%np + mesh%pv%np_ghost
421 select case (psib%status())
428 do ip = 1, min(mesh%np, np)
429 phase = conjg(this%phase(ip, psib%ik))
431 do ii = 1, psib%nst_linear
432 psib%zff_pack(ii, ip) = phase*src_%zff_pack(ii, ip)
440 phase = conjg(this%phase(ip, psib%ik))
442 do ii = 1, psib%nst_linear
443 psib%zff_pack(ii, ip) = phase*src_%zff_pack(ii, ip)
452 do ip = 1, min(mesh%np, np)
453 phase = this%phase(ip, psib%ik)
455 do ii = 1, psib%nst_linear
456 psib%zff_pack(ii, ip) = phase*src_%zff_pack(ii, ip)
464 phase = this%phase(ip, psib%ik)
466 do ii = 1, psib%nst_linear
467 psib%zff_pack(ii, ip) = phase*src_%zff_pack(ii, ip)
479 do ii = 1, psib%nst_linear
481 do ip = 1, min(mesh%np, np)
482 psib%zff_linear(ip, ii) = conjg(this%phase(ip, psib%ik))*src_%zff_linear(ip, ii)
489 psib%zff_linear(ip, ii) = conjg(this%phase(ip, psib%ik))*src_%zff_linear(ip, ii)
497 do ii = 1, psib%nst_linear
499 do ip = 1, min(mesh%np, np)
500 psib%zff_linear(ip, ii) = this%phase(ip, psib%ik)*src_%zff_linear(ip, ii)
507 psib%zff_linear(ip, ii) = this%phase(ip, psib%ik)*src_%zff_linear(ip, ii)
538 call accel_kernel_run(ker_phase, (/psib%pack_size(1), dim2, dim3/), (/psib%pack_size(1), wgsize, 1_int64/))
543 psib%has_phase = .not. conjugate
556 class(
phase_t),
intent(in) :: this
557 complex(real64),
intent(inout) :: psi(:, :)
558 integer,
intent(in) :: np
559 integer,
intent(in) :: dim
560 integer,
intent(in) :: ik
561 logical,
intent(in) :: conjugate
574 psi(ip, idim) = conjg(this%phase(ip, ik))*psi(ip, idim)
583 psi(ip, idim) = this%phase(ip, ik)*psi(ip, idim)
599 class(
phase_t),
intent(in) :: this
603 integer :: ip, ii, sp
604 integer,
allocatable :: spin_label(:)
606 integer(int64) :: wgsize
613 assert(der%boundaries%spiral)
617 if (der%mesh%parallel_in_domains) sp = der%mesh%np + der%mesh%pv%np_ghost
620 select case (psib%status())
624 do ip = sp + 1, der%mesh%np_part
625 do ii = 1, psib%nst_linear, 2
626 if (this%spin(3,psib%linear_to_ist(ii), psib%ik)>0)
then
627 psib%zff_pack(ii+1, ip) = psib%zff_pack(ii+1, ip)*this%phase_spiral(ip-sp, 1)
629 psib%zff_pack(ii, ip) = psib%zff_pack(ii, ip)*this%phase_spiral(ip-sp, 2)
638 do ii = 1, psib%nst_linear, 2
639 if (this%spin(3,psib%linear_to_ist(ii), psib%ik)>0)
then
641 do ip = sp + 1, der%mesh%np_part
642 psib%zff_linear(ip, ii+1) = psib%zff_linear(ip, ii+1)*this%phase_spiral(ip-sp, 1)
647 do ip = sp + 1, der%mesh%np_part
648 psib%zff_linear(ip, ii) = psib%zff_linear(ip, ii)*this%phase_spiral(ip-sp, 2)
664 safe_allocate(spin_label(1:psib%nst_linear))
666 do ii = 1, psib%nst_linear, 2
667 if (this%spin(3, psib%linear_to_ist(ii), psib%ik) > 0) spin_label(ii)=1
686 (/psib%pack_size(1)/2,
pad(der%mesh%np_part - sp, 2*wgsize)/), &
687 (/psib%pack_size(1)/2, 2*wgsize/))
693 safe_deallocate_a(spin_label)
703 logical pure function phase_is_allocated(this)
704 class(
phase_t),
intent(in) :: this
706 phase_is_allocated =
allocated(this%phase)
double exp(double __x) __attribute__((__nothrow__
subroutine, public accel_kernel_start_call(this, file_name, kernel_name, flags)
subroutine, public accel_finish()
integer pure function, public accel_max_size_per_dim(dim)
subroutine, public accel_release_buffer(this)
type(accel_kernel_t), target, save, public kernel_phase_spiral
pure logical function, public accel_is_enabled()
integer function, public accel_kernel_workgroup_size(kernel)
integer, parameter, public accel_mem_read_only
This module implements batches of mesh functions.
integer, parameter, public batch_not_packed
functions are stored in CPU memory, unpacked order
integer, parameter, public batch_device_packed
functions are stored in device memory in packed order
integer, parameter, public batch_packed
functions are stored in CPU memory, in transposed (packed) order
This module implements common operations on batches of mesh functions.
This module calculates the derivatives (gradients, Laplacians, etc.) of a function.
real(real64), parameter, public m_zero
complex(real64), parameter, public m_zi
real(real64), parameter, public m_one
This module implements the underlying real-space grid.
This module is intended to contain "only mathematical" functions and procedures.
This module defines the meshes, which are used in Octopus.
integer(int64) function, public mesh_periodic_point(mesh, space, ip)
This function returns the point inside the grid corresponding to a boundary point when PBCs are used....
real(real64) function, dimension(1:mesh%box%dim), public mesh_x_global(mesh, ipg)
subroutine phase_phase_spiral(this, der, psib)
apply spiral phase
subroutine phase_unset_phase_corr(phase, mesh, psib, async)
unset the phase correction (if necessary)
subroutine phase_init_phases(phase, gr, kpt, kpoints, d, space)
Initiliaze the phase arrays and copy to GPU the data.
subroutine phase_end(phase)
Releases the memory of the phase object.
subroutine phase_update_phases(phase, mesh, kpt, kpoints, d, space, uniform_vector_potential)
Update the phases.
logical pure function phase_is_allocated(this)
subroutine phase_apply_batch(this, mesh, np, conjugate, psib, src, async)
apply (remove) the phase to the wave functions before (after) applying the Hamiltonian
subroutine phase_set_phase_corr(phase, mesh, psib, async)
set the phase correction (if necessary)
subroutine phase_apply_mf(this, psi, np, dim, ik, conjugate)
apply (or remove) the phase to a wave function psi
subroutine, public profiling_out(label)
Increment out counter and sum up difference between entry and exit time.
subroutine, public profiling_in(label, exclude)
Increment in counter and save entry time.
This module handles spin dimensions of the states and the k-point distribution.
type(type_t), public type_cmplx
type(type_t), public type_integer
class representing derivatives
Distribution of N instances over mpi_grpsize processes, for the local rank mpi_grprank....
Describes mesh distribution to nodes.
A container for the phase.
class for organizing spins and k-points
batches of electronic states