Compare commits

...

8 Commits

Author SHA1 Message Date
Anthony Scemama ecd9531ee6 H_oo on GPU 2023-08-03 20:02:43 +02:00
Anthony Scemama 2ea84d4ffc Separate upload of t1, t2 and tau 2023-08-03 17:33:04 +02:00
Anthony Scemama d3c87d8181 r2 fully on GPU 2023-08-03 12:00:09 +02:00
Anthony Scemama a541abd04a r2 on GPU 2023-08-03 03:21:32 +02:00
Anthony Scemama 30e2e086a0 more dgemm 2023-08-03 02:57:43 +02:00
Anthony Scemama 803a59b586 Removed J1 from args 2023-08-03 02:12:15 +02:00
Anthony Scemama eec4af61e2 K1 on GPU 2023-08-03 01:21:16 +02:00
Anthony Scemama 2763954c19 Removed compute_J1 2023-08-02 23:50:28 +02:00
6 changed files with 922 additions and 762 deletions

View File

@ -120,19 +120,20 @@ subroutine run_ccsd_space_orb
gpu_data = gpu_init(nO, nV, cholesky_mo_num, &
cc_space_v_oo_chol, cc_space_v_ov_chol, cc_space_v_vo_chol, cc_space_v_vv_chol, &
cc_space_v_oooo, cc_space_v_vooo, cc_space_v_oovv, cc_space_v_vvoo, &
cc_space_v_oovo, cc_space_v_ovvo, cc_space_v_ovoo, cc_space_f_vo)
cc_space_v_oovo, cc_space_v_ovvo, cc_space_v_ovov, cc_space_v_ovoo, &
cc_space_f_oo, cc_space_f_vo, cc_space_f_vv)
do while (not_converged)
! Residue
if (do_ao_cholesky) then
! if (.False.) then
call compute_H_oo_chol(nO,nV,tau_x,H_oo)
call compute_H_vv_chol(nO,nV,tau_x,H_vv)
call compute_H_vo_chol(nO,nV,t1,H_vo)
call gpu_upload(gpu_data, nO, nV, t1, t2, tau, tau_x, H_vv);
call compute_H_oo_chol_gpu(gpu_data,nO,nV,0,H_oo)
call compute_r1_space_chol(nO,nV,t1,t2,tau,H_oo,H_vv,H_vo,r1,max_r1)
call compute_r2_space_chol(nO,nV,t1,t2,tau,H_oo,H_vv,H_vo,r2,max_r2,gpu_data)
call compute_r2_space_chol_gpu(gpu_data, nO, nV, t1, r2, max_r2)
else
call compute_H_oo(nO,nV,t1,t2,tau,H_oo)
call compute_H_vv(nO,nV,t1,t2,tau,H_vv)

View File

@ -439,477 +439,41 @@ subroutine compute_H_vo_chol(nO,nV,t1,H_vo)
end
! R2
subroutine compute_r2_space_chol(nO,nV,t1,t2,tau,H_oo,H_vv,H_vo,r2,max_r2,gpu_data)
use gpu_module
implicit none
! in
integer, intent(in) :: nO, nV
double precision, intent(in) :: t1(nO,nV), t2(nO,nO,nV,nV), tau(nO,nO,nV,nV)
double precision, intent(in) :: H_oo(nO,nO), H_vv(nV,nV), H_vo(nV,nO)
type(c_ptr), intent(in) :: gpu_data
! out
double precision, intent(out) :: r2(nO,nO,nV,nV), max_r2
! internal
integer :: u,v,i,j,beta,gam,a,b
double precision :: max_r2_local
double precision, allocatable :: Y_oovv(:,:,:,:)
integer :: block_size, iblock, k
block_size = 16
call set_multiple_levels_omp(.False.)
double precision, allocatable :: g_occ(:,:)
allocate(g_occ(nO,nO))
!$omp parallel &
!$omp shared(nO,nV,g_occ,H_oo, cc_space_v_ovoo,t1) &
!$omp private(i,j,a,u) &
!$omp default(none)
!$omp do
do i = 1, nO
do u = 1, nO
g_occ(u,i) = H_oo(u,i)
enddo
enddo
!$omp end do
!$omp do
do i = 1, nO
do j = 1, nO
do a = 1, nV
do u = 1, nO
g_occ(u,i) = g_occ(u,i) + (2d0 * cc_space_v_ovoo(u,a,i,j) - cc_space_v_ovoo(u,a,j,i)) * t1(j,a)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
double precision, allocatable :: J1(:,:,:,:)
allocate(J1(nO,nV,nV,nO))
J1 = 0.d0
call compute_r2_space_chol_gpu(nO,nV,cholesky_mo_num,gpu_data,t1,t2,tau, &
H_vv, g_occ, J1, r2)
!---
double precision, allocatable :: X_oovv(:,:,:,:)
double precision, allocatable :: X_vovv(:,:,:,:)
double precision, allocatable :: X_ovvo(:,:,:,:)
double precision, allocatable :: tcc(:,:,:), tcc2(:,:,:)
double precision, allocatable :: X_vovo(:,:,:,:), Y_oovo(:,:,:,:)
call compute_J1_chol(nO,nV,t1,t2,cc_space_v_ovvo,cc_space_v_ovoo, &
cc_space_v_vvoo,J1)
double precision, allocatable :: K1(:,:,:,:)
allocate(K1(nO,nV,nO,nV))
call compute_K1_chol(nO,nV,t1,t2,cc_space_v_ovoo,cc_space_v_vvoo, &
cc_space_v_ovov,K1)
allocate(X_ovvo(nO,nV,nV,nO))
!$omp parallel &
!$omp private(u,v,gam,beta,i,a) &
!$omp default(shared)
do i = 1, nO
!$omp do
do a = 1, nV
do beta = 1, nV
do u = 1, nO
X_ovvo(u,beta,a,i) = (J1(u,a,beta,i) - 0.5d0 * K1(u,a,i,beta))
enddo
enddo
enddo
!$omp end do nowait
enddo
!$omp end parallel
deallocate(J1)
double precision, allocatable :: Y_voov(:,:,:,:)
allocate(Y_voov(nV,nO,nO,nV))
!$omp parallel &
!$omp private(u,v,gam,beta,i,a) &
!$omp default(shared)
!$omp do
do gam = 1, nV
do v = 1, nO
do i = 1, nO
do a = 1, nV
Y_voov(a,i,v,gam) = 2d0 * t2(i,v,a,gam) - t2(i,v,gam,a)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
double precision, allocatable :: Z_ovov(:,:,:,:)
allocate(Z_ovov(nO,nV,nO,nV))
call dgemm('N','N', nO*nV,nO*nV,nV*nO, &
1d0, X_ovvo, size(X_ovvo,1) * size(X_ovvo,2), &
Y_voov, size(Y_voov,1) * size(Y_voov,2), &
0d0, Z_ovov, size(Z_ovov,1) * size(Z_ovov,2))
deallocate(X_ovvo,Y_voov)
!$omp parallel &
!$omp shared(nO,nV,r2,Z_ovov) &
!$omp private(u,v,gam,beta) &
!$omp default(none)
!$omp do
do gam = 1, nV
do beta = 1, nV
do v = 1, nO
do u = 1, nO
r2(u,v,beta,gam) = r2(u,v,beta,gam) + Z_ovov(u,beta,v,gam) + Z_ovov(v,gam,u,beta)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
deallocate(Z_ovov)
double precision, allocatable :: Y_ovov(:,:,:,:), X_ovov(:,:,:,:)
allocate(X_ovov(nO,nV,nO,nV))
allocate(Y_ovov(nO,nV,nO,nV))
!$omp parallel &
!$omp shared(nO,nV,r2,K1,X_ovov,Y_ovov,t2) &
!$omp private(u,a,i,beta,gam) &
!$omp default(none)
!$omp do
do beta = 1, nV
do u = 1, nO
do a = 1, nV
do i = 1, nO
X_ovov(i,a,u,beta) = 0.5d0 * K1(u,a,i,beta)
enddo
enddo
enddo
enddo
!$omp end do nowait
!$omp do
do gam = 1, nV
do v = 1, nO
do a = 1, nV
do i = 1, nO
Y_ovov(i,a,v,gam) = t2(i,v,gam,a)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
allocate(Z_ovov(nO,nV,nO,nV))
call dgemm('T','N',nO*nV,nO*nV,nO*nV, &
1d0, X_ovov, size(X_ovov,1) * size(X_ovov,2), &
Y_ovov, size(Y_ovov,1) * size(Y_ovov,2), &
0d0, Z_ovov, size(Y_ovov,1) * size(Y_ovov,2))
deallocate(X_ovov, Y_ovov)
!$omp parallel &
!$omp shared(nO,nV,r2,Z_ovov) &
!$omp private(u,v,gam,beta) &
!$omp default(none)
!$omp do
do gam = 1, nV
do beta = 1, nV
do v = 1, nO
do u = 1, nO
r2(u,v,beta,gam) = r2(u,v,beta,gam) - Z_ovov(u,beta,v,gam) - Z_ovov(v,gam,u,beta)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
deallocate(Z_ovov)
allocate(X_ovov(nO,nV,nO,nV),Y_ovov(nO,nV,nO,nV))
!$omp parallel &
!$omp shared(nO,nV,K1,X_ovov,Y_ovov,t2) &
!$omp private(u,v,gam,beta,i,a) &
!$omp default(none)
!$omp do
do a = 1, nV
do i = 1, nO
do gam = 1, nV
do u = 1, nO
X_ovov(u,gam,i,a) = K1(u,a,i,gam)
enddo
enddo
enddo
enddo
!$omp end do nowait
!$omp do
do beta = 1, nV
do v = 1, nO
do a = 1, nV
do i = 1, nO
Y_ovov(i,a,v,beta) = t2(i,v,beta,a)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
deallocate(K1)
allocate(Z_ovov(nO,nV,nO,nV))
call dgemm('N','N',nO*nV,nO*nV,nO*nV, &
1d0, X_ovov, size(X_ovov,1) * size(X_ovov,2), &
Y_ovov, size(Y_ovov,1) * size(Y_ovov,2), &
0d0, Z_ovov, size(Y_ovov,1) * size(Y_ovov,2))
deallocate(X_ovov,Y_ovov)
!$omp parallel &
!$omp shared(nO,nV,r2,Z_ovov) &
!$omp private(u,v,gam,beta) &
!$omp default(none)
!$omp do
do gam = 1, nV
do beta = 1, nV
do v = 1, nO
do u = 1, nO
r2(u,v,beta,gam) = r2(u,v,beta,gam) - Z_ovov(u,gam,v,beta) - Z_ovov(v,beta,u,gam)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
deallocate(Z_ovov)
! Change the sign for consistency with the code in spin orbitals
max_r2 = 0d0
!$omp parallel &
!$omp shared(nO,nV,r2,max_r2) &
!$omp private(i,j,a,b,max_r2_local) &
!$omp default(none)
max_r2_local = 0.d0
!$omp do
do b = 1, nV
do a = 1, nV
do j = 1, nO
do i = 1, nO
r2(i,j,a,b) = -r2(i,j,a,b)
max_r2_local = max(r2(i,j,a,b), max_r2_local)
enddo
enddo
enddo
enddo
!$omp end do nowait
!$omp critical
max_r2 = max(max_r2, max_r2_local)
!$omp end critical
!$omp end parallel
end
subroutine compute_J1_chol(nO,nV,t1,t2,v_ovvo,v_ovoo,v_vvoo,J1)
implicit none
integer, intent(in) :: nO,nV
double precision, intent(in) :: t1(nO, nV)
double precision, intent(in) :: t2(nO, nO, nV, nV)
double precision, intent(in) :: v_ovvo(nO,nV,nV,nO), v_ovoo(nO,nV,nO,nO)
double precision, intent(in) :: v_vvoo(nV,nV,nO,nO)
double precision, intent(inout) :: J1(nO, nV, nV, nO)
integer :: a,tmp_a,b,k,l,c,d,tmp_c,tmp_d,i,j,u,v, beta, gam
double precision, allocatable :: X_ovoo(:,:,:,:), Y_ovov(:,:,:,:)
allocate(X_ovoo(nO,nV,nO,nO),Y_ovov(nO,nV,nO,nV))
deallocate(X_ovoo)
double precision, allocatable :: tmp_cc(:,:,:), J1_tmp(:,:,:,:)
!TODO: I am here
!- cc_space_v_vvoo(a,b,i,j) * (0.5d0 * t2(u,j,b,beta) + t1(u,b) * t1(j,beta)) &
double precision, allocatable :: X_voov(:,:,:,:), Z_ovvo(:,:,:,:)
allocate(X_voov(nV,nO,nO,nV), Z_ovvo(nO,nV,nV,nO))
double precision, allocatable :: X_ovvo(:,:,:,:), Y_vovo(:,:,:,:)
allocate(X_ovvo(nO,nV,nV,nO))
allocate(Y_vovo(nV,nO,nV,nO))
!$omp parallel &
!$omp shared(nO,nV,J1,Z_ovvo,t2,Y_vovo,v_vvoo,X_ovvo) &
!$omp private(i,beta,a,u,j,b) &
!$omp default(none)
!+ 0.5d0 * (2d0 * cc_space_v_vvoo(a,b,i,j) - cc_space_v_vvoo(b,a,i,j)) * t2(u,j,beta,b)
do j = 1, nO
!$omp do
do b = 1, nV
do i = 1, nO
do a = 1, nV
Y_vovo(a,i,b,j) = 0.5d0 * (2d0 * v_vvoo(a,b,i,j) - v_vvoo(b,a,i,j))
enddo
enddo
enddo
!$omp end do nowait
enddo
do j = 1, nO
!$omp do
do b = 1, nV
do beta = 1, nV
do u = 1, nO
X_ovvo(u,beta,b,j) = t2(u,j,beta,b)
enddo
enddo
enddo
!$omp end do nowait
enddo
!$omp end parallel
call dgemm('N','T',nO*nV,nV*nO,nV*nO, &
1d0, X_ovvo, size(X_ovvo,1) * size(X_ovvo,2), &
Y_vovo, size(Y_vovo,1) * size(Y_vovo,2), &
0d0, Z_ovvo, size(Z_ovvo,1) * size(Z_ovvo,2))
!$omp parallel &
!$omp shared(nO,nV,J1,Z_ovvo) &
!$omp private(i,beta,a,u) &
!$omp default(none)
do i = 1, nO
!$omp do
do beta = 1, nV
do a = 1, nV
do u = 1, nO
J1(u,a,beta,i) = J1(u,a,beta,i) + Z_ovvo(u,beta,a,i)
enddo
enddo
enddo
!$omp end do nowait
enddo
!$omp end parallel
deallocate(X_ovvo,Z_ovvo)
end
! K1
subroutine compute_K1_chol(nO,nV,t1,t2,v_ovoo,v_vvoo,v_ovov,K1)
subroutine compute_H_oo_chol2(nO,nV,tau_x,H_oo)
implicit none
integer, intent(in) :: nO,nV
double precision, intent(in) :: t1(nO, nV)
double precision, intent(in) :: t2(nO, nO, nV, nV)
double precision, intent(in) :: v_vvoo(nV,nV,nO,nO), v_ovov(nO,nV,nO,nV)
double precision, intent(in) :: v_ovoo(nO,nV,nO,nO)
double precision, intent(out) :: K1(nO, nV, nO, nV)
double precision, intent(in) :: tau_x(nO, nO, nV, nV)
double precision, intent(out) :: H_oo(nO, nO)
double precision, allocatable :: X(:,:,:,:), Y(:,:,:,:), Z(:,:,:,:)
integer :: a,b,i,j,u,k
integer :: a,tmp_a,b,k,l,c,d,tmp_c,tmp_d,i,j,u,v, beta, gam
double precision, allocatable :: tau_kau(:,:,:), tmp_vov(:,:,:), tmp_ovv(:,:,:)
allocate(X(nV,nO,nV,nO),Y(nO,nV,nV,nO),Z(nO,nV,nV,nO))
!$omp parallel &
!$omp shared(nO,nV,K1,X,Y,v_vvoo,v_ovov,t1,t2) &
!$omp private(i,beta,a,u,j,b) &
!$omp default(none)
!$omp do
do beta = 1, nV
do i = 1, nO
do a = 1, nV
do u = 1, nO
K1(u,a,i,beta) = v_ovov(u,a,i,beta)
allocate(tau_kau(cholesky_mo_num,nV,nO))
allocate(tmp_vov(nV,nO,nV) )
allocate(tmp_ovv(nO,nV,nV) )
do u = 1, nO
call dcopy(nO*nV*nV, tau_x(u,1,1,1), nO, tmp_ovv, 1)
print *, u
print *, tmp_ovv
do b=1,nV
do j=1,nO
do a=1,nV
tmp_vov(a,j,b) = tmp_ovv(j,a,b)
enddo
enddo
enddo
call dgemm('N','T',cholesky_mo_num,nV,nO*nV,1.d0, &
cc_space_v_ov_chol, cholesky_mo_num, tmp_vov, nV, &
0.d0, tau_kau(1,1,u), cholesky_mo_num)
enddo
!$omp end do nowait
do i = 1, nO
!$omp do
do a = 1, nV
do j = 1, nO
do b = 1, nV
X(b,j,a,i) = - v_vvoo(b,a,i,j)
enddo
enddo
enddo
!$omp end do nowait
enddo
do j = 1, nO
!$omp do
do b = 1, nV
do beta = 1, nV
do u = 1, nO
Y(u,beta,b,j) = 0.5d0 * t2(u,j,b,beta) + t1(u,b) * t1(j,beta)
enddo
enddo
enddo
!$omp end do
enddo
!$omp end parallel
call dgemm('N','N',nO*nV*nO,nV,nO, &
-1d0, v_ovoo, size(v_ovoo,1) * size(v_ovoo,2) * size(v_ovoo,3), &
t1 , size(t1,1), &
1d0, K1 , size(K1,1) * size(K1,2) * size(K1,3))
double precision, allocatable :: K1tmp(:,:,:,:), t1v(:,:,:)
allocate(K1tmp(nO,nO,nV,nV), t1v(cholesky_mo_num,nO,nO))
call dgemm('N','T', cholesky_mo_num*nO, nO, nV, 1.d0, &
cc_space_v_ov_chol, cholesky_mo_num*nO, t1, nO, 0.d0, &
t1v, cholesky_mo_num*nO)
call dgemm('T','N', nO*nO, nV*nV, cholesky_mo_num, 1.d0, &
t1v, cholesky_mo_num, cc_space_v_vv_chol, cholesky_mo_num, 0.d0, &
K1tmp, nO*nO)
deallocate(t1v)
! Y(u,beta,b,j) * X(b,j,a,i) = Z(u,beta,a,i)
call dgemm('N','N',nV*nO,nO*nV,nV*nO, &
1d0, Y, size(Y,1) * size(Y,2), &
X, size(X,1) * size(X,2), &
0d0, Z, size(Z,1) * size(Z,2))
!$omp parallel &
!$omp shared(nO,nV,K1,Z,K1tmp) &
!$omp private(i,beta,a,u) &
!$omp default(none)
!$omp do
do beta = 1, nV
do i = 1, nO
do a = 1, nV
do u = 1, nO
K1(u,a,i,beta) = K1(u,a,i,beta) + K1tmp(u,i,a,beta) + Z(u,beta,a,i)
enddo
enddo
enddo
enddo
!$omp end do
!$omp end parallel
deallocate(K1tmp,X,Y,Z)
deallocate(tmp_vov)
call dcopy(nO*nO, cc_space_f_oo, 1, H_oo, 1);
call dgemm('T', 'N', nO, nO, cholesky_mo_num*nV, 1.d0, &
tau_kau, cholesky_mo_num*nV, cc_space_v_vo_chol, cholesky_mo_num*nV, &
1.d0, H_oo, nO)
end

File diff suppressed because it is too large Load Diff

View File

@ -9,7 +9,20 @@ typedef struct {
double* cc_space_v_vvoo;
double* cc_space_v_oovo;
double* cc_space_v_ovvo;
double* cc_space_v_ovov;
double* cc_space_v_ovoo;
double* cc_space_f_oo;
double* cc_space_f_vo;
double* cc_space_f_vv;
double* tau;
double* tau_x;
double* t1;
double* t2;
double* H_oo;
double* H_vv;
int nO;
int nV;
int cholesky_mo_num;
} gpu_data;
#define MULTIGPU 1

View File

@ -6,8 +6,6 @@
#include <cuda_runtime.h>
#include "gpu.h"
#define BLOCK_SIZE 16
gpu_data* gpu_init(
int nO, int nV, int cholesky_mo_num,
double* cc_space_v_oo_chol, double* cc_space_v_ov_chol,
@ -15,7 +13,9 @@ gpu_data* gpu_init(
double* cc_space_v_oooo, double* cc_space_v_vooo,
double* cc_space_v_oovv, double* cc_space_v_vvoo,
double* cc_space_v_oovo, double* cc_space_v_ovvo,
double* cc_space_v_ovoo, double* cc_space_f_vo)
double* cc_space_v_ovov, double* cc_space_v_ovoo,
double* cc_space_f_oo, double* cc_space_f_vo,
double* cc_space_f_vv)
{
int ngpus = 1;
cudaGetDeviceCount(&ngpus);
@ -77,15 +77,52 @@ gpu_data* gpu_init(
cudaMalloc((void **)&d_cc_space_v_ovvo, nO*nV*nV*nO * sizeof(double));
cublasSetMatrix(lda, nV*nO, sizeof(double), cc_space_v_ovvo, lda, d_cc_space_v_ovvo, lda);
double* d_cc_space_v_ovov;
lda = nO*nV;
cudaMalloc((void **)&d_cc_space_v_ovov, nO*nV*nV*nO * sizeof(double));
cublasSetMatrix(lda, nV*nO, sizeof(double), cc_space_v_ovov, lda, d_cc_space_v_ovov, lda);
double* d_cc_space_v_ovoo;
lda = nO*nV;
cudaMalloc((void **)&d_cc_space_v_ovoo, nO*nV*nO*nO * sizeof(double));
cublasSetMatrix(lda, nO*nO, sizeof(double), cc_space_v_ovoo, lda, d_cc_space_v_ovoo, lda);
double* d_cc_space_f_oo;
cudaMalloc((void**)&d_cc_space_f_oo, nO*nO*sizeof(double));
cublasSetMatrix(nO, nO, sizeof(double), cc_space_f_oo, nO, d_cc_space_f_oo, nO);
double* d_cc_space_f_vo;
cudaMalloc((void**)&d_cc_space_f_vo, nV*nO*sizeof(double));
cublasSetMatrix(nV, nO, sizeof(double), cc_space_f_vo, nV, d_cc_space_f_vo, nV);
double* d_cc_space_f_vv;
cudaMalloc((void**)&d_cc_space_f_vv, nV*nV*sizeof(double));
cublasSetMatrix(nV, nV, sizeof(double), cc_space_f_vv, nV, d_cc_space_f_vv, nV);
double* d_tau;
lda = nO * nO;
cudaMalloc((void **)&d_tau, lda * nV * nV * sizeof(double));
double* d_tau_x;
lda = nO * nO;
cudaMalloc((void **)&d_tau_x, lda * nV * nV * sizeof(double));
double* d_t1;
lda = nO;
cudaMalloc((void **)&d_t1, nO * nV * sizeof(double));
double* d_t2;
lda = nO*nO;
cudaMalloc((void **)&d_t2, nO*nO*nV*nV * sizeof(double));
double* d_H_oo;
lda = nO;
cudaMalloc((void **)&d_H_oo, nO * nO * sizeof(double));
double* d_H_vv;
lda = nV;
cudaMalloc((void **)&d_H_vv, nV * nV * sizeof(double));
data[igpu].cc_space_v_oo_chol = d_cc_space_v_oo_chol;
data[igpu].cc_space_v_ov_chol = d_cc_space_v_ov_chol;
data[igpu].cc_space_v_vo_chol = d_cc_space_v_vo_chol;
@ -96,8 +133,22 @@ gpu_data* gpu_init(
data[igpu].cc_space_v_vvoo = d_cc_space_v_vvoo;
data[igpu].cc_space_v_oovo = d_cc_space_v_oovo;
data[igpu].cc_space_v_ovvo = d_cc_space_v_ovvo;
data[igpu].cc_space_v_ovov = d_cc_space_v_ovov;
data[igpu].cc_space_v_ovoo = d_cc_space_v_ovoo;
data[igpu].cc_space_f_vo = d_cc_space_f_vo;
data[igpu].cc_space_f_oo = d_cc_space_f_oo;
data[igpu].cc_space_f_vo = d_cc_space_f_vo;
data[igpu].cc_space_f_vv = d_cc_space_f_vv;
data[igpu].tau = d_tau;
data[igpu].tau_x = d_tau_x;
data[igpu].t1 = d_t1;
data[igpu].t2 = d_t2;
data[igpu].H_oo = d_H_oo;
data[igpu].H_vv = d_H_vv;
data[igpu].nO = nO;
data[igpu].nV = nV;
data[igpu].cholesky_mo_num = cholesky_mo_num;
}
return data;
}

View File

@ -7,7 +7,8 @@ module gpu_module
type(c_ptr) function gpu_init(nO, nV, cholesky_mo_num, &
cc_space_v_oo_chol, cc_space_v_ov_chol, cc_space_v_vo_chol, cc_space_v_vv_chol, &
cc_space_v_oooo, cc_space_v_vooo, cc_space_v_oovv, cc_space_v_vvoo, &
cc_space_v_oovo, cc_space_v_ovvo, cc_space_v_ovoo, cc_space_f_vo) bind(C)
cc_space_v_oovo, cc_space_v_ovvo, cc_space_v_ovov, cc_space_v_ovoo, &
cc_space_f_oo, cc_space_f_vo, cc_space_f_vv) bind(C)
import c_int, c_double, c_ptr
integer(c_int), intent(in), value :: nO, nV, cholesky_mo_num
real(c_double), intent(in) :: cc_space_v_oo_chol(cholesky_mo_num,nO,nO)
@ -20,22 +21,38 @@ module gpu_module
real(c_double), intent(in) :: cc_space_v_vvoo(nV,nV,nO,nO)
real(c_double), intent(in) :: cc_space_v_oovo(nO,nO,nV,nO)
real(c_double), intent(in) :: cc_space_v_ovvo(nO,nV,nV,nO)
real(c_double), intent(in) :: cc_space_v_ovov(nO,nV,nO,nV)
real(c_double), intent(in) :: cc_space_v_ovoo(nO,nV,nO,nO)
real(c_double), intent(in) :: cc_space_f_oo(nO,nO)
real(c_double), intent(in) :: cc_space_f_vo(nV,nO)
real(c_double), intent(in) :: cc_space_f_vv(nV,nV)
end function
subroutine compute_r2_space_chol_gpu(nO,nV,cholesky_mo_num, gpu_data, t1, t2, tau,&
H_vv, g_occ, J1, r2) bind(C)
subroutine gpu_upload(gpu_data, nO, nV, t1, t2, tau, tau_x, H_vv) bind(C)
import c_int, c_double, c_ptr
integer(c_int), intent(in), value :: nO, nV, cholesky_mo_num
type(c_ptr), value :: gpu_data
integer(c_int), intent(in), value :: nO, nV
real(c_double), intent(in) :: t1(nO,nV)
real(c_double), intent(in) :: t2(nO,nO,nV,nV)
real(c_double), intent(in) :: tau(nO,nO,nV,nV)
real(c_double), intent(in) :: tau_x(nO,nO,nV,nV)
real(c_double), intent(in) :: H_vv(nV,nV)
end subroutine
subroutine compute_H_oo_chol_gpu(gpu_data, nO, nV, igpu, H_oo) bind(C)
import c_int, c_double, c_ptr
type(c_ptr), value :: gpu_data
integer(c_int), intent(in), value :: nO, nV, igpu
real(c_double), intent(out) :: H_oo(nO,nO)
end subroutine
subroutine compute_r2_space_chol_gpu(gpu_data, nO, nV, t1, r2, max_r2) bind(C)
import c_int, c_double, c_ptr
type(c_ptr), value :: gpu_data
integer(c_int), intent(in), value :: nO, nV
real(c_double), intent(in) :: t1(nO,nV)
real(c_double), intent(in) :: t2(nO,nO,nV,nV)
real(c_double), intent(in) :: tau(nO,nO,nV,nV)
real(c_double), intent(in) :: H_vv(nV,nV)
real(c_double), intent(in) :: g_occ(nO,nO)
real(c_double) :: J1(nO,nV,nV,nO)
real(c_double), intent(out) :: r2(nO,nO,nV,nV)
real(c_double), intent(out) :: max_r2
end subroutine