Changes for HIP

PLEGMAConfig.cmake.in

@@ -0,0 +1,14 @@
+@PACKAGE_INIT@
+
+include(CMakeFindDependencyMacro)
+
+set(PLEGMA_TARGET_CUDA @PLEGMA_TARGET_CUDA@)
+set(PLEGMA_TARGET_HIP  @PLEGMA_TARGET_HIP@)
+
+if( PLEGMA_TARGET_CUDA )
+  include(${CMAKE_CURRENT_LIST_DIR}/find_target_cuda_dependencies.cmake)
+elseif(PLEGMA_TARGET_HIP )
+  include(${CMAKE_CURRENT_LIST_DIR}/find_target_hip_dependencies.cmake )
+endif()
+
+include(${CMAKE_CURRENT_LIST_DIR}/PLEGMATargets.cmake)

cmake/find_target_cuda_dependencies.cmake

@@ -0,0 +1,6 @@
+# CUDA Specific CMake
+
+enable_language(CUDA)
+
+find_dependency(CUDAToolkit REQUIRED)
+

cmake/find_target_hip_dependencies.cmake

@@ -0,0 +1,20 @@
+# HIP Specific CMake
+enable_language(HIP)
+
+if (NOT DEFINED ROCM_PATH )
+  if (NOT DEFINED ENV{ROCM_PATH} )
+    set(ROCM_PATH "/opt/rocm" CACHE PATH "ROCm path")
+  else()
+    set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "ROCm path")
+  endif()
+endif()
+
+set(CMAKE_MODULE_PATH "${ROCM_PATH}/lib/cmake" ${CMAKE_MODULE_PATH})
+find_dependency(HIP REQUIRED)
+find_dependency(hipfft REQUIRED)
+find_dependency(hiprand REQUIRED)
+find_dependency(rocrand REQUIRED)
+find_dependency(hipblas REQUIRED)
+find_dependency(rocblas REQUIRED)
+find_dependency(hipcub REQUIRED)
+find_dependency(rocprim REQUIRED)

include/PLEGMA_BLAS.h

@@ -5,13 +5,19 @@
 #if defined(HAVE_MKL)
 #include <mkl.h>
 #elif defined(HAVE_OPENBLAS)
-#include <cblas.h>
+#include <openblas/cblas.h>
 //#include <common.h> // do not know when is needed or not
 #else
 #error Neither mkl nor openBLAS have been defined
 #endif
 
+#if defined (__HIP__)
+#include <hipblas.h>
+#else
 #include <cublas_v2.h>
+//#else
+//#error Neither HIP nor NVCC have been defined
+#endif
 #include <mpi.h>
 #include <PLEGMA_utils.h>
 #include <quda_api.h>
@@ -97,160 +103,12 @@ namespace cBLAS{
 }
 //=================================================================//
 
+#if defined (__HIP__)
+#include "target/hip/PLEGMA_hipBLAS.h"
+#else
+#include "target/cuda/PLEGMA_cuBLAS.h"
+#endif
 
-namespace cuBLAS{
-
-  template<typename Float>
-  inline void axpy(int NN, Float val[2], Float *x, Float *y){}
-  template<>
-  inline void axpy<float>(int NN, float val[2], float *x, float *y){
-    cuComplex cu_val = make_cuComplex(val[0],val[1]);
-    cublasStatus_t error =  cublasCaxpy(HGC_cublas_handle,NN, &cu_val, (cuComplex*) x, 1, (cuComplex*) y, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCaxpy failed with error %d", error);
-  }
-  template<>
-  inline void axpy<double>(int NN, double val[2], double *x, double *y){
-    cuDoubleComplex cu_val = make_cuDoubleComplex(val[0],val[1]);
-    cublasStatus_t error =  cublasZaxpy(HGC_cublas_handle, NN, &cu_val,(cuDoubleComplex*) x, 1, (cuDoubleComplex*) y, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasZaxpy failed with error %d", error);
-  }
-
-  //------------------------------------------------------------------
-  template<typename Float>
-  inline void scal(int NN, const Float val, Float *x);
-  template<>
-  inline void scal<float>(int NN, const float val, float *x){
-    cublasStatus_t error = cublasCsscal(HGC_cublas_handle, NN, &val, (cuComplex*) x, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCsscal failed with error %d", error);
-  }
-  template<>
-  inline void scal<double>(int NN, const double val, double *x){
-    cublasStatus_t error = cublasZdscal(HGC_cublas_handle, NN, &val, (cuDoubleComplex*) x, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasZdscal failed with error %d", error);
-  }
-  //------------------------------------------------------------------
-  template<typename Float>
-  inline void cscal(int NN, const Float val[2], Float *x);
-  template<>
-  inline void cscal<float>(int NN, const float val[2], float *x){
-    cuComplex cu_val = make_cuComplex(val[0],val[1]);
-    cublasStatus_t error = cublasCscal(HGC_cublas_handle, NN, &cu_val, (cuComplex*) x, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCscal failed with error %d", error);
-  }
-  template<>
-  inline void cscal<double>(int NN, const double val[2], double *x){
-    cuDoubleComplex cu_val = make_cuDoubleComplex(val[0],val[1]);
-    cublasStatus_t error = cublasZscal(HGC_cublas_handle, NN, &cu_val, (cuDoubleComplex*) x, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasZscal failed with error %d", error);
-  }
-  //-----------------------------------------------------------------
-  template<typename Float>
-  inline std::complex<Float> dot(int NN, const Float *x, const Float *y);
-  template<>
-  inline std::complex<float> dot<float>(int NN, const float *x, const float *y){
-    cuComplex cu_res;
-    cublasStatus_t error = cublasCdotc(HGC_cublas_handle, NN,(cuComplex*)x, 1, (cuComplex*)y,1,&cu_res);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCdotc failed with error %d", error);
-    return std::complex<float>(cu_res.x,cu_res.y);
-  }
-  template<>
-  inline std::complex<double> dot<double>(int NN, const double *x, const double *y){
-    cuDoubleComplex cu_res;
-    cublasStatus_t error = cublasZdotc(HGC_cublas_handle, NN,(cuDoubleComplex*)x, 1, (cuDoubleComplex*)y,1,&cu_res);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasZdotc failed with error %d", error);
-    return std::complex<double>(cu_res.x,cu_res.y);
-  }  
-  template<typename Float>
-  inline std::complex<Float> dot(int NN, const Float *x, const Float *y, MPI_Comm comm) {
-    if(comm == MPI_COMM_NULL) PLEGMA_error("Communicator is NULL and cannot be used for MPI reduction");
-    std::complex<Float> result, res = cuBLAS::dot(NN, x, y);
-    int mpiErr = MPI_Allreduce((Float*) &res, (Float*) &result, 2,
-  			       MPI_Type<Float>(), MPI_SUM, comm);
-    if(mpiErr != MPI_SUCCESS) PLEGMA_error("MPI_Allreduce failed with error %d\n", mpiErr);
-    return result;
-  }  
-  //-----------------------------------------------------------------
-  template<typename Float>
-  inline Float norm(int NN, const Float *x);
-  template<>
-  inline float norm<float>(int NN, const float *x){
-    float res;
-    cublasStatus_t error = cublasScnrm2(HGC_cublas_handle, NN, (cuComplex*)x, 1, &res);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCdotc failed with error %d", error);
-    return res;
-  }
-  template<>
-  inline double norm<double>(int NN, const double *x){
-    double res;
-    cublasStatus_t error = cublasDznrm2(HGC_cublas_handle, NN, (cuDoubleComplex*)x, 1, &res);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCdotc failed with error %d", error);
-    return res;
-  }
-  template<typename Float>
-  inline Float norm(int NN, const Float *x, MPI_Comm comm) {
-    if(comm == MPI_COMM_NULL) PLEGMA_error("Communicator is NULL and cannot be used for MPI reduction");
-    Float result, loc_res = std::pow(cuBLAS::norm(NN, x),2);
-    int mpiErr = MPI_Allreduce(&loc_res, &result, 1, MPI_Type<Float>(), MPI_SUM,
-			       comm);
-    if(mpiErr != MPI_SUCCESS) PLEGMA_error("MPI_Allreduce failed with error %d\n", mpiErr);
-    return sqrt(result);
-  }
-
-
-  //---------------------------------------------------------
-  template<typename Float>
-  inline void gemv_(OPER_MATR_BLAS trans, int m, int n, Float alpha[2], Float* A, Float* x, Float beta[2], Float* y){}
-
-  template<>
-  inline void gemv_<float>(OPER_MATR_BLAS trans, int m, int n, float alpha[2], float* A, float* x, float beta[2], float* y){
-    cublasOperation_t Oper;
-    cuComplex cu_alpha = make_cuComplex(alpha[0],alpha[1]);
-    cuComplex cu_beta = make_cuComplex(beta[0],beta[1]);
-    switch(trans){case(NOTRANS): Oper=CUBLAS_OP_N; break; case(TRANS): Oper=CUBLAS_OP_T; break; case(DAGGER): Oper=CUBLAS_OP_C; break;}
-    cublasStatus_t error = cublasCgemv(HGC_cublas_handle, Oper, m, n, &cu_alpha, (cuComplex*) A, m, (cuComplex*) x, 1, &cu_beta, (cuComplex*) y, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasCgemv failed with error %d", error);
-  }
-
-  template<>
-  inline void gemv_<double>(OPER_MATR_BLAS trans, int m, int n, double alpha[2], double* A, double* x, double beta[2], double* y){
-    cublasOperation_t Oper;
-    cuDoubleComplex cu_alpha = make_cuDoubleComplex(alpha[0],alpha[1]);
-    cuDoubleComplex cu_beta = make_cuDoubleComplex(beta[0],beta[1]);
-    switch(trans){case(NOTRANS): Oper=CUBLAS_OP_N; break; case(TRANS): Oper=CUBLAS_OP_T; break; case(DAGGER): Oper=CUBLAS_OP_C; break;}
-    cublasStatus_t error = cublasZgemv(HGC_cublas_handle, Oper, m, n, &cu_alpha, (cuDoubleComplex*) A, m,(cuDoubleComplex*) x, 1, &cu_beta,(cuDoubleComplex*) y, 1);
-    if(error != CUBLAS_STATUS_SUCCESS) PLEGMA_error("cublasZgemv failed with error %d", error);
-  }
-
-  // In case of m is partitioned and trans=NOTRANS OR m is not partitioned one can use whatever trans
-  // Cannot work if n is partitioned
-  template<typename Float>
-  inline void gemv(OPER_MATR_BLAS trans, int m, int n, Float alpha[2], Float* A, Float* x, Float beta[2], Float* y){
-    gemv_(trans, m, n, alpha, A, x, beta, y);
-  }
-  // In case that m is partitioned and we take trans of dagger then reduction is needed
-  // Cannot work if n is partitioned
-  template<typename Float>
-  inline void gemv(OPER_MATR_BLAS trans, int m, int n, Float alpha[2], Float* A, Float* x, Float beta[2], Float* y, Float *yHost, MPI_Comm comm){
-    if(trans == NOTRANS) PLEGMA_error("Use gemv without MPI comm");
-    if(comm == MPI_COMM_NULL) PLEGMA_error("Communicator is NULL and cannot be used for MPI reduction");
-    cuBLAS::gemv_(trans, m, n, alpha, A, x, beta, y);
-    qudaMemcpy(yHost,y,n*2*sizeof(Float),qudaMemcpyDeviceToHost);
-    checkQudaError();
-    int mpiErr = MPI_Allreduce(MPI_IN_PLACE,yHost,n*2,MPI_Type(yHost),MPI_SUM,comm);
-    if(mpiErr != MPI_SUCCESS) PLEGMA_error("MPI_Allreduce failed with error %d\n", mpiErr);
-  }
-
-  // In case that m is partitioned and we take trans of dagger then reduction is needed
-  // Cannot work if n is partitioned
-  template<typename Float>
-  inline void gemv(OPER_MATR_BLAS trans, int m, int n, Float alpha[2], Float* A, Float* x, Float beta[2], Float* y, MPI_Comm comm){
-    Float yHost[n*2];
-    cuBLAS::gemv(trans,m, n, alpha, A, x, beta, y, yHost,comm);
-    qudaMemcpy(y,yHost,sizeof(yHost),qudaMemcpyHostToDevice);
-    checkQudaError();
-  }
-}
-//=================================================================//
 
 namespace plegma{
   template<typename Float>

include/PLEGMA_FT.h

@@ -63,11 +63,17 @@ namespace plegma {
        @params bool accum = false: In case we want to accumulation results from each transformation on the class buffer
        @params bool dimT = HGC_localL[DIM_T]: Size of the time dimension in case we want to transform only part of the vector
      **/
-    template<typename T>
-    PLEGMA_FT(std::vector<T> mom, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+    PLEGMA_FT(std::vector<int> mom, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
 
-    template<typename T>
-    PLEGMA_FT( std::vector<std::vector<T>> &moms, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+    PLEGMA_FT( std::vector<std::vector<int>> &moms, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+
+    PLEGMA_FT(std::vector<float> mom, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+
+    PLEGMA_FT( std::vector<std::vector<float>> &moms, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+
+    PLEGMA_FT(std::vector<double> mom, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
+
+    PLEGMA_FT( std::vector<std::vector<double>> &moms, int D3D4 = 3, bool accum = false, int dimT = HGC_localL[DIM_T]);
 
     ~PLEGMA_FT() {};
 

include/PLEGMA_Field.h

@@ -113,11 +113,20 @@ namespace plegma {
     /**
        @brief Creates a texture object which binds on field elements on GPU
      */
+#ifdef __NVCC__
     cudaTextureObject_t createTexObject() const;
+#elif defined (__HIP__)
+    hipTextureObject_t createTexObject() const;
+#endif
     /**
        @brief Destroys the texture object which binds on field elements on GPU
      */
+#ifdef __NVCC__
     void destroyTexObject(cudaTextureObject_t tex) const;
+#elif defined (__HIP__)
+    void destroyTexObject(hipTextureObject_t tex) const;
+#endif
+
     /**
      * @return a pointer to access Host elements of the field
      */

include/PLEGMA_Hprobing.h

@@ -1,4 +1,5 @@
 #pragma once
+#include <malloc_quda.h>
 //#include <PLEGMA_utils.h>
 #include <tune_quda.h>
 
@@ -101,15 +102,15 @@ namespace plegma {
       createElemColBlock();
       createColLattice();
       //  if(check)checkColoring();
-      cudaMalloc((void**)&d_arrVc, HGC_localVolume*sizeof(int));
+      d_arrVc=(int*)device_malloc(HGC_localVolume*sizeof(int));
 //      checkQudaError();
-      cudaMemcpy(d_arrVc, h_arrVc, HGC_localVolume*sizeof(int), cudaMemcpyHostToDevice);
+      qudaMemcpy(d_arrVc, h_arrVc, HGC_localVolume*sizeof(int), qudaMemcpyHostToDevice);
 //      checkQudaError();    
     }
     ~PLEGMA_Hprobing(){
       delete[] h_arrVc;
       delete[] arrlc;
-      cudaFree(d_arrVc);
+      device_free(d_arrVc);
       //checkQudaError();
     }
     int* H_arrVc() const{return h_arrVc;}