From 5c1d3987a3443dd6a1c3c5cd118782ce31327e9d Mon Sep 17 00:00:00 2001
From: Anthony Scemama <scemama@irsamc.ups-tlse.fr>
Date: Sun, 16 Jul 2023 13:47:44 +0200
Subject: [PATCH] GPU acceleration

---
 devel/ccsd_gpu/NEED  |  2 +-
 devel/ccsd_gpu/gpu.c | 23 ++++++++++++++++-------
 2 files changed, 17 insertions(+), 8 deletions(-)

diff --git a/devel/ccsd_gpu/NEED b/devel/ccsd_gpu/NEED
index e6e6bc5..cf913f8 100644
--- a/devel/ccsd_gpu/NEED
+++ b/devel/ccsd_gpu/NEED
@@ -1,2 +1,2 @@
 hartree_fock
-utils_cc
+utils_cc_gpu
diff --git a/devel/ccsd_gpu/gpu.c b/devel/ccsd_gpu/gpu.c
index ec8a100..c77a890 100644
--- a/devel/ccsd_gpu/gpu.c
+++ b/devel/ccsd_gpu/gpu.c
@@ -124,17 +124,26 @@ void compute_r2_space_chol_gpu(const int nO, const int nV, const int cholesky_mo
  //               cublasGetMatrix(m, n, sizeof(double), d_tmpB1, ldc, tmpB1, ldc);
 
 
-                for (size_t beta=iblock ; beta<(nV < iblock+BLOCK_SIZE ? nV : iblock+BLOCK_SIZE) ; ++beta) {
-                  for (size_t b=0 ; b<nV ; ++b) {
-                    cublasDcopy(handle, nV, &(d_tmpB1[nV*(beta-iblock + BLOCK_SIZE*b)]), 1,
-                       &(d_B1[nV*(b + nV*(beta-iblock))]),1);
+                for (size_t bet=iblock ; bet<(nV < iblock+BLOCK_SIZE ? nV : iblock+BLOCK_SIZE) ; ++bet) {
+
+           alpha = 1.0;
+           beta = 0.0;
+           A = &(d_tmpB1[nV*(bet-iblock)]); lda = nV*BLOCK_SIZE;
+           B = d_tmpB1; ldb = nV;
+           C = &(d_B1[nV*nV*(bet-iblock)]) ; ldc = nV;
+           cublasDgeam(handle, CUBLAS_OP_N, CUBLAS_OP_N, nV, nV, &alpha, A, lda, &beta, B, ldb, C, ldc);
+//                  for (size_t b=0 ; b<nV ; ++b) {
+/*
+                    cublasDcopy(handle, nV, &(d_tmpB1[nV*(bet-iblock + BLOCK_SIZE*b)]), 1,
+                       &(d_B1[nV*(b + nV*(bet-iblock))]),1);
+*/
 /*
                     for (size_t a=0 ; a<nV ; ++a) {
-                       B1[a + nV*(b + nV*(beta-iblock))] =
-                           tmpB1[a + nV*(beta-iblock + BLOCK_SIZE*b)];
+                       B1[a + nV*(b + nV*(bet-iblock))] =
+                           tmpB1[a + nV*(bet-iblock + BLOCK_SIZE*b)];
                     }
 */
-                  }
+//                  }
                 }
 //                cublasSetMatrix(nV*nV, mbs, sizeof(double), B1, nV*nV, d_B1, nV*nV);