From 53bacf193f4240d62492378c27319f517779bd4d Mon Sep 17 00:00:00 2001 From: LwhJesse <256257451+LwhJesse@users.noreply.github.com> Date: Mon, 18 May 2026 23:13:56 +0800 Subject: [PATCH] Cache CUDA SpMV cuSPARSE resources --- Common/include/linear_algebra/CSysMatrix.hpp | 11 +++ Common/src/linear_algebra/CSysMatrix.cpp | 6 ++ Common/src/linear_algebra/CSysMatrixGPU.cu | 79 +++++++++++++++----- 3 files changed, 76 insertions(+), 20 deletions(-) diff --git a/Common/include/linear_algebra/CSysMatrix.hpp b/Common/include/linear_algebra/CSysMatrix.hpp index ecb26959a06..1b8a7565ce5 100644 --- a/Common/include/linear_algebra/CSysMatrix.hpp +++ b/Common/include/linear_algebra/CSysMatrix.hpp @@ -108,6 +108,16 @@ struct CSysMatrixComms { MPI_QUANTITIES commType = MPI_QUANTITIES::SOLUTION_MATRIX); }; +/*! + * \brief Opaque storage for CUDA/cuSPARSE resources used by the GPU SpMV path. + */ +struct CudaSpMVResources; + +/*! + * \brief Release cached CUDA/cuSPARSE resources used by the GPU SpMV path. + */ +void ReleaseCudaSpMVResources(CudaSpMVResources*& resources); + /*! * \class CSysMatrix * \ingroup SpLinSys @@ -150,6 +160,7 @@ class CSysMatrix { const unsigned long* d_col_ind; /*!< \brief Device Column index for each of the elements in val(). */ bool useCuda = false; /*!< \brief Boolean that indicates whether user has enabled CUDA or not. Mainly used to conditionally free GPU memory in the class destructor. */ + mutable CudaSpMVResources* spmv_resources = nullptr; /*!< \brief Cached cuSPARSE resources for GPU SpMV. */ ScalarType* ILU_matrix; /*!< \brief Entries of the ILU sparse matrix. */ unsigned long nnz_ilu; /*!< \brief Number of possible nonzero entries in the matrix (ILU). */ diff --git a/Common/src/linear_algebra/CSysMatrix.cpp b/Common/src/linear_algebra/CSysMatrix.cpp index 00ef51c2023..30c5d52e014 100644 --- a/Common/src/linear_algebra/CSysMatrix.cpp +++ b/Common/src/linear_algebra/CSysMatrix.cpp @@ -67,6 +67,8 @@ template CSysMatrix::~CSysMatrix() { SU2_ZONE_SCOPED + ReleaseCudaSpMVResources(spmv_resources); + delete[] omp_partitions; MemoryAllocation::aligned_free(ILU_matrix); MemoryAllocation::aligned_free(matrix); @@ -86,6 +88,10 @@ CSysMatrix::~CSysMatrix() { #endif } +#ifndef HAVE_CUDA +void ReleaseCudaSpMVResources(CudaSpMVResources*& resources) { resources = nullptr; } +#endif + template void CSysMatrix::Initialize(unsigned long npoint, unsigned long npointdomain, unsigned short nvar, unsigned short neqn, bool EdgeConnect, CGeometry* geometry, diff --git a/Common/src/linear_algebra/CSysMatrixGPU.cu b/Common/src/linear_algebra/CSysMatrixGPU.cu index a3f1a77ca40..af31b9d8ec1 100644 --- a/Common/src/linear_algebra/CSysMatrixGPU.cu +++ b/Common/src/linear_algebra/CSysMatrixGPU.cu @@ -55,6 +55,38 @@ inline cusparseIndexType_t GetCusparseIndexType() { } } +struct CudaSpMVResources { + cusparseHandle_t handle = nullptr; + cusparseConstSpMatDescr_t mat = nullptr; + size_t buffer_size = 0; + void* buffer = nullptr; +}; + +void ReleaseCudaSpMVResources(CudaSpMVResources*& resources) { + if (resources == nullptr) { + return; + } + + if (resources->buffer != nullptr) { + gpuErrChk(cudaFree(resources->buffer)); + resources->buffer = nullptr; + resources->buffer_size = 0; + } + + if (resources->mat != nullptr) { + cusparseErrChk(cusparseDestroySpMat(resources->mat)); + resources->mat = nullptr; + } + + if (resources->handle != nullptr) { + cusparseErrChk(cusparseDestroy(resources->handle)); + resources->handle = nullptr; + } + + delete resources; + resources = nullptr; +} + template constexpr cudaDataType GetCudaDataType() { if constexpr (std::is_same::value) { @@ -100,40 +132,47 @@ void CSysMatrix::GPUMatrixVectorProduct(const CSysVector const ScalarType alpha = 1.0; const ScalarType beta = 0.0; - cusparseHandle_t handle = nullptr; - cusparseConstSpMatDescr_t matA = nullptr; - cusparseDnVecDescr_t vecX = nullptr; - cusparseDnVecDescr_t vecY = nullptr; + if (spmv_resources == nullptr) { + spmv_resources = new CudaSpMVResources; + + cusparseErrChk(cusparseCreate(&spmv_resources->handle)); - cusparseErrChk(cusparseCreate(&handle)); + cusparseErrChk(cusparseCreateConstBsr(&spmv_resources->mat, brows, bcols, bnnz, blockSize, blockSize, d_row_ptr, + d_col_ind, d_matrix, indexType, indexType, CUSPARSE_INDEX_BASE_ZERO, + valueType, CUSPARSE_ORDER_ROW)); + } - cusparseErrChk(cusparseCreateConstBsr(&matA, brows, bcols, bnnz, blockSize, blockSize, d_row_ptr, d_col_ind, d_matrix, - indexType, indexType, CUSPARSE_INDEX_BASE_ZERO, valueType, CUSPARSE_ORDER_ROW)); + cusparseDnVecDescr_t vecX = nullptr; + cusparseDnVecDescr_t vecY = nullptr; cusparseErrChk(cusparseCreateDnVec(&vecX, xSize, d_vec, valueType)); cusparseErrChk(cusparseCreateDnVec(&vecY, ySize, d_prod, valueType)); - size_t bufferSize = 0; - void* dBuffer = nullptr; + size_t required_buffer_size = 0; - cusparseErrChk(cusparseSpMV_bufferSize(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matA, vecX, &beta, vecY, - valueType, CUSPARSE_SPMV_BSR_ALG1, &bufferSize)); + cusparseErrChk(cusparseSpMV_bufferSize(spmv_resources->handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, + spmv_resources->mat, vecX, &beta, vecY, valueType, CUSPARSE_SPMV_BSR_ALG1, + &required_buffer_size)); - if (bufferSize > 0) { - gpuErrChk(cudaMalloc(&dBuffer, bufferSize)); - } + if (required_buffer_size > spmv_resources->buffer_size) { + if (spmv_resources->buffer != nullptr) { + gpuErrChk(cudaFree(spmv_resources->buffer)); + } - cusparseErrChk(cusparseSpMV(handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, matA, vecX, &beta, vecY, valueType, - CUSPARSE_SPMV_BSR_ALG1, dBuffer)); + if (required_buffer_size > 0) { + gpuErrChk(cudaMalloc(&spmv_resources->buffer, required_buffer_size)); + } else { + spmv_resources->buffer = nullptr; + } - if (dBuffer != nullptr) { - gpuErrChk(cudaFree(dBuffer)); + spmv_resources->buffer_size = required_buffer_size; } + cusparseErrChk(cusparseSpMV(spmv_resources->handle, CUSPARSE_OPERATION_NON_TRANSPOSE, &alpha, spmv_resources->mat, + vecX, &beta, vecY, valueType, CUSPARSE_SPMV_BSR_ALG1, spmv_resources->buffer)); + cusparseErrChk(cusparseDestroyDnVec(vecY)); cusparseErrChk(cusparseDestroyDnVec(vecX)); - cusparseErrChk(cusparseDestroySpMat(matA)); - cusparseErrChk(cusparseDestroy(handle)); prod.DtHTransfer(); }