From 525809632ef3f198c182c9350db2f7c5d3afd832 Mon Sep 17 00:00:00 2001 From: Harmen Stoppels Date: Mon, 11 Dec 2023 10:30:14 +0100 Subject: petsc: improve hipsparse compat (#40311) Co-authored-by: Satish Balay --- .../builtin/packages/petsc/hip-5.6.0-for-3.18.diff | 85 ++++++++++++++++++++++ .../packages/petsc/hip-5.7-plus-for-3.18.diff | 20 +++++ var/spack/repos/builtin/packages/petsc/package.py | 37 ++++++---- 3 files changed, 127 insertions(+), 15 deletions(-) create mode 100644 var/spack/repos/builtin/packages/petsc/hip-5.6.0-for-3.18.diff create mode 100644 var/spack/repos/builtin/packages/petsc/hip-5.7-plus-for-3.18.diff (limited to 'var') diff --git a/var/spack/repos/builtin/packages/petsc/hip-5.6.0-for-3.18.diff b/var/spack/repos/builtin/packages/petsc/hip-5.6.0-for-3.18.diff new file mode 100644 index 0000000000..c587e3451d --- /dev/null +++ b/var/spack/repos/builtin/packages/petsc/hip-5.6.0-for-3.18.diff @@ -0,0 +1,85 @@ +commit 9b52b1224039b470f0f450943ce503af1df37b00 +Author: Satish Balay +Date: Fri Oct 6 15:19:34 2023 -0500 + + hip-6.0 fix + +diff --git a/src/mat/impls/aij/seq/seqhipsparse/aijhipsparse.hip.cpp b/src/mat/impls/aij/seq/seqhipsparse/aijhipsparse.hip.cpp +index e6be2076975..0c388c90ca3 100644 +--- a/src/mat/impls/aij/seq/seqhipsparse/aijhipsparse.hip.cpp ++++ b/src/mat/impls/aij/seq/seqhipsparse/aijhipsparse.hip.cpp +@@ -1259,14 +1259,22 @@ static PetscErrorCode MatSolve_SeqAIJHIPSPARSE_ILU0(Mat fact, Vec b, Vec x) + /* Solve L*y = b */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, (void *)barray)); + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_Y, fs->Y)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* L Y = X */ ++ fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_L)); // hipsparseSpSV_solve() secretely uses the external buffer used in hipsparseSpSV_analysis()! ++ #else + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* L Y = X */ + fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_L, fs->spsvBuffer_L)); // hipsparseSpSV_solve() secretely uses the external buffer used in hipsparseSpSV_analysis()! +- ++ #endif + /* Solve U*x = y */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, xarray)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_U, /* U X = Y */ ++ fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_U)); ++ #else ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_U, /* U X = Y */ + fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_U, fs->spsvBuffer_U)); +- ++ #endif + PetscCall(VecHIPRestoreArrayRead(b, &barray)); + PetscCall(VecHIPRestoreArrayWrite(x, &xarray)); + +@@ -1309,14 +1317,22 @@ static PetscErrorCode MatSolveTranspose_SeqAIJHIPSPARSE_ILU0(Mat fact, Vec b, Ve + /* Solve Ut*y = b */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, (void *)barray)); + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_Y, fs->Y)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_U, /* Ut Y = X */ ++ fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Ut)); ++ #else + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_U, /* Ut Y = X */ + fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Ut, fs->spsvBuffer_Ut)); +- ++ #endif + /* Solve Lt*x = y */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, xarray)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* Lt X = Y */ ++ fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Lt)); ++ #else + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* Lt X = Y */ + fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Lt, fs->spsvBuffer_Lt)); +- ++ #endif + PetscCall(VecHIPRestoreArrayRead(b, &barray)); + PetscCall(VecHIPRestoreArrayWrite(x, &xarray)); + PetscCall(PetscLogGpuTimeEnd()); +@@ -1544,14 +1560,22 @@ static PetscErrorCode MatSolve_SeqAIJHIPSPARSE_ICC0(Mat fact, Vec b, Vec x) + /* Solve L*y = b */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, (void *)barray)); + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_Y, fs->Y)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* L Y = X */ ++ fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_L)); ++ #else + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_NON_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* L Y = X */ + fs->dnVecDescr_X, fs->dnVecDescr_Y, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_L, fs->spsvBuffer_L)); +- ++ #endif + /* Solve Lt*x = y */ + PetscCallHIPSPARSE(hipsparseDnVecSetValues(fs->dnVecDescr_X, xarray)); ++ #if PETSC_PKG_HIP_VERSION_EQ(5, 6, 0) ++ PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* Lt X = Y */ ++ fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Lt)); ++ #else + PetscCallHIPSPARSE(hipsparseSpSV_solve(fs->handle, HIPSPARSE_OPERATION_TRANSPOSE, &PETSC_HIPSPARSE_ONE, fs->spMatDescr_L, /* Lt X = Y */ + fs->dnVecDescr_Y, fs->dnVecDescr_X, hipsparse_scalartype, HIPSPARSE_SPSV_ALG_DEFAULT, fs->spsvDescr_Lt, fs->spsvBuffer_Lt)); +- ++ #endif + PetscCall(VecHIPRestoreArrayRead(b, &barray)); + PetscCall(VecHIPRestoreArrayWrite(x, &xarray)); + diff --git a/var/spack/repos/builtin/packages/petsc/hip-5.7-plus-for-3.18.diff b/var/spack/repos/builtin/packages/petsc/hip-5.7-plus-for-3.18.diff new file mode 100644 index 0000000000..22555963b3 --- /dev/null +++ b/var/spack/repos/builtin/packages/petsc/hip-5.7-plus-for-3.18.diff @@ -0,0 +1,20 @@ +diff --git a/src/vec/is/sf/impls/basic/hip/sfhip.hip.cpp b/src/vec/is/sf/impls/basic/hip/sfhip.hip.cpp +index a39933c6893..6ef9f513bd6 100644 +--- a/src/vec/is/sf/impls/basic/hip/sfhip.hip.cpp ++++ b/src/vec/is/sf/impls/basic/hip/sfhip.hip.cpp +@@ -471,6 +471,7 @@ __device__ static float atomicMax(float *address, float val) + #endif + + /* As of ROCm 3.10 llint atomicMin/Max(llint*, llint) is not supported */ ++#if PETSC_PKG_HIP_VERSION_LT(5, 7, 0) + __device__ static llint atomicMin(llint *address, llint val) + { + ullint *address_as_ull = (ullint *)(address); +@@ -492,6 +493,7 @@ __device__ static llint atomicMax(llint *address, llint val) + } while (assumed != old); + return (llint)old; + } ++#endif + + template + struct AtomicMin { diff --git a/var/spack/repos/builtin/packages/petsc/package.py b/var/spack/repos/builtin/packages/petsc/package.py index 403e2c2cb4..51d65f229c 100644 --- a/var/spack/repos/builtin/packages/petsc/package.py +++ b/var/spack/repos/builtin/packages/petsc/package.py @@ -161,12 +161,17 @@ class Petsc(Package, CudaPackage, ROCmPackage): variant("kokkos", default=False, description="Activates support for kokkos and kokkos-kernels") variant("fortran", default=True, description="Activates fortran support") - # https://github.com/spack/spack/issues/37416 - conflicts("^rocprim@5.3.0:5.3.2", when="+rocm") - # petsc 3.20 has workaround for breaking change in hipsparseSpSV_solve api, - # but it seems to misdetect hipsparse@5.6.1 as 5.6.0, so the workaround - # only makes things worse - conflicts("^hipsparse@5.6", when="+rocm @3.20.0") + with when("+rocm"): + # https://github.com/spack/spack/issues/37416 + conflicts("^rocprim@5.3.0:5.3.2") + # hipsparse@5.6.0 broke hipsparseSpSV_solve() API, reverted in 5.6.1. + patch( + "https://gitlab.com/petsc/petsc/-/commit/ef7140cce45367033b48bbd2624dfd2b6aa4b997.diff", + when="@3.20.0", + sha256="ba327f8b2a0fa45209dfb7a4278f3e9a323965b5a668be204c1c77c17a963a7f", + ) + patch("hip-5.6.0-for-3.18.diff", when="@3.18:3.19 ^hipsparse@5.6.0") + patch("hip-5.7-plus-for-3.18.diff", when="@3.18:3.19 ^hipsparse@5.7:") # 3.8.0 has a build issue with MKL - so list this conflict explicitly conflicts("^intel-mkl", when="@3.8.0") @@ -225,15 +230,17 @@ class Petsc(Package, CudaPackage, ROCmPackage): depends_on("mpi", when="+mpi") depends_on("cuda", when="+cuda") depends_on("hip", when="+rocm") - depends_on("hipblas", when="+rocm") - depends_on("hipsparse", when="+rocm") - depends_on("hipsolver", when="+rocm") - depends_on("rocsparse", when="+rocm") - depends_on("rocsolver", when="+rocm") - depends_on("rocblas", when="+rocm") - depends_on("rocrand", when="+rocm") - depends_on("rocthrust", when="+rocm") - depends_on("rocprim", when="+rocm") + + with when("+rocm"): + depends_on("hipblas") + depends_on("hipsparse") + depends_on("hipsolver") + depends_on("rocsparse") + depends_on("rocsolver") + depends_on("rocblas") + depends_on("rocrand") + depends_on("rocthrust") + depends_on("rocprim") # Build dependencies depends_on("python@2.6:2.8,3.4:3.8", when="@:3.13", type="build") -- cgit v1.2.3-60-g2f50