add ispin to transferT0MatrixToGPUHIP

Author: meisenbach

src/MultipleScattering/calculateTauMatrix.cpp

@@ -464,7 +464,7 @@ void calculateTauMatrix(LSMSSystemParameters &lsms, LocalTypeInfo &local,
           devM = deviceStorage->getDevM();
           transferMatrixToGPUHip(devM, m);
           devT0 = deviceStorage->getDevT0();
-          transferT0MatrixToGPUHip(devT0, lsms, local, atom, iie);
+          transferT0MatrixToGPUHip(devT0, lsms, local, atom, iie, ispin);
           break;
 #endif
         default:
@@ -520,7 +520,7 @@ void calculateTauMatrix(LSMSSystemParameters &lsms, LocalTypeInfo &local,
           break;
         case MST_LINEAR_SOLVER_ZGETRF_ROCSOLVER:
           devT0 = deviceStorage->getDevT0();
-          transferT0MatrixToGPUHip(devT0, lsms, local, atom, iie);
+          transferT0MatrixToGPUHip(devT0, lsms, local, atom, iie, ispin);
           break;
 #ifdef ACCELERATOR_CUDA_C
         case MST_LINEAR_SOLVER_ZGETRF_CUBLAS:

src/MultipleScattering/linearSolvers.hpp

@@ -145,7 +145,7 @@ void unitMatrixCuda(T *devM, int lDim, int nCol) {
 #ifdef ACCELERATOR_HIP
 void transferMatrixToGPUHip(Complex *devM, Matrix<Complex> &m);
 void transferMatrixFromGPUHip(Matrix<Complex> &m, hipDoubleComplex *devM);
-void transferT0MatrixToGPUHip(Complex *devT0, LSMSSystemParameters &lsms, LocalTypeInfo &local, AtomData &atom, int iie);
+void transferT0MatrixToGPUHip(Complex *devT0, LSMSSystemParameters &lsms, LocalTypeInfo &local, AtomData &atom, int iie, int ispin);
 void transferFullTMatrixToGPUHip(Complex *devT, LSMSSystemParameters &lsms, LocalTypeInfo &local,
                                   AtomData &atom, int ispin);
 

src/MultipleScattering/linearSolvers_HIP.cpp

@@ -71,10 +71,12 @@ __global__ void zeroDiagonalBlocksKernelHip(T *devM, int lDim, int nCol,
 }
 
 void transferT0MatrixToGPUHip(Complex *devT0, LSMSSystemParameters &lsms,
-                              LocalTypeInfo &local, AtomData &atom, int iie) {
+                              LocalTypeInfo &local, AtomData &atom, int iie, int ispin) {
   int kkrsz_ns = lsms.n_spin_cant * atom.kkrsz;
+  int jsm =  kkrsz_ns * kkrsz_ns * ispin;
+
   hipError_t ret = hipMemcpy(devT0,
-            &local.tmatStore(iie * local.blkSizeTmatStore, atom.LIZStoreIdx[0]),
+            &local.tmatStore(iie * local.blkSizeTmatStore + jsm, atom.LIZStoreIdx[0]),
             kkrsz_ns * kkrsz_ns * sizeof(hipDoubleComplex),
             hipMemcpyHostToDevice);
 }