hypre-space · abagusetty · Jul 24, 2021 · Jul 27, 2021 · Jul 27, 2021 · Jul 28, 2021
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -106,6 +106,9 @@ option(HYPRE_WITH_CUDA               "Use CUDA. Require cuda-8.0 or higher" OFF)
 option(HYPRE_WITH_SYCL               "Use SYCL" OFF)
 option(HYPRE_ENABLE_UNIFIED_MEMORY   "Use unified memory for allocating the memory" OFF)
 # CUDA options
+option(HYPRE_ENABLE_ONEMKLSPARSE     "Use oneMKLSPARSE" OFF)
+option(HYPRE_ENABLE_ONEMKLBLAS       "Use oneMKLBLAS" OFF)
+option(HYPRE_ENABLE_ONEMKLRAND        "Use oneMKLRNG" ON)
 option(HYPRE_ENABLE_CUDA_STREAMS     "Use CUDA streams" ON)
 option(HYPRE_ENABLE_CUSPARSE         "Use cuSPARSE" ON)
 option(HYPRE_ENABLE_DEVICE_POOL      "Use device memory pool" OFF)
@@ -289,7 +292,7 @@ endif (HYPRE_WITH_CUDA)
 # SYCL
 if (HYPRE_WITH_SYCL)
   enable_language(CXX)
-  message(STATUS "Enabled support for CXX.")
+  message(STATUS "Enabled support for SYCL.")
 
   # Enforce C++17
   if (NOT CMAKE_CXX_STANDARD OR CMAKE_CXX_STANDARD LESS 17)
@@ -299,17 +302,14 @@ if (HYPRE_WITH_SYCL)
 
   message(STATUS "Using CXX standard: c++${CMAKE_CXX_STANDARD}")
 
-  # Set CXX compiler to dpcpp
-  set(CMAKE_CXX_COMPILER "dpcpp")
-
   # Add any extra CXX compiler flags HYPRE_WITH_EXTRA_CXXFLAGS
   if (NOT HYPRE_WITH_EXTRA_CXXFLAGS STREQUAL "")
     string(REPLACE " " ";" HYPRE_WITH_EXTRA_CXXFLAGS "${HYPRE_WITH_EXTRA_CXXFLAGS}")
     add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:${HYPRE_WITH_EXTRA_CXXFLAGS}>")
   endif ()
 
-  set(HYPRE_USING_SYCL ON CACHE BOOL "" FORCE)
-  set(HYPRE_USING_GPU ON CACHE BOOL "" FORCE)
+  set(HYPRE_USING_SYCL         ON CACHE BOOL "" FORCE)
+  set(HYPRE_USING_GPU          ON CACHE BOOL "" FORCE)
 
   if (HYPRE_ENABLE_UNIFIED_MEMORY)
     set(HYPRE_USING_UNIFIED_MEMORY ON CACHE BOOL "" FORCE)
@@ -324,12 +324,17 @@ if (HYPRE_WITH_SYCL)
     set(HYPRE_BUILD_EXAMPLES OFF CACHE BOOL "" FORCE)
   endif ()
 
+  add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-D_GLIBCXX_USE_TBB_PAR_BACKEND=0>")
+  add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-sycl-std=2020>")
   add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-fsycl>")
   add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-fsycl-unnamed-lambda>")
+  add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-fsycl-device-only>")
   add_compile_options("$<$<COMPILE_LANGUAGE:CXX>:-fsycl-device-code-split=per_kernel>")
 
   set(HYPRE_USING_HOST_MEMORY OFF CACHE BOOL "" FORCE)
 
+#  include(HYPRE_SetupSYCLToolkit)
+#  set(CMAKE_CXX_FLAGS "-I/home/abagusetty/build/oneDPL/build_sycl/include ${CMAKE_CXX_FLAGS}")
 endif (HYPRE_WITH_SYCL)
 
 # Add any extra C compiler flags HYPRE_WITH_EXTRA_CFLAGS
@@ -477,10 +482,6 @@ if (HYPRE_USING_CUDA)
   set_source_files_properties(${HYPRE_GPU_SOURCES} PROPERTIES LANGUAGE CUDA)
 endif ()
 
-if (HYPRE_USING_SYCL)
-  set_source_files_properties(${HYPRE_GPU_SOURCES} PROPERTIES LANGUAGE CXX)
-endif ()
-
 # Set MPI compile flags
 if (NOT HYPRE_SEQUENTIAL)
   find_program(MPIEXEC_EXECUTABLE NAMES mpiexec mpirun)

diff --git a/src/IJ_mv/HYPRE_IJVector.c b/src/IJ_mv/HYPRE_IJVector.c
@@ -240,7 +240,7 @@ HYPRE_IJVectorSetValues( HYPRE_IJVector        vector,
 
    if ( hypre_IJVectorObjectType(vec) == HYPRE_PARCSR )
    {
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
       HYPRE_ExecutionPolicy exec = hypre_GetExecPolicy1( hypre_IJVectorMemoryLocation(vector) );
 
       if (exec == HYPRE_EXEC_DEVICE)
@@ -295,7 +295,7 @@ HYPRE_IJVectorAddToValues( HYPRE_IJVector        vector,
 
    if ( hypre_IJVectorObjectType(vec) == HYPRE_PARCSR )
    {
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
       HYPRE_ExecutionPolicy exec = hypre_GetExecPolicy1( hypre_IJVectorMemoryLocation(vector) );
 
       if (exec == HYPRE_EXEC_DEVICE)
@@ -333,7 +333,7 @@ HYPRE_IJVectorAssemble( HYPRE_IJVector vector )
 
    if ( hypre_IJVectorObjectType(vec) == HYPRE_PARCSR )
    {
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
       HYPRE_ExecutionPolicy exec = hypre_GetExecPolicy1( hypre_IJVectorMemoryLocation(vector) );
 
       if (exec == HYPRE_EXEC_DEVICE)

diff --git a/src/IJ_mv/IJMatrix_parcsr_device.c b/src/IJ_mv/IJMatrix_parcsr_device.c
@@ -164,8 +164,8 @@ hypre_IJMatrixSetAddValuesParCSRDevice( hypre_IJMatrix       *matrix,
       /* mark unwanted elements as -1 */
       dim3 bDim = hypre_GetDefaultDeviceBlockDimension();
       dim3 gDim = hypre_GetDefaultDeviceGridDimension(len1, "thread", bDim);
-      HYPRE_CUDA_LAUNCH( hypreCUDAKernel_IJMatrixValues_dev1, gDim, bDim, len1, indicator,
-                         (HYPRE_Int *) row_indexes, ncols, indicator );
+      HYPRE_GPU_LAUNCH( hypreCUDAKernel_IJMatrixValues_dev1, gDim, bDim, len1, indicator,
+                        (HYPRE_Int *) row_indexes, ncols, indicator );
 
       auto new_end = HYPRE_THRUST_CALL(
                         copy_if,
@@ -233,7 +233,7 @@ hypre_IJMatrixAssembleSortAndReduce1(HYPRE_Int  N0, HYPRE_BigInt  *I0, HYPRE_Big
    /*
    dim3 bDim = hypre_GetDefaultDeviceBlockDimension();
    dim3 gDim = hypre_GetDefaultDeviceGridDimension(N0, "thread", bDim);
-   HYPRE_CUDA_LAUNCH( hypreCUDAKernel_IJMatrixAssembleSortAndReduce1, gDim, bDim, N0, I0, J0, X0, A0 );
+   HYPRE_GPU_LAUNCH( hypreCUDAKernel_IJMatrixAssembleSortAndReduce1, gDim, bDim, N0, I0, J0, X0, A0 );
    */
 
    /* output X: 0: keep, 1: zero-out */

diff --git a/src/IJ_mv/IJVector_parcsr.c b/src/IJ_mv/IJVector_parcsr.c
@@ -124,7 +124,7 @@ HYPRE_Int
 hypre_IJVectorSetMaxOffProcElmtsPar(hypre_IJVector *vector,
                                     HYPRE_Int       max_off_proc_elmts)
 {
-   hypre_AuxParVector *aux_vector;
+   hypre_AuxParVector *aux_vector = NULL;
 
    aux_vector = (hypre_AuxParVector*) hypre_IJVectorTranslator(vector);
    if (!aux_vector)
@@ -134,7 +134,7 @@ hypre_IJVectorSetMaxOffProcElmtsPar(hypre_IJVector *vector,
    }
    hypre_AuxParVectorMaxOffProcElmts(aux_vector) = max_off_proc_elmts;
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
    hypre_AuxParVectorUsrOffProcElmts(aux_vector) = max_off_proc_elmts;
 #endif
 
@@ -156,7 +156,7 @@ hypre_IJVectorDistributePar(hypre_IJVector  *vector,
                             const HYPRE_Int *vec_starts)
 {
    hypre_ParVector *old_vector = (hypre_ParVector*) hypre_IJVectorObject(vector);
-   hypre_ParVector *par_vector;
+   hypre_ParVector *par_vector = NULL;
    HYPRE_Int print_level = hypre_IJVectorPrintLevel(vector);
 
    if (!old_vector)
@@ -208,8 +208,8 @@ hypre_IJVectorZeroValuesPar(hypre_IJVector *vector)
 
    hypre_ParVector *par_vector = (hypre_ParVector*) hypre_IJVectorObject(vector);
    MPI_Comm comm = hypre_IJVectorComm(vector);
-   HYPRE_BigInt *partitioning;
-   hypre_Vector *local_vector;
+   HYPRE_BigInt *partitioning = NULL;
+   hypre_Vector *local_vector = NULL;
    HYPRE_Int print_level = hypre_IJVectorPrintLevel(vector);
 
    hypre_MPI_Comm_rank(comm, &my_id);
@@ -281,13 +281,13 @@ hypre_IJVectorSetValuesPar(hypre_IJVector       *vector,
    HYPRE_Int my_id;
    HYPRE_Int j, k;
    HYPRE_BigInt i, vec_start, vec_stop;
-   HYPRE_Complex *data;
+   HYPRE_Complex *data = NULL;
    HYPRE_Int print_level = hypre_IJVectorPrintLevel(vector);
 
    HYPRE_BigInt *IJpartitioning = hypre_IJVectorPartitioning(vector);
    hypre_ParVector *par_vector = (hypre_ParVector*) hypre_IJVectorObject(vector);
    MPI_Comm comm = hypre_IJVectorComm(vector);
-   hypre_Vector *local_vector;
+   hypre_Vector *local_vector = NULL;
 
    /* If no components are to be set, perform no checking and return */
    if (num_values < 1) { return 0; }
@@ -394,14 +394,14 @@ hypre_IJVectorAddToValuesPar(hypre_IJVector       *vector,
 {
    HYPRE_Int my_id;
    HYPRE_Int i, j, vec_start, vec_stop;
-   HYPRE_Complex *data;
+   HYPRE_Complex *data = NULL;
    HYPRE_Int print_level = hypre_IJVectorPrintLevel(vector);
 
    HYPRE_BigInt *IJpartitioning = hypre_IJVectorPartitioning(vector);
    hypre_ParVector *par_vector = (hypre_ParVector*) hypre_IJVectorObject(vector);
    hypre_AuxParVector *aux_vector = (hypre_AuxParVector*) hypre_IJVectorTranslator(vector);
    MPI_Comm comm = hypre_IJVectorComm(vector);
-   hypre_Vector *local_vector;
+   hypre_Vector *local_vector = NULL;
 
    /* If no components are to be retrieved, perform no checking and return */
    if (num_values < 1) { return 0; }
@@ -557,8 +557,8 @@ hypre_IJVectorAssemblePar(hypre_IJVector *vector)
    {
       HYPRE_Int off_proc_elmts, current_num_elmts;
       HYPRE_Int max_off_proc_elmts;
-      HYPRE_BigInt *off_proc_i;
-      HYPRE_Complex *off_proc_data;
+      HYPRE_BigInt *off_proc_i = NULL;
+      HYPRE_Complex *off_proc_data = NULL;
       current_num_elmts = hypre_AuxParVectorCurrentOffProcElmts(aux_vector);
       hypre_MPI_Allreduce(&current_num_elmts, &off_proc_elmts, 1, HYPRE_MPI_INT,
                           hypre_MPI_SUM, comm);
@@ -704,12 +704,12 @@ hypre_IJVectorAssembleOffProcValsPar( hypre_IJVector       *vector,
    HYPRE_Int  first_index;
 
    void *void_contact_buf = NULL;
-   void *index_ptr;
-   void *recv_data_ptr;
+   void *index_ptr = NULL;
+   void *recv_data_ptr = NULL;
 
    HYPRE_Complex tmp_complex;
    HYPRE_BigInt *ex_contact_buf = NULL;
-   HYPRE_Complex *vector_data;
+   HYPRE_Complex *vector_data = NULL;
    HYPRE_Complex value;
 
    hypre_DataExchangeResponse      response_obj1, response_obj2;
@@ -718,7 +718,7 @@ hypre_IJVectorAssembleOffProcValsPar( hypre_IJVector       *vector,
    MPI_Comm comm = hypre_IJVectorComm(vector);
    hypre_ParVector *par_vector = (hypre_ParVector*) hypre_IJVectorObject(vector);
 
-   hypre_IJAssumedPart   *apart;
+   hypre_IJAssumedPart   *apart = NULL;
 
    hypre_MPI_Comm_rank(comm, &myid);
 
@@ -1119,7 +1119,7 @@ hypre_IJVectorAssembleOffProcValsPar( hypre_IJVector       *vector,
       hypre_TMemcpy(off_proc_data_recv_d, off_proc_data_recv, HYPRE_Complex, off_proc_nelm_recv_cur,
                     HYPRE_MEMORY_DEVICE, HYPRE_MEMORY_HOST);
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
       hypre_IJVectorSetAddValuesParDevice(vector, off_proc_nelm_recv_cur, off_proc_i_recv_d,
                                           off_proc_data_recv_d, "add");
 #endif

diff --git a/src/IJ_mv/IJVector_parcsr_device.c b/src/IJ_mv/IJVector_parcsr_device.c
@@ -251,9 +251,9 @@ hypre_IJVectorAssembleParDevice(hypre_IJVector *vector)
       /* set/add to local vector */
       dim3 bDim = hypre_GetDefaultDeviceBlockDimension();
       dim3 gDim = hypre_GetDefaultDeviceGridDimension(new_nnz, "thread", bDim);
-      HYPRE_CUDA_LAUNCH( hypreCUDAKernel_IJVectorAssemblePar, gDim, bDim, new_nnz, new_data, new_i,
-                         vec_start, new_sora,
-                         hypre_VectorData(hypre_ParVectorLocalVector(par_vector)) );
+      HYPRE_GPU_LAUNCH( hypreCUDAKernel_IJVectorAssemblePar, gDim, bDim, new_nnz, new_data, new_i,
+                        vec_start, new_sora,
+                        hypre_VectorData(hypre_ParVectorLocalVector(par_vector)) );
 
       hypre_TFree(new_i,    HYPRE_MEMORY_DEVICE);
       hypre_TFree(new_data, HYPRE_MEMORY_DEVICE);

diff --git a/src/IJ_mv/_hypre_IJ_mv.h b/src/IJ_mv/_hypre_IJ_mv.h
@@ -170,7 +170,7 @@ typedef struct
 
    HYPRE_MemoryLocation memory_location;
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
    HYPRE_Int            max_stack_elmts;      /* length of stash for SetValues and AddToValues*/
    HYPRE_Int            current_stack_elmts;  /* current no. of elements stored in stash */
    HYPRE_BigInt        *stack_i;              /* contains row indices */
@@ -193,7 +193,7 @@ typedef struct
 
 #define hypre_AuxParVectorMemoryLocation(vector)       ((vector) -> memory_location)
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
 #define hypre_AuxParVectorMaxStackElmts(vector)        ((vector) -> max_stack_elmts)
 #define hypre_AuxParVectorCurrentStackElmts(vector)    ((vector) -> current_stack_elmts)
 #define hypre_AuxParVectorStackI(vector)               ((vector) -> stack_i)

diff --git a/src/IJ_mv/aux_par_vector.c b/src/IJ_mv/aux_par_vector.c
@@ -32,7 +32,7 @@ hypre_AuxParVectorCreate( hypre_AuxParVector **aux_vector)
    hypre_AuxParVectorOffProcI(vector) = NULL;
    hypre_AuxParVectorOffProcData(vector) = NULL;
    hypre_AuxParVectorMemoryLocation(vector) = HYPRE_MEMORY_HOST;
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
    hypre_AuxParVectorMaxStackElmts(vector) = 0;
    hypre_AuxParVectorCurrentStackElmts(vector) = 0;
    hypre_AuxParVectorStackI(vector) = NULL;
@@ -61,7 +61,7 @@ hypre_AuxParVectorDestroy( hypre_AuxParVector *vector )
       hypre_TFree(hypre_AuxParVectorOffProcI(vector),    HYPRE_MEMORY_HOST);
       hypre_TFree(hypre_AuxParVectorOffProcData(vector), HYPRE_MEMORY_HOST);
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
       hypre_TFree(hypre_AuxParVectorStackI(vector),    hypre_AuxParVectorMemoryLocation(vector));
       hypre_TFree(hypre_AuxParVectorStackData(vector), hypre_AuxParVectorMemoryLocation(vector));
       hypre_TFree(hypre_AuxParVectorStackSorA(vector), hypre_AuxParVectorMemoryLocation(vector));

diff --git a/src/IJ_mv/aux_par_vector.h b/src/IJ_mv/aux_par_vector.h
@@ -30,7 +30,7 @@ typedef struct
 
    HYPRE_MemoryLocation memory_location;
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
    HYPRE_Int            max_stack_elmts;      /* length of stash for SetValues and AddToValues*/
    HYPRE_Int            current_stack_elmts;  /* current no. of elements stored in stash */
    HYPRE_BigInt        *stack_i;              /* contains row indices */
@@ -53,7 +53,7 @@ typedef struct
 
 #define hypre_AuxParVectorMemoryLocation(vector)       ((vector) -> memory_location)
 
-#if defined(HYPRE_USING_CUDA) || defined(HYPRE_USING_HIP)
+#if defined(HYPRE_USING_GPU)
 #define hypre_AuxParVectorMaxStackElmts(vector)        ((vector) -> max_stack_elmts)
 #define hypre_AuxParVectorCurrentStackElmts(vector)    ((vector) -> current_stack_elmts)
 #define hypre_AuxParVectorStackI(vector)               ((vector) -> stack_i)

diff --git a/src/config/HYPRE_config.h.in b/src/config/HYPRE_config.h.in
@@ -220,7 +220,7 @@
 /* onemkl::BLAS being used */
 #undef HYPRE_USING_ONEMKLBLAS
 
-/* onemkl::rng being used */
+/* onemkl::rng::device being used */
 #undef HYPRE_USING_ONEMKLRAND
 
 /* onemkl::SPARSE being used */

diff --git a/src/config/configure.in b/src/config/configure.in
@@ -209,7 +209,7 @@ dnl *********************************************************************
 hypre_using_sycl=no
 hypre_using_onemklsparse=no
 hypre_using_onemklblas=no
-hypre_using_onemklrand=no
+hypre_using_onemklrand=yes
 
 hypre_found_mkl=no
 
@@ -2341,20 +2341,20 @@ dnl *********************************************************************
 dnl * Set SYCL options
 dnl *********************************************************************
 AS_IF([test x"$hypre_using_sycl" == x"yes"],
-      [
+[
         AC_DEFINE(HYPRE_USING_GPU,  1, [Define to 1 if executing on GPU device])
         AC_DEFINE(HYPRE_USING_SYCL, 1, [SYCL being used])
 
         dnl (Ab)Using CUCC when compiling HIP
-	     LINK_CC=${CUCC}
+        LINK_CC=${CUCC}
         LINK_CXX=${CUCC}
 
-        SYCLFLAGS="-fsycl -fsycl-unnamed-lambda -fsycl-device-code-split=per_kernel"
+        SYCLFLAGS="-Wunknown-pragmas -Wunused-local-typedef -Wunused-lambda-capture -D_GLIBCXX_USE_TBB_PAR_BACKEND=0 -std=c++17 -sycl-std=2020 -fsycl -fsycl-unnamed-lambda -fsycl-device-code-split=per_kernel"
         if test "$hypre_using_debug" = "yes"
         then
-           SYCLFLAGS="-Wall -g ${SYCLFLAGS}"
+           SYCLFLAGS="-O0 -Wall -g -fno-sycl-early-optimizations ${SYCLFLAGS}"
         else
-           SYCLFLAGS="-g -O3 ${SYCLFLAGS}"
+           SYCLFLAGS="-g -O2 ${SYCLFLAGS}"
         fi
 
         dnl (Ab)Use CUFLAGS to capture SYCL compilation flags