mirror of
https://gitlab.com/scemama/qp_plugins_scemama.git
synced 2024-12-22 04:13:40 +01:00
more dgemm
This commit is contained in:
parent
803a59b586
commit
30e2e086a0
@ -507,62 +507,6 @@ subroutine compute_r2_space_chol(nO,nV,t1,t2,tau,H_oo,H_vv,H_vo,r2,max_r2,gpu_da
|
||||
double precision, allocatable :: Y_voov(:,:,:,:)
|
||||
double precision, allocatable :: 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 &
|
||||
|
||||
@ -62,6 +62,9 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
|
||||
double* d_t2;
|
||||
double* d_tmp_cc;
|
||||
|
||||
double* d_K1;
|
||||
cudaMalloc((void **)&d_K1, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
lda = nO * nO;
|
||||
cudaMalloc((void **)&d_tau, lda * nV * nV * sizeof(double));
|
||||
cublasSetMatrix(nO*nO, nV*nV, sizeof(double), tau, lda, d_tau, lda);
|
||||
@ -82,6 +85,377 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
|
||||
#pragma omp sections
|
||||
{
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_J1;
|
||||
cudaMalloc((void **)&d_J1, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
A = d_cc_space_v_ovvo; lda = nO*nV;
|
||||
B = d_cc_space_v_ovvo; ldb = nO*nV;
|
||||
C = d_J1; ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO*nV, nV*nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
|
||||
|
||||
double* d_X_ovoo;
|
||||
cudaMalloc((void **)&d_X_ovoo, nO*nV*nO*nO * sizeof(double));
|
||||
alpha = 0.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_X_ovoo[nO*nV*(i+nO*j)]); lda = nO;
|
||||
B = &(d_cc_space_v_ovoo[nO*nV*(j+nO*i)]); ldb = nO;
|
||||
C = &(d_X_ovoo[nO*nV*(i+nO*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
|
||||
double* d_Y_ovov;
|
||||
cudaMalloc((void **)&d_Y_ovov, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV*nO; n=nV; k=nO;
|
||||
A=d_X_ovoo; lda=nO*nV*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_Y_ovov; ldc=nO*nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_ovoo);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = -1.0;
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_J1[nO*nV*(i+nV*j)]); lda = nO;
|
||||
B = &(d_Y_ovov[nO*nV*(j+nO*i)]); ldb = nO;
|
||||
C = &(d_J1[nO*nV*(i+nV*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_tmp_cc;
|
||||
cudaMalloc((void **)&d_tmp_cc, cholesky_mo_num*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=cholesky_mo_num*nV; n=nO; k=nV;
|
||||
A=d_cc_space_v_vv_chol; lda=cholesky_mo_num*nV;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_tmp_cc; ldc=cholesky_mo_num*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_J1_tmp;
|
||||
cudaMalloc((void **)&d_J1_tmp, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nV*nO; n=nV*nO; k=cholesky_mo_num;
|
||||
A=d_tmp_cc; lda=cholesky_mo_num;
|
||||
B=d_cc_space_v_vo_chol; ldb=cholesky_mo_num;
|
||||
C=d_J1_tmp; ldc=nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
cudaFree(d_tmp_cc);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_J1[nO*nV*nV*i]); lda = nO*nV;
|
||||
B = &(d_J1_tmp[nV*nO*nV*i]); ldb = nV;
|
||||
C = &(d_J1[nO*nV*nV*i]); ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_T, nO*nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_J1_tmp);
|
||||
|
||||
double* d_X_voov;
|
||||
cudaMalloc((void **)&d_X_voov, nV*nO*nO*nV * sizeof(double));
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
beta = t1[j+b*nO];
|
||||
A = &(d_t2[nO*(j + nO*nV*b)]); lda = nO*nO;
|
||||
B = d_t1; ldb = nO;
|
||||
C = &(d_Y_ovov[nO*(b+nV*j)]); ldc = nO*nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_cc_space_v_vvoo[nV*(b+nV*nO*j)]); lda = nV*nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*(b+nV*nO*j)]); ldb = nV*nV;
|
||||
C = &(d_X_voov[nV*nO*(j+nO*b)]); ldc = nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nV, nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_Z_ovvo;
|
||||
cudaMalloc((void **)&d_Z_ovvo, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = -1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV; n=nV*nO; k=nO*nV;
|
||||
A=d_Y_ovov; lda=nO*nV;
|
||||
B=d_X_voov; ldb=nV*nO;
|
||||
C=d_Z_ovvo; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_voov);
|
||||
cudaFree(d_Y_ovov);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_J1[nO*nV*(b+nV*i)]); lda = nO;
|
||||
B = &(d_Z_ovvo[nO*(b+nV*nV*i)]); ldb=nO*nV;
|
||||
C = &(d_J1[nO*nV*(b+nV*i)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
double* d_Y_vovo;
|
||||
cudaMalloc((void **)&d_Y_vovo, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = -0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); lda = nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); ldb = nV;
|
||||
C = &(d_Y_vovo[nV*(i+nO*nV*j)]); ldc = nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_T, nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
double* d_X_ovvo;
|
||||
cudaMalloc((void **)&d_X_ovvo, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_t2[nO*(j+nO*nV*b)]); lda = nO*nO;
|
||||
B = &(d_t2[nO*(j+nO*nV*b)]); ldb = nO*nO;
|
||||
C = &(d_X_ovvo[nO*nV*(b+nV*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV; n=nV*nO; k=nV*nO;
|
||||
A=d_X_ovvo; lda=nO*nV;
|
||||
B=d_Y_vovo; ldb=nV*nO;
|
||||
C=d_Z_ovvo; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_ovvo);
|
||||
cudaFree(d_Y_vovo);
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_J1[nO*nV*(b+nV*i)]); lda = nO;
|
||||
B = &(d_Z_ovvo[nO*(b+nV*nV*i)]); ldb = nO*nV;
|
||||
C = &(d_J1[nO*nV*(b+nV*i)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
cudaFree(d_Z_ovvo);
|
||||
|
||||
lda = nO*nV;
|
||||
cublasGetMatrix(nO*nV, nV*nO, sizeof(double), d_J1, lda, J1, lda);
|
||||
cudaFree(d_J1);
|
||||
|
||||
}
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
A = d_cc_space_v_ovov; lda = nO*nV;
|
||||
B = d_cc_space_v_ovov; ldb = nO*nV;
|
||||
C = d_K1; ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO*nV, nO*nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
|
||||
alpha = -1.0;
|
||||
beta = 1.0;
|
||||
m=nO*nV*nO; n=nV; k=nO;
|
||||
A=d_cc_space_v_ovoo; lda=nO*nV*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_K1; ldc=nO*nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_X;
|
||||
cudaMalloc((void **)&d_X, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
double* d_Y;
|
||||
cudaMalloc((void **)&d_Y, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha =-1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); lda = nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); ldb = nV;
|
||||
C = &(d_X[nV*(j+nO*nV*i)]); ldc = nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
alpha = 0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
beta = t1[j+b*nO];
|
||||
A = &(d_t2[nO*(j+nO*nV*b)]); lda = nO*nO;
|
||||
B = d_t1; ldb = nO;
|
||||
C = &(d_Y[nO*(b+nV*nV*j)]); ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_Z;
|
||||
cudaMalloc((void **)&d_Z, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nV*nO; n=nO*nV; k=nV*nO;
|
||||
A=d_Y; lda=nO*nV;
|
||||
B=d_X; ldb=nV*nO;
|
||||
C=d_Z; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
cudaFree(d_X);
|
||||
cudaFree(d_Y);
|
||||
|
||||
|
||||
double* d_t1v;
|
||||
cudaMalloc((void **)&d_t1v, cholesky_mo_num*nO*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=cholesky_mo_num*nO; n=nO; k=nV;
|
||||
A=d_cc_space_v_ov_chol; lda=cholesky_mo_num*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_t1v; ldc=cholesky_mo_num*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_K1tmp;
|
||||
cudaMalloc((void **)&d_K1tmp, nO*nO*nV*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nO; n=nV*nV; k=cholesky_mo_num;
|
||||
A=d_t1v; lda=cholesky_mo_num;
|
||||
B=d_cc_space_v_vv_chol; ldb=cholesky_mo_num;
|
||||
C=d_K1tmp; ldc=nO*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_t1v);
|
||||
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_K1[nO*nV*(i+nO*b)]); lda = nO;
|
||||
B = &(d_K1tmp[nO*(i+nO*nV*b)]); ldb = nO*nO;
|
||||
C = &(d_K1[nO*nV*(i+nO*b)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_K1[nO*nV*(i+nO*b)]); lda = nO;
|
||||
B = &(d_Z[nO*(b+nV*nV*i)]); ldb = nO*nV;
|
||||
C = &(d_K1[nO*nV*(i+nO*b)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_K1tmp);
|
||||
cudaFree(d_Z);
|
||||
|
||||
lda = nO*nV;
|
||||
cublasGetMatrix(nO*nV, nO*nV, sizeof(double), d_K1, lda, K1, lda);
|
||||
|
||||
}
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_Y_oooo;
|
||||
@ -255,6 +629,13 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_Y_oovv);
|
||||
}
|
||||
} // end sections
|
||||
|
||||
lda = nO*nV;
|
||||
cublasSetMatrix(lda, nO*nV, sizeof(double), K1, lda, d_K1, lda);
|
||||
|
||||
#pragma omp sections
|
||||
{
|
||||
|
||||
// g_occ
|
||||
#pragma omp section
|
||||
@ -557,391 +938,6 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_J1;
|
||||
cudaMalloc((void **)&d_J1, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
A = d_cc_space_v_ovvo; lda = nO*nV;
|
||||
B = d_cc_space_v_ovvo; ldb = nO*nV;
|
||||
C = d_J1; ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO*nV, nV*nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
|
||||
|
||||
double* d_X_ovoo;
|
||||
cudaMalloc((void **)&d_X_ovoo, nO*nV*nO*nO * sizeof(double));
|
||||
alpha = 0.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_X_ovoo[nO*nV*(i+nO*j)]); lda = nO;
|
||||
B = &(d_cc_space_v_ovoo[nO*nV*(j+nO*i)]); ldb = nO;
|
||||
C = &(d_X_ovoo[nO*nV*(i+nO*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
|
||||
double* d_Y_ovov;
|
||||
cudaMalloc((void **)&d_Y_ovov, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV*nO; n=nV; k=nO;
|
||||
A=d_X_ovoo; lda=nO*nV*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_Y_ovov; ldc=nO*nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_ovoo);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = -1.0;
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_J1[nO*nV*(i+nV*j)]); lda = nO;
|
||||
B = &(d_Y_ovov[nO*nV*(j+nO*i)]); ldb = nO;
|
||||
C = &(d_J1[nO*nV*(i+nV*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_tmp_cc;
|
||||
cudaMalloc((void **)&d_tmp_cc, cholesky_mo_num*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=cholesky_mo_num*nV; n=nO; k=nV;
|
||||
A=d_cc_space_v_vv_chol; lda=cholesky_mo_num*nV;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_tmp_cc; ldc=cholesky_mo_num*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_J1_tmp;
|
||||
cudaMalloc((void **)&d_J1_tmp, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nV*nO; n=nV*nO; k=cholesky_mo_num;
|
||||
A=d_tmp_cc; lda=cholesky_mo_num;
|
||||
B=d_cc_space_v_vo_chol; ldb=cholesky_mo_num;
|
||||
C=d_J1_tmp; ldc=nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
cudaFree(d_tmp_cc);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_J1[nO*nV*nV*i]); lda = nO*nV;
|
||||
B = &(d_J1_tmp[nV*nO*nV*i]); ldb = nV;
|
||||
C = &(d_J1[nO*nV*nV*i]); ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_T, nO*nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_J1_tmp);
|
||||
|
||||
double* d_X_voov;
|
||||
cudaMalloc((void **)&d_X_voov, nV*nO*nO*nV * sizeof(double));
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
beta = t1[j+b*nO];
|
||||
A = &(d_t2[nO*(j + nO*nV*b)]); lda = nO*nO;
|
||||
B = d_t1; ldb = nO;
|
||||
C = &(d_Y_ovov[nO*(b+nV*j)]); ldc = nO*nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_cc_space_v_vvoo[nV*(b+nV*nO*j)]); lda = nV*nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*(b+nV*nO*j)]); ldb = nV*nV;
|
||||
C = &(d_X_voov[nV*nO*(j+nO*b)]); ldc = nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nV, nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_Z_ovvo;
|
||||
cudaMalloc((void **)&d_Z_ovvo, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = -1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV; n=nV*nO; k=nO*nV;
|
||||
A=d_Y_ovov; lda=nO*nV;
|
||||
B=d_X_voov; ldb=nV*nO;
|
||||
C=d_Z_ovvo; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_voov);
|
||||
cudaFree(d_Y_ovov);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_J1[nO*nV*(b+nV*i)]); lda = nO;
|
||||
B = &(d_Z_ovvo[nO*(b+nV*nV*i)]); ldb=nO*nV;
|
||||
C = &(d_J1[nO*nV*(b+nV*i)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
double* d_Y_vovo;
|
||||
cudaMalloc((void **)&d_Y_vovo, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = -0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); lda = nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); ldb = nV;
|
||||
C = &(d_Y_vovo[nV*(i+nO*nV*j)]); ldc = nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_T, nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
double* d_X_ovvo;
|
||||
cudaMalloc((void **)&d_X_ovvo, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_t2[nO*(j+nO*nV*b)]); lda = nO*nO;
|
||||
B = &(d_t2[nO*(j+nO*nV*b)]); ldb = nO*nO;
|
||||
C = &(d_X_ovvo[nO*nV*(b+nV*j)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV; n=nV*nO; k=nV*nO;
|
||||
A=d_X_ovvo; lda=nO*nV;
|
||||
B=d_Y_vovo; ldb=nV*nO;
|
||||
C=d_Z_ovvo; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_X_ovvo);
|
||||
cudaFree(d_Y_vovo);
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_J1[nO*nV*(b+nV*i)]); lda = nO;
|
||||
B = &(d_Z_ovvo[nO*(b+nV*nV*i)]); ldb = nO*nV;
|
||||
C = &(d_J1[nO*nV*(b+nV*i)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
cudaFree(d_Z_ovvo);
|
||||
|
||||
lda = nO*nV;
|
||||
cublasGetMatrix(nO*nV, nV*nO, sizeof(double), d_J1, lda, J1, lda);
|
||||
cudaFree(d_J1);
|
||||
|
||||
}
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_K1;
|
||||
cudaMalloc((void **)&d_K1, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
A = d_cc_space_v_ovov; lda = nO*nV;
|
||||
B = d_cc_space_v_ovov; ldb = nO*nV;
|
||||
C = d_K1; ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO*nV, nO*nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
|
||||
alpha = -1.0;
|
||||
beta = 1.0;
|
||||
m=nO*nV*nO; n=nV; k=nO;
|
||||
A=d_cc_space_v_ovoo; lda=nO*nV*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_K1; ldc=nO*nV*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_X;
|
||||
cudaMalloc((void **)&d_X, nV*nO*nV*nO * sizeof(double));
|
||||
|
||||
double* d_Y;
|
||||
cudaMalloc((void **)&d_Y, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha =-1.0;
|
||||
beta = 0.0;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
cublasSetStream(handle, stream[i]);
|
||||
A = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); lda = nV;
|
||||
B = &(d_cc_space_v_vvoo[nV*nV*(i+nO*j)]); ldb = nV;
|
||||
C = &(d_X[nV*(j+nO*nV*i)]); ldc = nV*nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
alpha = 0.5;
|
||||
for (int j=0 ; j<nO ; ++j) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
beta = t1[j+b*nO];
|
||||
A = &(d_t2[nO*(j+nO*nV*b)]); lda = nO*nO;
|
||||
B = d_t1; ldb = nO;
|
||||
C = &(d_Y[nO*(b+nV*nV*j)]); ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_Z;
|
||||
cudaMalloc((void **)&d_Z, nO*nV*nV*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nV*nO; n=nO*nV; k=nV*nO;
|
||||
A=d_Y; lda=nO*nV;
|
||||
B=d_X; ldb=nV*nO;
|
||||
C=d_Z; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
cudaFree(d_X);
|
||||
cudaFree(d_Y);
|
||||
|
||||
|
||||
double* d_t1v;
|
||||
cudaMalloc((void **)&d_t1v, cholesky_mo_num*nO*nO * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=cholesky_mo_num*nO; n=nO; k=nV;
|
||||
A=d_cc_space_v_ov_chol; lda=cholesky_mo_num*nO;
|
||||
B=d_t1; ldb=nO;
|
||||
C=d_t1v; ldc=cholesky_mo_num*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_N, CUBLAS_OP_T, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
double* d_K1tmp;
|
||||
cudaMalloc((void **)&d_K1tmp, nO*nO*nV*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nO; n=nV*nV; k=cholesky_mo_num;
|
||||
A=d_t1v; lda=cholesky_mo_num;
|
||||
B=d_cc_space_v_vv_chol; ldb=cholesky_mo_num;
|
||||
C=d_K1tmp; ldc=nO*nO;
|
||||
cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
cudaFree(d_t1v);
|
||||
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = 1.0;
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_K1[nO*nV*(i+nO*b)]); lda = nO;
|
||||
B = &(d_K1tmp[nO*(i+nO*nV*b)]); ldb = nO*nO;
|
||||
C = &(d_K1[nO*nV*(i+nO*b)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i=0 ; i<nO ; ++i) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_K1[nO*nV*(i+nO*b)]); lda = nO;
|
||||
B = &(d_Z[nO*(b+nV*nV*i)]); ldb = nO*nV;
|
||||
C = &(d_K1[nO*nV*(i+nO*b)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_K1tmp);
|
||||
cudaFree(d_Z);
|
||||
|
||||
lda = nO*nV;
|
||||
cublasGetMatrix(nO*nV, nO*nV, sizeof(double), d_K1, lda, K1, lda);
|
||||
cudaFree(d_K1);
|
||||
|
||||
}
|
||||
|
||||
} // end sections
|
||||
|
||||
#pragma omp barrier
|
||||
|
||||
#pragma omp sections
|
||||
{
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_K1;
|
||||
lda = nO*nV;
|
||||
cudaMalloc((void **)&d_K1, nO*nV*nO*nV * sizeof(double));
|
||||
cublasSetMatrix(lda, nO*nV, sizeof(double), K1, lda, d_K1, lda);
|
||||
|
||||
double* d_J1;
|
||||
lda = nO*nV;
|
||||
cudaMalloc((void **)&d_J1, nO*nV*nV*nO * sizeof(double));
|
||||
@ -1034,7 +1030,89 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
|
||||
|
||||
}
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
double* d_X_ovov;
|
||||
cudaMalloc((void **)&d_X_ovov, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
double* d_Y_ovov;
|
||||
cudaMalloc((void **)&d_Y_ovov, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 0.5;
|
||||
beta = 0.0;
|
||||
for (int a=0 ; a<nV ; ++a) {
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
cublasSetStream(handle, stream[b]);
|
||||
A = &(d_K1[nO*(a+nV*nO*b)]); lda = nO*nV;
|
||||
B = &(d_K1[nO*(a+nV*nO*b)]); ldb = nO*nV;
|
||||
C = &(d_X_ovov[nO*(a+nV*nO*b)]); ldc = nO*nV;
|
||||
cublasDgeam(handle, CUBLAS_OP_T, CUBLAS_OP_N, nO, nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
alpha = 1.0;
|
||||
for (int v=0 ; v<nO ; ++v) {
|
||||
for (int g=0 ; g<nV ; ++g) {
|
||||
cublasSetStream(handle, stream[g]);
|
||||
A = &(d_t2[nO*(v+nO*g)]); lda = nO*nO*nV;
|
||||
B = &(d_t2[nO*(v+nO*g)]); ldb = nO*nO*nV;
|
||||
C = &(d_Y_ovov[nO*nV*(v+nO*g)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
|
||||
double* d_Z_ovov;
|
||||
cudaMalloc((void **)&d_Z_ovov, nO*nV*nO*nV * sizeof(double));
|
||||
|
||||
alpha = 1.0;
|
||||
beta = 0.0;
|
||||
m=nO*nV; n=nO*nV; k=nO*nV;
|
||||
A=d_X_ovov; lda=nO*nV;
|
||||
B=d_Y_ovov; ldb=nO*nV;
|
||||
C=d_Z_ovov; ldc=nO*nV;
|
||||
cublasDgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, k, &alpha, A, lda, B, ldb, &beta, C, ldc);
|
||||
|
||||
cudaFree(d_X_ovov);
|
||||
cudaFree(d_Y_ovov);
|
||||
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamCreate(&(stream[i]));
|
||||
}
|
||||
alpha = 1.0;
|
||||
beta = -1.0;
|
||||
for (int b=0 ; b<nV ; ++b) {
|
||||
for (int g=0 ; g<nV ; ++g) {
|
||||
cublasSetStream(handle, stream[g]);
|
||||
A = &(d_r2[nO*nO*(b+nV*g)]); lda = nO;
|
||||
B = &(d_Z_ovov[nO*(b+nV*nO*g)]); ldb = nO*nV;
|
||||
C = &(d_r2[nO*nO*(b+nV*g)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nO, nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
for (int g=0 ; g<nV ; ++g) {
|
||||
cublasSetStream(handle, stream[g]);
|
||||
A = &(d_r2[nO*nO*(b+nV*g)]); lda = nO;
|
||||
B = &(d_Z_ovov[nO*(g+nV*nO*b)]); ldb = nO*nV;
|
||||
C = &(d_r2[nO*nO*(b+nV*g)]); ldc = nO;
|
||||
cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_T, nO, nO, &alpha, A, lda, &beta, B, ldb, C, ldc);
|
||||
}
|
||||
}
|
||||
for (int i=0 ; i<nV ; ++i) {
|
||||
cudaStreamDestroy(stream[i]);
|
||||
}
|
||||
cublasSetStream(handle, NULL);
|
||||
cudaFree(d_Z_ovov);
|
||||
|
||||
}
|
||||
|
||||
#pragma omp section
|
||||
{
|
||||
}
|
||||
} // end sections
|
||||
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user