diff --git a/.github/vcpkg-overlays/libaec/portfile.cmake b/.github/vcpkg-overlays/libaec/portfile.cmake deleted file mode 100644 index e04dabea..00000000 --- a/.github/vcpkg-overlays/libaec/portfile.cmake +++ /dev/null @@ -1,64 +0,0 @@ -# Overlay of the upstream vcpkg `libaec` port, identical to the port at our -# pinned vcpkg baseline except for where the sources come from. -# -# Upstream fetches from https://gitlab.dkrz.de/k202009/libaec. That host now -# answers every request to its /-/archive/ tarball endpoint with a bare -# nginx "429 Too Many Requests" (no Retry-After), so any cold vcpkg build -# fails there and no amount of retrying gets past it. The repository web -# pages still serve 200 — it is specifically the archive endpoint that is -# rate limited. -# -# libaec reaches us only as a transitive dependency: vtk -> hdf5 -> (szip -# feature) -> libaec, so it cannot simply be dropped. -# -# MathisRosenhauer/libaec on GitHub is the upstream author's mirror of the -# same project. The SHA512 below is of GitHub's v1.1.6 tarball, which is a -# different archive from GitLab's and therefore has a different hash than -# upstream's portfile records. -# -# Remove this overlay once gitlab.dkrz.de serves archives again. - -vcpkg_from_github( - OUT_SOURCE_PATH - SOURCE_PATH - REPO - MathisRosenhauer/libaec - REF - "v${VERSION}" - SHA512 - 76df7501d1b7d91a43b525ba828f092f18d83f8ab09a9331e5758f93942a9758ad580baca8f9316b92a98639bde2e23cacbc2f33f52d0dd98ce7efe412cf43cd - HEAD_REF - master) - -string(COMPARE EQUAL "${VCPKG_LIBRARY_LINKAGE}" "static" BUILD_STATIC) - -vcpkg_cmake_configure(SOURCE_PATH "${SOURCE_PATH}" OPTIONS -DBUILD_STATIC_LIBS=${BUILD_STATIC} - -Dlibaec_INSTALL_CMAKEDIR=share/${PORT}) -vcpkg_cmake_install() -vcpkg_copy_pdbs() -vcpkg_cmake_config_fixup() -vcpkg_replace_string("${CURRENT_PACKAGES_DIR}/share/libaec/libaec-config.cmake" - "if(libaec_USE_STATIC_LIBS)" "if(\"${BUILD_STATIC}\") # forced by vcpkg") - -# Compatibility with user's CMake < 3.18 (vcpkg claims support for >= 3.16): -# Make imported targets global so that libaec-config.cmake can create ALIAS targets. -set(_target_file "libaec_shared-targets") -if(BUILD_STATIC) - set(_target_file "libaec_static-targets") -endif() -file(READ "${CURRENT_PACKAGES_DIR}/share/libaec/${_target_file}.cmake" libaec_targets) -string(REGEX REPLACE " (SHARED|STATIC) IMPORTED" " \\1 IMPORTED \${libaec_maybe_global}" - libaec_targets "${libaec_targets}") -file( - WRITE "${CURRENT_PACKAGES_DIR}/share/libaec/${_target_file}.cmake" - "set(libaec_maybe_global \"\") -if(CMAKE_VERSION VERSION_LESS 3.18) - set(libaec_maybe_global \"GLOBAL\") -endif() -${libaec_targets} -") - -file(REMOVE_RECURSE "${CURRENT_PACKAGES_DIR}/debug/include") - -file(INSTALL "${CURRENT_PORT_DIR}/usage" DESTINATION "${CURRENT_PACKAGES_DIR}/share/${PORT}") -vcpkg_install_copyright(FILE_LIST "${SOURCE_PATH}/LICENSE.txt") diff --git a/.github/vcpkg-overlays/libaec/usage b/.github/vcpkg-overlays/libaec/usage deleted file mode 100644 index a6aeb5f6..00000000 --- a/.github/vcpkg-overlays/libaec/usage +++ /dev/null @@ -1,7 +0,0 @@ -libaec provides CMake targets: - - find_package(libaec CONFIG REQUIRED) - # libaec API - target_link_libraries(main PRIVATE libaec::aec) - # szip compatible API - target_link_libraries(main PRIVATE libaec::sz) diff --git a/.github/vcpkg-overlays/libaec/vcpkg.json b/.github/vcpkg-overlays/libaec/vcpkg.json deleted file mode 100644 index 299cd632..00000000 --- a/.github/vcpkg-overlays/libaec/vcpkg.json +++ /dev/null @@ -1,17 +0,0 @@ -{ - "name": "libaec", - "version": "1.1.6", - "description": "Adaptive Entropy Coding library", - "homepage": "https://github.com/MathisRosenhauer/libaec", - "license": "BSD-2-Clause", - "dependencies": [ - { - "name": "vcpkg-cmake", - "host": true - }, - { - "name": "vcpkg-cmake-config", - "host": true - } - ] -} diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 6ffad122..5d24b799 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -78,11 +78,7 @@ jobs: if: ${{ matrix.os == 'windows-latest' && steps.vcpkg-installed-cache.outputs.cache-hit != 'true' }} shell: pwsh run: | - # The libaec overlay sources that port from GitHub instead of - # gitlab.dkrz.de, whose archive endpoint returns 429 to everyone. - # See .github/vcpkg-overlays/libaec/portfile.cmake. - ./vcpkg/vcpkg install --triplet x64-windows ` - --overlay-ports="${{ github.workspace }}/.github/vcpkg-overlays" + ./vcpkg/vcpkg install --triplet x64-windows - name: 🏗️ Compile (Windows) if: ${{ matrix.os == 'windows-latest' }} diff --git a/.github/workflows/python.yml b/.github/workflows/python.yml index 6ad060dc..91b66aec 100644 --- a/.github/workflows/python.yml +++ b/.github/workflows/python.yml @@ -75,11 +75,7 @@ jobs: if: ${{ matrix.os == 'windows-latest' && steps.vcpkg-installed-cache.outputs.cache-hit != 'true' }} shell: pwsh run: | - # The libaec overlay sources that port from GitHub instead of - # gitlab.dkrz.de, whose archive endpoint returns 429 to everyone. - # See .github/vcpkg-overlays/libaec/portfile.cmake. - ./vcpkg/vcpkg install --triplet x64-windows ` - --overlay-ports="${{ github.workspace }}/.github/vcpkg-overlays" + ./vcpkg/vcpkg install --triplet x64-windows - name: 🐍 Build and check Python Module (Windows) if: ${{ matrix.os == 'windows-latest' }} @@ -157,21 +153,30 @@ jobs: env: CIBW_SKIP: ${{ matrix.skip }} CIBW_MANYLINUX_X86_64_IMAGE: ghcr.io/viennatools/vienna-builder:cuda-python + # The CUDA libraries are far too large to vendor (libcusparse alone is + # >160 MB, over the PyPI file limit), so they are left unbundled. + # That is safe here only because _core.so does not link them: they are + # needed solely by libViennaLS_GPU.so, which ViennaLS dlopen()s on + # demand and does without when it cannot be loaded. Keep it that way — + # linking CUDA into _core.so is what broke the 5.8.3 wheel. + CIBW_REPAIR_WHEEL_COMMAND_LINUX: >- + auditwheel repair + --exclude libcuda.so.1 + --exclude libcudart.so.12 + --exclude libcusparse.so.12 + --exclude libnvJitLink.so.12 + -w {dest_dir} {wheel} CIBW_TEST_COMMAND: >- python -c "import viennals; print(viennals.__file__); print(viennals.d2); print(viennals.d3)" - # VIENNALS_USE_GPU is deliberately OFF for wheels: it puts hard - # DT_NEEDED entries for libcudart/libcusparse on _core.so, and those - # are far too large to vendor (libcusparse alone is >160 MB, over the - # PyPI file limit). The wheel would then fail to import on any machine - # without a CUDA toolkit. GPU users build from source with - # -DVIENNALS_USE_GPU=ON. CIBW_CONFIG_SETTINGS: >- cmake.define.VIENNALS_PACKAGE_PYTHON=ON cmake.define.VIENNALS_IS_CI=ON cmake.define.VIENNALS_VTK_RENDERING=ON + cmake.define.VIENNALS_USE_GPU=ON + cmake.define.VIENNACORE_LINK_CUDA_DRIVER=OFF cmake.define.USE_IPO=OFF - name: 🏗️ Build Wheels (Other) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7c3ba02b..00e3d9d7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.20 FATAL_ERROR) project( ViennaLS LANGUAGES CXX - VERSION 5.8.4) + VERSION 5.8.5) # -------------------------------------------------------------------------------------------------------- # Library options @@ -158,9 +158,16 @@ if(VIENNALS_USE_GPU) "[ViennaLS] cuSPARSE not found in CUDA ${CUDAToolkit_VERSION}; disabling GPU BiCGSTAB solver." ) else() - # Compile the CUDA kernels into a small static library so that consuming - # .cpp files (compiled by g++) never see __global__ / <<< >>> syntax. - add_library(ViennaLS_GPU STATIC ${PROJECT_SOURCE_DIR}/gpu/lsOxidationBiCGSTABKernels.cu) + # Compile the CUDA kernels into a SHARED library so that consuming .cpp + # files (compiled by g++) never see __global__ / <<< >>> syntax, and — just + # as importantly — so that this is the *only* binary linking CUDA. + # + # It is deliberately not linked into ViennaLS. ViennaLS opens it lazily with + # dlopen() at first use (see lsOxidationBiCGSTABInterface.hpp). Linking it + # would put DT_NEEDED entries for libcudart/libcusparse onto every consumer, + # and the dynamic loader would then refuse to load ViennaLS at all on a + # machine with no CUDA runtime — which is what broke the 5.8.3 PyPI wheel. + add_library(ViennaLS_GPU SHARED ${PROJECT_SOURCE_DIR}/gpu/lsOxidationBiCGSTABKernels.cu) target_compile_definitions(ViennaLS_GPU PRIVATE VIENNALS_GPU_BICGSTAB=1) @@ -170,28 +177,37 @@ if(VIENNALS_USE_GPU) set_target_properties(ViennaLS_GPU PROPERTIES CUDA_SEPARABLE_COMPILATION OFF POSITION_INDEPENDENT_CODE ON) - target_link_libraries(ViennaLS_GPU PUBLIC CUDA::cudart CUDA::cusparse) + target_link_libraries(ViennaLS_GPU PRIVATE CUDA::cudart CUDA::cusparse) - if(UNIX AND NOT APPLE) + # A locally built library has to find the CUDA runtime, which often lives + # outside the default loader search path (e.g. /usr/local/cuda/lib64). + # Wheel builds skip this: auditwheel rewrites RPATHs, and a build-machine + # path baked into a redistributable wheel would be meaningless anyway. + if(UNIX + AND NOT APPLE + AND NOT VIENNALS_PACKAGE_PYTHON) + set(_VIENNALS_CUDA_RPATHS) foreach(_VIENNALS_CUDA_TARGET CUDA::cudart CUDA::cusparse) get_target_property(_VIENNALS_CUDA_LIBRARY ${_VIENNALS_CUDA_TARGET} IMPORTED_LOCATION) if(_VIENNALS_CUDA_LIBRARY) get_filename_component(_VIENNALS_CUDA_LIBRARY_DIR "${_VIENNALS_CUDA_LIBRARY}" DIRECTORY) - target_link_options(ViennaLS_GPU INTERFACE - "$") + list(APPEND _VIENNALS_CUDA_RPATHS "${_VIENNALS_CUDA_LIBRARY_DIR}") endif() endforeach() + list(REMOVE_DUPLICATES _VIENNALS_CUDA_RPATHS) + set_target_properties(ViennaLS_GPU PROPERTIES BUILD_RPATH "${_VIENNALS_CUDA_RPATHS}" + INSTALL_RPATH "${_VIENNALS_CUDA_RPATHS}") endif() - # Expose the define and the GPU library to downstream targets. - # BUILD_INTERFACE limits the ViennaLS_GPU dependency to the build tree so it - # is never exported to the install tree (avoiding the "not in export set" error - # from packageProject when ViennaLS_GPU is a non-installed local target). + # Expose only the define to downstream targets — never the GPU library and + # never the CUDA include dirs. The dispatch layer needs nothing from the + # CUDA headers; it talks to ViennaLS_GPU through the plain C ABI in + # lsOxidationBiCGSTABAbi.hpp. Consumers do need the dl functions. target_compile_definitions(${PROJECT_NAME} INTERFACE VIENNALS_GPU_BICGSTAB=1) - target_link_libraries(${PROJECT_NAME} INTERFACE $) - target_include_directories(${PROJECT_NAME} INTERFACE ${CUDAToolkit_INCLUDE_DIRS}) + target_link_libraries(${PROJECT_NAME} INTERFACE ${CMAKE_DL_LIBS}) - message(STATUS "[ViennaLS] GPU BiCGSTAB solver enabled (CUDA ${CUDAToolkit_VERSION})") + message(STATUS "[ViennaLS] GPU BiCGSTAB solver enabled (CUDA ${CUDAToolkit_VERSION}, " + "loaded at runtime)") endif() endif() diff --git a/README.md b/README.md index b6517412..63197824 100644 --- a/README.md +++ b/README.md @@ -175,7 +175,7 @@ We recommend using [CPM.cmake](https://github.com/cpm-cmake/CPM.cmake) to consum * Installation with CPM ```cmake - CPMAddPackage("gh:viennatools/viennals@5.8.4") + CPMAddPackage("gh:viennatools/viennals@5.8.5") ``` * With a local installation diff --git a/gpu/lsOxidationBiCGSTABKernels.cu b/gpu/lsOxidationBiCGSTABKernels.cu index 99eb8778..033633ce 100644 --- a/gpu/lsOxidationBiCGSTABKernels.cu +++ b/gpu/lsOxidationBiCGSTABKernels.cu @@ -1,199 +1,213 @@ -// Compiled by nvcc. Provides C++ wrapper functions that delegate to the -// CUDA kernel implementations in lsOxidationBiCGSTAB.cuh. +// Compiled by nvcc into the ViennaLS_GPU *shared* library. Provides the +// stable C ABI declared in lsOxidationBiCGSTABAbi.hpp, delegating to the CUDA +// kernel implementations in lsOxidationBiCGSTAB.cuh. // // This translation unit is the ONLY place where lsOxidationBiCGSTAB.cuh // is included; no .hpp/.cpp file compiled by g++ ever sees the __global__ // kernel syntax. +// +// This library is the only binary that links against the CUDA runtime. +// ViennaLS opens it lazily with dlopen() (see lsOxidationBiCGSTABInterface.hpp) +// so that a GPU-enabled build still loads where no CUDA runtime is installed. +// Do NOT include lsOxidationBiCGSTABInterface.hpp here: that header defines +// inline dispatch functions for the *caller* side. #include -#include +#include -#include #include +#include #include #include -namespace viennals { -namespace gpu { - namespace { + +using viennals::gpu::GpuBiCGSTABBuffers; + thread_local std::string lastGpuError; -void setLastGpuError(const std::string& message) { - lastGpuError = message; - fprintf(stderr, "GpuBiCGSTABBuffers: %s\n", message.c_str()); +void setLastGpuError(const std::string &message) { + lastGpuError = message; + fprintf(stderr, "GpuBiCGSTABBuffers: %s\n", message.c_str()); } -void clearLastGpuError() { - lastGpuError.clear(); -} -} // namespace +void clearLastGpuError() { lastGpuError.clear(); } -GpuBiCGSTABBuffers* allocGpuBuffers(uint32_t n, int nFaces, - bool useIlu0Preconditioner) { - clearLastGpuError(); - // Bail out if no device exists or if a CUDA context cannot actually be - // created (driver mismatch, permissions, etc.). cudaGetDeviceCount only - // enumerates devices; cudaFree(nullptr) forces real context initialization. - int deviceCount = 0; - const cudaError_t countErr = cudaGetDeviceCount(&deviceCount); - if (countErr != cudaSuccess) { - setLastGpuError(std::string("cudaGetDeviceCount failed: ") + - cudaGetErrorString(countErr)); - cudaGetLastError(); - return nullptr; - } - if (deviceCount == 0) { - setLastGpuError("cudaGetDeviceCount reported zero CUDA devices"); - return nullptr; - } - int deviceId = 0; - if (cudaGetDevice(&deviceId) != cudaSuccess || deviceId < 0 || - deviceId >= deviceCount) { - cudaGetLastError(); - deviceId = 0; - } - const cudaError_t setErr = cudaSetDevice(deviceId); - if (setErr != cudaSuccess) { - setLastGpuError(std::string("cudaSetDevice(") + - std::to_string(deviceId) + ") failed: " + - cudaGetErrorString(setErr)); - cudaGetLastError(); - return nullptr; - } - const cudaError_t ctxErr = cudaFree(nullptr); - if (ctxErr != cudaSuccess) { - setLastGpuError(std::string("CUDA context initialization failed on " - "device ") + - std::to_string(deviceId) + ": " + - cudaGetErrorString(ctxErr)); - cudaGetLastError(); // consume so later CUDA calls start clean - return nullptr; - } - - auto* b = new (std::nothrow) GpuBiCGSTABBuffers(); - if (!b) return nullptr; - b->deviceId = deviceId; - b->allocate(n, nFaces); - b->useIlu0Preconditioner = useIlu0Preconditioner; - if (!b->valid) { - delete b; - return nullptr; - } - return b; +GpuBiCGSTABBuffers *asBuffers(void *handle) { + return static_cast(handle); } -void freeGpuBuffers(GpuBiCGSTABBuffers* gpu) { - delete gpu; // calls ~GpuBiCGSTABBuffers() which calls gpu->free() +const GpuBiCGSTABBuffers *asBuffers(const void *handle) { + return static_cast(handle); } -bool gpuIsValid(const GpuBiCGSTABBuffers* gpu) { - return gpu != nullptr && gpu->valid; +} // namespace + +extern "C" { + +int viennalsGpuAbiVersion(void) { return VIENNALS_GPU_ABI_VERSION; } + +void *viennalsGpuAllocBuffers(uint32_t n, int nFaces, + int useIlu0Preconditioner) { + clearLastGpuError(); + // Bail out if no device exists or if a CUDA context cannot actually be + // created (driver mismatch, permissions, etc.). cudaGetDeviceCount only + // enumerates devices; cudaFree(nullptr) forces real context initialization. + int deviceCount = 0; + const cudaError_t countErr = cudaGetDeviceCount(&deviceCount); + if (countErr != cudaSuccess) { + setLastGpuError(std::string("cudaGetDeviceCount failed: ") + + cudaGetErrorString(countErr)); + cudaGetLastError(); + return nullptr; + } + if (deviceCount == 0) { + setLastGpuError("cudaGetDeviceCount reported zero CUDA devices"); + return nullptr; + } + int deviceId = 0; + if (cudaGetDevice(&deviceId) != cudaSuccess || deviceId < 0 || + deviceId >= deviceCount) { + cudaGetLastError(); + deviceId = 0; + } + const cudaError_t setErr = cudaSetDevice(deviceId); + if (setErr != cudaSuccess) { + setLastGpuError(std::string("cudaSetDevice(") + std::to_string(deviceId) + + ") failed: " + cudaGetErrorString(setErr)); + cudaGetLastError(); + return nullptr; + } + const cudaError_t ctxErr = cudaFree(nullptr); + if (ctxErr != cudaSuccess) { + setLastGpuError(std::string("CUDA context initialization failed on " + "device ") + + std::to_string(deviceId) + ": " + + cudaGetErrorString(ctxErr)); + cudaGetLastError(); // consume so later CUDA calls start clean + return nullptr; + } + + auto *b = new (std::nothrow) GpuBiCGSTABBuffers(); + if (!b) + return nullptr; + b->deviceId = deviceId; + b->allocate(n, nFaces); + b->useIlu0Preconditioner = useIlu0Preconditioner != 0; + if (!b->valid) { + delete b; + return nullptr; + } + return b; } -const char* gpuGetLastErrorMessage() { - return lastGpuError.c_str(); +void viennalsGpuFreeBuffers(void *handle) { + delete asBuffers(handle); // calls ~GpuBiCGSTABBuffers() -> gpu->free() } -bool gpuUploadNeighborIds(GpuBiCGSTABBuffers* gpu, - const uint32_t* nb, - std::size_t count) { - const std::size_t expected = - static_cast(gpu->n) * static_cast(gpu->nFaces); - if (count != expected) { - setLastGpuError(std::string("neighbor ID upload length mismatch (got ") + - std::to_string(count) + ", expected " + - std::to_string(expected) + ")"); - return false; - } - if (!gpu->activateDevice("gpuUploadNeighborIds")) - return false; - return gpu->checkCuda( - cudaMemcpy(gpu->d_nb, nb, count * sizeof(uint32_t), - cudaMemcpyHostToDevice), - "cudaMemcpy(d_nb)"); +int viennalsGpuIsValid(const void *handle) { + const auto *gpu = asBuffers(handle); + return (gpu != nullptr && gpu->valid) ? 1 : 0; } -bool gpuUploadSolverArrays(GpuBiCGSTABBuffers* gpu, - const double* diag, - const double* b, - const double* coeff, - uint32_t diagLen, - std::size_t coeffLen) { - const std::size_t expectedCoeff = - static_cast(gpu->n) * static_cast(gpu->nFaces); - if (diagLen != gpu->n || coeffLen != expectedCoeff) { - setLastGpuError(std::string("solver array upload length mismatch " - "(diag got ") + - std::to_string(diagLen) + ", expected " + - std::to_string(gpu->n) + "; coeff got " + - std::to_string(coeffLen) + ", expected " + - std::to_string(expectedCoeff) + ")"); - return false; - } - if (!gpu->activateDevice("gpuUploadSolverArrays")) - return false; - if (!gpu->checkCuda(cudaMemcpy(gpu->d_diag, diag, diagLen * sizeof(double), - cudaMemcpyHostToDevice), - "cudaMemcpy(d_diag)")) - return false; - if (!gpu->checkCuda(cudaMemcpy(gpu->d_b, b, diagLen * sizeof(double), - cudaMemcpyHostToDevice), - "cudaMemcpy(d_b)")) - return false; - if (!gpu->checkCuda(cudaMemcpy(gpu->d_coeff, coeff, - coeffLen * sizeof(double), +const char *viennalsGpuGetLastErrorMessage(void) { return lastGpuError.c_str(); } + +int viennalsGpuUploadNeighborIds(void *handle, const uint32_t *nb, + std::size_t count) { + auto *gpu = asBuffers(handle); + const std::size_t expected = + static_cast(gpu->n) * static_cast(gpu->nFaces); + if (count != expected) { + setLastGpuError(std::string("neighbor ID upload length mismatch (got ") + + std::to_string(count) + ", expected " + + std::to_string(expected) + ")"); + return 0; + } + if (!gpu->activateDevice("gpuUploadNeighborIds")) + return 0; + return gpu->checkCuda(cudaMemcpy(gpu->d_nb, nb, count * sizeof(uint32_t), cudaMemcpyHostToDevice), - "cudaMemcpy(d_coeff)")) - return false; - if (!gpu->useIlu0Preconditioner) - return true; - // Refresh ILU(0) values and re-factorize whenever the coefficients change. - if (gpu->iluReady) - return gpu->fillAndFactorizeILU(diag, coeff, gpu->nFaces); - return false; + "cudaMemcpy(d_nb)") + ? 1 + : 0; } -bool gpuUploadRhs(GpuBiCGSTABBuffers* gpu, const double* b, uint32_t n) { - if (n != gpu->n) { - setLastGpuError(std::string("RHS upload length mismatch (got ") + - std::to_string(n) + ", expected " + - std::to_string(gpu->n) + ")"); - return false; - } - if (!gpu->activateDevice("gpuUploadRhs")) - return false; - return gpu->checkCuda( - cudaMemcpy(gpu->d_b, b, static_cast(n) * sizeof(double), - cudaMemcpyHostToDevice), - "cudaMemcpy(d_b)"); +int viennalsGpuUploadSolverArrays(void *handle, const double *diag, + const double *b, const double *coeff, + uint32_t diagLen, std::size_t coeffLen) { + auto *gpu = asBuffers(handle); + const std::size_t expectedCoeff = + static_cast(gpu->n) * static_cast(gpu->nFaces); + if (diagLen != gpu->n || coeffLen != expectedCoeff) { + setLastGpuError(std::string("solver array upload length mismatch " + "(diag got ") + + std::to_string(diagLen) + ", expected " + + std::to_string(gpu->n) + "; coeff got " + + std::to_string(coeffLen) + ", expected " + + std::to_string(expectedCoeff) + ")"); + return 0; + } + if (!gpu->activateDevice("gpuUploadSolverArrays")) + return 0; + if (!gpu->checkCuda(cudaMemcpy(gpu->d_diag, diag, diagLen * sizeof(double), + cudaMemcpyHostToDevice), + "cudaMemcpy(d_diag)")) + return 0; + if (!gpu->checkCuda(cudaMemcpy(gpu->d_b, b, diagLen * sizeof(double), + cudaMemcpyHostToDevice), + "cudaMemcpy(d_b)")) + return 0; + if (!gpu->checkCuda(cudaMemcpy(gpu->d_coeff, coeff, + coeffLen * sizeof(double), + cudaMemcpyHostToDevice), + "cudaMemcpy(d_coeff)")) + return 0; + if (!gpu->useIlu0Preconditioner) + return 1; + // Refresh ILU(0) values and re-factorize whenever the coefficients change. + if (gpu->iluReady) + return gpu->fillAndFactorizeILU(diag, coeff, gpu->nFaces) ? 1 : 0; + return 0; +} + +int viennalsGpuUploadRhs(void *handle, const double *b, uint32_t n) { + auto *gpu = asBuffers(handle); + if (n != gpu->n) { + setLastGpuError(std::string("RHS upload length mismatch (got ") + + std::to_string(n) + ", expected " + + std::to_string(gpu->n) + ")"); + return 0; + } + if (!gpu->activateDevice("gpuUploadRhs")) + return 0; + return gpu->checkCuda(cudaMemcpy(gpu->d_b, b, + static_cast(n) * sizeof(double), + cudaMemcpyHostToDevice), + "cudaMemcpy(d_b)") + ? 1 + : 0; } -bool gpuSetupCSR(GpuBiCGSTABBuffers* gpu, - const uint32_t* h_nb, - uint32_t n, - int nFaces) { - if (!gpu->activateDevice("gpuSetupCSR")) - return false; - return gpu->setupCSR(h_nb, n, nFaces); +int viennalsGpuSetupCSR(void *handle, const uint32_t *hNb, uint32_t n, + int nFaces) { + auto *gpu = asBuffers(handle); + if (!gpu->activateDevice("gpuSetupCSR")) + return 0; + return gpu->setupCSR(hNb, n, nFaces) ? 1 : 0; } -bool gpuSolveBiCGSTAB(GpuBiCGSTABBuffers* gpu, - double* x, - double diagEps, - unsigned maxIter, - double tolerance, - unsigned& outIterations, - double& outResidual) { - if (!gpu->activateDevice("gpuSolveBiCGSTAB")) - return false; - std::vector xVec(x, x + gpu->n); - const bool ok = solveBiCGSTAB(*gpu, xVec, diagEps, maxIter, tolerance, - outIterations, outResidual); - if (ok) - std::copy(xVec.begin(), xVec.end(), x); - return ok; +int viennalsGpuSolveBiCGSTAB(void *handle, double *x, double diagEps, + unsigned maxIter, double tolerance, + unsigned *outIterations, double *outResidual) { + auto *gpu = asBuffers(handle); + if (!gpu->activateDevice("gpuSolveBiCGSTAB")) + return 0; + std::vector xVec(x, x + gpu->n); + const bool ok = viennals::gpu::solveBiCGSTAB( + *gpu, xVec, diagEps, maxIter, tolerance, *outIterations, *outResidual); + if (ok) + std::copy(xVec.begin(), xVec.end(), x); + return ok ? 1 : 0; } -} // namespace gpu -} // namespace viennals +} // extern "C" diff --git a/include/viennals/lsOxidationBiCGSTABAbi.hpp b/include/viennals/lsOxidationBiCGSTABAbi.hpp new file mode 100644 index 00000000..dbd37ab7 --- /dev/null +++ b/include/viennals/lsOxidationBiCGSTABAbi.hpp @@ -0,0 +1,78 @@ +// Stable C ABI exported by the ViennaLS_GPU shared library. +// +// The GPU BiCGSTAB solver lives in a separate shared library that links +// against the CUDA runtime. ViennaLS itself must NOT link against CUDA, so +// that a build with GPU support still loads on machines with no CUDA runtime +// installed. The library is therefore opened lazily with dlopen() and these +// symbols are resolved with dlsym() -- see lsOxidationBiCGSTABInterface.hpp. +// +// Everything here is deliberately plain C: opaque void* handles, no +// references, int instead of bool. That keeps the boundary independent of +// C++ name mangling and of the standard library, so the two sides stay +// compatible even when built by different compilers (nvcc vs g++). +// +// Bump VIENNALS_GPU_ABI_VERSION on any incompatible change to the signatures +// below; the loader refuses a library reporting a different version. + +#pragma once + +#include +#include + +#define VIENNALS_GPU_ABI_VERSION 1 + +// The exporting side (lsOxidationBiCGSTABKernels.cu) must keep these visible +// even under -fvisibility=hidden, or dlsym() cannot find them. +#if defined(_WIN32) +#define VIENNALS_GPU_ABI __declspec(dllexport) +#else +#define VIENNALS_GPU_ABI __attribute__((visibility("default"))) +#endif + +extern "C" { + +/// Returns the ABI version this library was built against. +VIENNALS_GPU_ABI int viennalsGpuAbiVersion(void); + +/// Allocate GPU buffers for `n` nodes and `nFaces` (2*D) faces. +/// Returns nullptr if no usable CUDA device or context is available. +VIENNALS_GPU_ABI void *viennalsGpuAllocBuffers(uint32_t n, int nFaces, + int useIlu0Preconditioner); + +/// Free previously allocated buffers. Safe to call with nullptr. +VIENNALS_GPU_ABI void viennalsGpuFreeBuffers(void *handle); + +/// Non-zero if the handle is valid (non-null and successfully allocated). +VIENNALS_GPU_ABI int viennalsGpuIsValid(const void *handle); + +/// Human-readable detail for the last failure on this thread. +VIENNALS_GPU_ABI const char *viennalsGpuGetLastErrorMessage(void); + +/// Upload the geometry-fixed neighbor-ID array; `count` must equal nFaces*n. +VIENNALS_GPU_ABI int viennalsGpuUploadNeighborIds(void *handle, + const uint32_t *nb, + std::size_t count); + +/// Build the CSR sparsity pattern and run the cuSPARSE symbolic analysis. +VIENNALS_GPU_ABI int viennalsGpuSetupCSR(void *handle, const uint32_t *hNb, + uint32_t n, int nFaces); + +/// Upload per-solve arrays (diag, b, faceCoeffs) and re-factorize ILU(0). +VIENNALS_GPU_ABI int +viennalsGpuUploadSolverArrays(void *handle, const double *diag, const double *b, + const double *coeff, uint32_t diagLen, + std::size_t coeffLen); + +/// Upload only the RHS vector, when the matrix geometry is already resident. +VIENNALS_GPU_ABI int viennalsGpuUploadRhs(void *handle, const double *b, + uint32_t n); + +/// Run GPU BiCGSTAB. `x` is the initial guess on entry, the solution on exit. +/// Non-zero only when the solve converged and produced finite values. +VIENNALS_GPU_ABI int viennalsGpuSolveBiCGSTAB(void *handle, double *x, + double diagEps, unsigned maxIter, + double tolerance, + unsigned *outIterations, + double *outResidual); + +} // extern "C" diff --git a/include/viennals/lsOxidationBiCGSTABInterface.hpp b/include/viennals/lsOxidationBiCGSTABInterface.hpp index 7c892196..6dd71f92 100644 --- a/include/viennals/lsOxidationBiCGSTABInterface.hpp +++ b/include/viennals/lsOxidationBiCGSTABInterface.hpp @@ -7,13 +7,36 @@ // // The actual CUDA kernels live in lsOxidationBiCGSTAB.cuh. That file // must never be included from a .cpp compiled by g++ — only from .cu files. +// +// ── Why dlopen instead of linking ────────────────────────────────────────── +// ViennaLS must not carry a hard (DT_NEEDED) dependency on libcudart / +// libcusparse. If it did, the dynamic loader would fail to load ViennaLS at +// all on a machine without a CUDA runtime — before a single line of our code +// runs, so no amount of runtime checking could recover. That is exactly what +// broke the 5.8.3 PyPI wheel. +// +// Instead, the CUDA-linked code lives in a separate ViennaLS_GPU shared +// library which is opened lazily the first time the GPU path is requested. +// If it (or the CUDA runtime it needs) cannot be loaded, allocGpuBuffers() +// returns nullptr and callers transparently fall back to the CPU solver, with +// the reason available from gpuGetLastErrorMessage(). #pragma once #ifdef VIENNALS_GPU_BICGSTAB +#include "lsOxidationBiCGSTABAbi.hpp" + #include #include +#include +#include + +#if defined(_WIN32) +#include +#else +#include +#endif namespace viennals { namespace gpu { @@ -26,50 +49,270 @@ static constexpr uint32_t kNoNode = 0xFFFFFFFFu; // lsOxidationBiCGSTABKernels.cu. Consumers hold a raw pointer only. struct GpuBiCGSTABBuffers; +namespace detail { + +#if defined(_WIN32) +using LibraryHandle = HMODULE; +inline LibraryHandle openLibrary(const char *path) { + return LoadLibraryA(path); +} +inline void *findSymbol(LibraryHandle lib, const char *name) { + return reinterpret_cast(GetProcAddress(lib, name)); +} +inline std::string lastLoadError() { + return "LoadLibrary failed (error " + std::to_string(GetLastError()) + ")"; +} +inline const char *gpuLibraryName() { return "ViennaLS_GPU.dll"; } +inline char pathSeparator() { return '\\'; } +#else +using LibraryHandle = void *; +inline LibraryHandle openLibrary(const char *path) { + return dlopen(path, RTLD_NOW | RTLD_LOCAL); +} +inline void *findSymbol(LibraryHandle lib, const char *name) { + return dlsym(lib, name); +} +inline std::string lastLoadError() { + const char *err = dlerror(); + return err ? std::string(err) : std::string("unknown dynamic loader error"); +} +inline const char *gpuLibraryName() { return "libViennaLS_GPU.so"; } +inline char pathSeparator() { return '/'; } +#endif + +/// Directory containing the binary this code is compiled into (e.g. the +/// Python extension module). The GPU library is shipped alongside it. +inline std::string currentModuleDirectory() { + std::string path; +#if defined(_WIN32) + HMODULE module = nullptr; + if (GetModuleHandleExA(GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | + GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT, + reinterpret_cast(¤tModuleDirectory), + &module) && + module) { + char buffer[MAX_PATH] = {}; + if (GetModuleFileNameA(module, buffer, MAX_PATH)) + path = buffer; + } +#else + Dl_info info{}; + if (dladdr(reinterpret_cast(¤tModuleDirectory), &info) && + info.dli_fname) + path = info.dli_fname; +#endif + const auto pos = path.find_last_of(pathSeparator()); + return pos == std::string::npos ? std::string() : path.substr(0, pos); +} + +/// Lazily-loaded handle to ViennaLS_GPU plus the resolved entry points. +/// Constructed exactly once, on first use, by the function-local static in +/// runtime() below. +struct GpuRuntime { + bool available = false; + std::string status; + + int (*abiVersion)(void) = nullptr; + void *(*allocBuffers)(uint32_t, int, int) = nullptr; + void (*freeBuffers)(void *) = nullptr; + int (*isValid)(const void *) = nullptr; + const char *(*lastErrorMessage)(void) = nullptr; + int (*uploadNeighborIds)(void *, const uint32_t *, std::size_t) = nullptr; + int (*setupCSR)(void *, const uint32_t *, uint32_t, int) = nullptr; + int (*uploadSolverArrays)(void *, const double *, const double *, + const double *, uint32_t, std::size_t) = nullptr; + int (*uploadRhs)(void *, const double *, uint32_t) = nullptr; + int (*solveBiCGSTAB)(void *, double *, double, unsigned, double, unsigned *, + double *) = nullptr; + + GpuRuntime() { load(); } + + void load() { + // Search order: explicit override, next to the calling binary, the + // wheel's bundled library folder, then the default loader search path. + std::string candidates[4]; + int count = 0; + if (const char *override = std::getenv("VIENNALS_GPU_LIBRARY")) + candidates[count++] = override; + const std::string moduleDir = currentModuleDirectory(); + if (!moduleDir.empty()) { + candidates[count++] = moduleDir + pathSeparator() + gpuLibraryName(); + candidates[count++] = moduleDir + pathSeparator() + ".." + + pathSeparator() + "viennals.libs" + + pathSeparator() + gpuLibraryName(); + } + candidates[count++] = gpuLibraryName(); + + LibraryHandle lib = nullptr; + std::string attempts; + for (int i = 0; i < count && !lib; ++i) { + lib = openLibrary(candidates[i].c_str()); + if (!lib) { + // Report every attempt. The interesting failure is usually not the + // last one: when the library is found but the CUDA runtime is not, + // that candidate reports "libcudart.so.12: cannot open shared object + // file", while the final bare-soname attempt only reports that the + // library itself is missing. Keeping all of them means the message + // names the real cause instead of the last symptom. + attempts += "\n " + candidates[i] + ": " + lastLoadError(); + } + } + + if (!lib) { + status = "ViennaLS_GPU could not be loaded, so the GPU solver is " + "unavailable; using the CPU solver instead. Attempts:" + + attempts; + return; + } + + if (!resolveAll(lib)) { + status = "ViennaLS_GPU is missing expected entry points; it is probably " + "from a different ViennaLS version. Using the CPU solver " + "instead."; + return; + } + + if (abiVersion() != VIENNALS_GPU_ABI_VERSION) { + status = "ViennaLS_GPU reports ABI version " + + std::to_string(abiVersion()) + " but ViennaLS expects " + + std::to_string(VIENNALS_GPU_ABI_VERSION) + + ". Using the CPU solver instead."; + return; + } + + available = true; + status = "ViennaLS_GPU loaded."; + } + + /// Resolves every entry point; all-or-nothing so we never call through a + /// null pointer once `available` is set. + bool resolveAll(LibraryHandle lib) { + return resolve(lib, "viennalsGpuAbiVersion", abiVersion) && + resolve(lib, "viennalsGpuAllocBuffers", allocBuffers) && + resolve(lib, "viennalsGpuFreeBuffers", freeBuffers) && + resolve(lib, "viennalsGpuIsValid", isValid) && + resolve(lib, "viennalsGpuGetLastErrorMessage", lastErrorMessage) && + resolve(lib, "viennalsGpuUploadNeighborIds", uploadNeighborIds) && + resolve(lib, "viennalsGpuSetupCSR", setupCSR) && + resolve(lib, "viennalsGpuUploadSolverArrays", uploadSolverArrays) && + resolve(lib, "viennalsGpuUploadRhs", uploadRhs) && + resolve(lib, "viennalsGpuSolveBiCGSTAB", solveBiCGSTAB); + } + + template + static bool resolve(LibraryHandle lib, const char *name, Fn &target) { + target = reinterpret_cast(findSymbol(lib, name)); + return target != nullptr; + } +}; + +/// The single process-wide runtime instance. Initialization of a +/// function-local static is thread-safe since C++11, so the dlopen happens +/// exactly once even under concurrent first use. +inline const GpuRuntime &runtime() { + static GpuRuntime instance; + return instance; +} + +inline void *toHandle(GpuBiCGSTABBuffers *gpu) { + return reinterpret_cast(gpu); +} +inline const void *toHandle(const GpuBiCGSTABBuffers *gpu) { + return reinterpret_cast(gpu); +} + +} // namespace detail + +/// True when the GPU library and its CUDA runtime were loaded successfully. +/// Callers that want to warn before doing GPU-specific setup can check this. +inline bool gpuRuntimeAvailable() { return detail::runtime().available; } + +/// Human-readable reason describing why the GPU runtime is (un)available. +inline const char *gpuRuntimeStatusMessage() { + return detail::runtime().status.c_str(); +} + // Allocate GPU buffers for a solver with `n` nodes and `nFaces` (2*D) faces. // Returns nullptr if CUDA is unavailable. -GpuBiCGSTABBuffers *allocGpuBuffers(uint32_t n, int nFaces, - bool useIlu0Preconditioner); +inline GpuBiCGSTABBuffers *allocGpuBuffers(uint32_t n, int nFaces, + bool useIlu0Preconditioner) { + const auto &rt = detail::runtime(); + if (!rt.available) + return nullptr; + return reinterpret_cast( + rt.allocBuffers(n, nFaces, useIlu0Preconditioner ? 1 : 0)); +} // Free previously allocated GPU buffers. Safe to call with nullptr. -void freeGpuBuffers(GpuBiCGSTABBuffers *gpu); +inline void freeGpuBuffers(GpuBiCGSTABBuffers *gpu) { + const auto &rt = detail::runtime(); + if (rt.available) + rt.freeBuffers(detail::toHandle(gpu)); +} // Human-readable detail for the last GPU wrapper failure on this thread. -const char *gpuGetLastErrorMessage(); +// Falls back to the loader status when the library never loaded at all. +inline const char *gpuGetLastErrorMessage() { + const auto &rt = detail::runtime(); + return rt.available ? rt.lastErrorMessage() : rt.status.c_str(); +} // Is the buffer handle valid (non-null and successfully allocated)? -bool gpuIsValid(const GpuBiCGSTABBuffers *gpu); +inline bool gpuIsValid(const GpuBiCGSTABBuffers *gpu) { + const auto &rt = detail::runtime(); + return rt.available && rt.isValid(detail::toHandle(gpu)) != 0; +} // Upload geometry-fixed neighbor-ID array (face-major, kNoNode = 0xFFFFFFFF). // `count` must equal nFaces * n. -bool gpuUploadNeighborIds(GpuBiCGSTABBuffers *gpu, const uint32_t *nb, - std::size_t count); +inline bool gpuUploadNeighborIds(GpuBiCGSTABBuffers *gpu, const uint32_t *nb, + std::size_t count) { + const auto &rt = detail::runtime(); + return rt.available && + rt.uploadNeighborIds(detail::toHandle(gpu), nb, count) != 0; +} // Build the CSR sparsity pattern from h_nb (face-major, length nFaces*n), // upload to the device, and run CUSPARSE symbolic analysis for ILU(0) and // the two triangular solves. Must be called after gpuUploadNeighborIds and // before the first gpuUploadSolverArrays / gpuSolveBiCGSTAB call. -bool gpuSetupCSR(GpuBiCGSTABBuffers *gpu, const uint32_t *h_nb, uint32_t n, - int nFaces); +inline bool gpuSetupCSR(GpuBiCGSTABBuffers *gpu, const uint32_t *h_nb, + uint32_t n, int nFaces) { + const auto &rt = detail::runtime(); + return rt.available && + rt.setupCSR(detail::toHandle(gpu), h_nb, n, nFaces) != 0; +} // Upload per-solve arrays (diag, b, faceCoeffs) and re-factorize ILU(0). // `diagLen` == n, `coeffLen` == nFaces * n. -bool gpuUploadSolverArrays(GpuBiCGSTABBuffers *gpu, const double *diag, - const double *b, const double *coeff, - uint32_t diagLen, std::size_t coeffLen); +inline bool gpuUploadSolverArrays(GpuBiCGSTABBuffers *gpu, const double *diag, + const double *b, const double *coeff, + uint32_t diagLen, std::size_t coeffLen) { + const auto &rt = detail::runtime(); + return rt.available && rt.uploadSolverArrays(detail::toHandle(gpu), diag, b, + coeff, diagLen, coeffLen) != 0; +} // Upload only the RHS vector (d_b). Use when the matrix geometry is already // uploaded and only the right-hand side changes (e.g. successive Stokes // component solves that share the same stiffness matrix). -bool gpuUploadRhs(GpuBiCGSTABBuffers *gpu, const double *b, uint32_t n); +inline bool gpuUploadRhs(GpuBiCGSTABBuffers *gpu, const double *b, uint32_t n) { + const auto &rt = detail::runtime(); + return rt.available && rt.uploadRhs(detail::toHandle(gpu), b, n) != 0; +} // Run GPU BiCGSTAB. // x (length n, host): initial guess on entry, solution on exit. // outResidual is the raw (unnormalized) max-abs residual on exit. // Returns true only when the GPU solve converged and produced finite values. -bool gpuSolveBiCGSTAB(GpuBiCGSTABBuffers *gpu, double *x, double diagEps, - unsigned maxIter, double tolerance, - unsigned &outIterations, double &outResidual); +inline bool gpuSolveBiCGSTAB(GpuBiCGSTABBuffers *gpu, double *x, double diagEps, + unsigned maxIter, double tolerance, + unsigned &outIterations, double &outResidual) { + const auto &rt = detail::runtime(); + return rt.available && + rt.solveBiCGSTAB(detail::toHandle(gpu), x, diagEps, maxIter, tolerance, + &outIterations, &outResidual) != 0; +} } // namespace gpu } // namespace viennals diff --git a/include/viennals/lsOxidationDeformation.hpp b/include/viennals/lsOxidationDeformation.hpp index a4e49d4d..f8e44f82 100644 --- a/include/viennals/lsOxidationDeformation.hpp +++ b/include/viennals/lsOxidationDeformation.hpp @@ -894,37 +894,37 @@ class OxidationDeformation final : public VelocityField, gpuStokesBufs_ = nullptr; return; } - const bool tryGpu = (gpuMode_ == GpuMode::Gpu); + const bool tryGpu = (gpuMode_ == GpuMode::Gpu || gpuMode_ == GpuMode::Auto); const bool useIlu0 = (gpuPreconditioner_ == GpuPreconditioner::ILU0); + // Each setup step is chained with else-if: once one fails the handle is + // released and set to null, and every later step must be skipped rather + // than called with a null handle. (In GpuMode::Gpu reportGpuUnavailable + // throws, but in GpuMode::Auto it returns and execution continues here.) gpu::freeGpuBuffers(gpuPressBufs_); gpuPressBufs_ = nullptr; if (tryGpu) { gpuPressBufs_ = gpu::allocGpuBuffers(static_cast(n), 2 * D, useIlu0); if (!gpuPressBufs_) { - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "pressure solver CUDA buffers could not be " "allocated or the CUDA context could not be " - "initialized." + - gpuErrorDetail()); - } - if (!gpu::gpuUploadNeighborIds(gpuPressBufs_, pressNeighId32_.data(), - 2u * D * n)) { + "initialized."); + } else if (!gpu::gpuUploadNeighborIds( + gpuPressBufs_, pressNeighId32_.data(), 2u * D * n)) { gpu::freeGpuBuffers(gpuPressBufs_); gpuPressBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " - "uploading pressure GPU neighbor IDs failed." + - gpuErrorDetail()); - } - if (useIlu0 && !gpu::gpuSetupCSR(gpuPressBufs_, pressNeighId32_.data(), - static_cast(n), 2 * D)) { + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " + "uploading pressure GPU neighbor IDs failed."); + } else if (useIlu0 && + !gpu::gpuSetupCSR(gpuPressBufs_, pressNeighId32_.data(), + static_cast(n), 2 * D)) { gpu::freeGpuBuffers(gpuPressBufs_); gpuPressBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "CUSPARSE setup for the pressure GPU BiCGSTAB " - "solver failed." + - gpuErrorDetail()); + "solver failed."); } } @@ -934,28 +934,24 @@ class OxidationDeformation final : public VelocityField, gpuStokesBufs_ = gpu::allocGpuBuffers(static_cast(n), 2 * D, useIlu0); if (!gpuStokesBufs_) { - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "Stokes solver CUDA buffers could not be " "allocated or the CUDA context could not be " - "initialized." + - gpuErrorDetail()); - } - if (!gpu::gpuUploadNeighborIds(gpuStokesBufs_, stokesNeighId32_.data(), - 2u * D * n)) { + "initialized."); + } else if (!gpu::gpuUploadNeighborIds( + gpuStokesBufs_, stokesNeighId32_.data(), 2u * D * n)) { gpu::freeGpuBuffers(gpuStokesBufs_); gpuStokesBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " - "uploading Stokes GPU neighbor IDs failed." + - gpuErrorDetail()); - } - if (useIlu0 && !gpu::gpuSetupCSR(gpuStokesBufs_, stokesNeighId32_.data(), - static_cast(n), 2 * D)) { + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " + "uploading Stokes GPU neighbor IDs failed."); + } else if (useIlu0 && + !gpu::gpuSetupCSR(gpuStokesBufs_, stokesNeighId32_.data(), + static_cast(n), 2 * D)) { gpu::freeGpuBuffers(gpuStokesBufs_); gpuStokesBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "CUSPARSE setup for the Stokes GPU BiCGSTAB " - "solver failed." + - gpuErrorDetail()); + "solver failed."); } } @@ -966,32 +962,29 @@ class OxidationDeformation final : public VelocityField, gpuHarmonicBufs_ = gpu::allocGpuBuffers(static_cast(n), 2 * D, useIlu0); if (!gpuHarmonicBufs_) { - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "harmonic solver CUDA buffers could not be " - "allocated." + - gpuErrorDetail()); - } - if (!gpu::gpuUploadNeighborIds(gpuHarmonicBufs_, stokesNeighId32_.data(), - 2u * D * n)) { + "allocated."); + } else if (!gpu::gpuUploadNeighborIds( + gpuHarmonicBufs_, stokesNeighId32_.data(), 2u * D * n)) { gpu::freeGpuBuffers(gpuHarmonicBufs_); gpuHarmonicBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " - "uploading harmonic GPU neighbor IDs failed." + - gpuErrorDetail()); - } - if (useIlu0 && - !gpu::gpuSetupCSR(gpuHarmonicBufs_, stokesNeighId32_.data(), - static_cast(n), 2 * D)) { + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " + "uploading harmonic GPU neighbor IDs failed."); + } else if (useIlu0 && + !gpu::gpuSetupCSR(gpuHarmonicBufs_, stokesNeighId32_.data(), + static_cast(n), 2 * D)) { gpu::freeGpuBuffers(gpuHarmonicBufs_); gpuHarmonicBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDeformation: GPU mode was selected, but " + reportGpuUnavailable("OxidationDeformation: GPU mode was selected, but " "CUSPARSE setup for the harmonic GPU BiCGSTAB " - "solver failed." + - gpuErrorDetail()); + "solver failed."); } } - if (tryGpu) + // Only claim the GPU backend when every solver actually has usable + // buffers; under GpuMode::Auto any of them may have degraded to the CPU. + if (tryGpu && gpuPressBufs_ && gpuStokesBufs_ && gpuHarmonicBufs_) logDeformationBackend("GPU BiCGSTAB", "pressure/Stokes/harmonic, preconditioner=" + std::string(useIlu0 ? "ILU0" : "Jacobi")); @@ -1005,6 +998,19 @@ class OxidationDeformation final : public VelocityField, return std::string(" Detail: ") + detail; return {}; } + + /// Reports that the GPU solver cannot be used. + /// GpuMode::Gpu means "GPU or nothing", so this raises an error (which + /// aborts). GpuMode::Auto asked for the GPU only if it happens to work, so + /// it warns and lets the caller continue on to the CPU solver. + void reportGpuUnavailable(const std::string &message) const { + if (gpuMode_ == GpuMode::Auto) { + VIENNACORE_LOG_WARNING(message + gpuErrorDetail() + + " Falling back to the CPU solver."); + } else { + VIENNACORE_LOG_ERROR(message + gpuErrorDetail()); + } + } #endif void logDeformationBackend(const std::string &backend, @@ -1098,6 +1104,10 @@ class OxidationDeformation final : public VelocityField, VIENNACORE_LOG_ERROR("OxidationDeformation: explicit GPU mode was " "requested, but ViennaLS was built without " "VIENNALS_GPU_BICGSTAB."); + } else if (gpuMode_ == GpuMode::Auto) { + VIENNACORE_LOG_WARNING("OxidationDeformation: GPU mode Auto was " + "requested, but ViennaLS was built without " + "VIENNALS_GPU_BICGSTAB. Using the CPU solver."); } #endif } diff --git a/include/viennals/lsOxidationDiffusion.hpp b/include/viennals/lsOxidationDiffusion.hpp index c1cbc1b8..293a36fd 100644 --- a/include/viennals/lsOxidationDiffusion.hpp +++ b/include/viennals/lsOxidationDiffusion.hpp @@ -25,7 +25,13 @@ namespace viennals { /// GPU failures are reported and not silently fallen back to CPU. enum class GpuMode { Cpu, ///< Always use CPU (default) - Gpu ///< Always use GPU; fail if unavailable or unsuccessful + Gpu, ///< Always use GPU; fail if unavailable or unsuccessful + /// Use the GPU when it is usable, otherwise warn and fall back to the CPU. + /// This covers the GPU being *unavailable*: no CUDA runtime installed, no + /// device, or device setup failing. A GPU that is present and initialises + /// but then fails mid-solve is still reported as an error, because that + /// indicates a malfunction rather than an absent GPU. + Auto }; /// Selects the preconditioner used by the GPU BiCGSTAB solver. @@ -647,12 +653,20 @@ class OxidationDiffusion final : public VelocityField, gpu::freeGpuBuffers(gpuBufs_); gpuBufs_ = nullptr; - const bool tryGpu = (gpuMode_ == GpuMode::Gpu); + const bool tryGpu = (gpuMode_ == GpuMode::Gpu || gpuMode_ == GpuMode::Auto); if (tryGpu) { const bool useIlu0 = gpuPreconditioner_ == GpuPreconditioner::ILU0; gpuBufs_ = gpu::allocGpuBuffers(static_cast(n), 2 * D, useIlu0); - if (gpuBufs_) { + // Each setup step is chained with else-if: once one fails the handle is + // released and set to null, and every later step must be skipped rather + // than called with a null handle. (In GpuMode::Gpu reportGpuUnavailable + // throws, but in GpuMode::Auto it returns and execution continues here.) + if (!gpuBufs_) { + reportGpuUnavailable("OxidationDiffusion: GPU mode was selected, but " + "CUDA buffers could not be allocated or the CUDA " + "context could not be initialized."); + } else { // Convert std::size_t neighbor IDs to uint32_t for the device const std::size_t nf = 2u * D * n; std::vector nb32(nf); @@ -663,35 +677,32 @@ class OxidationDiffusion final : public VelocityField, if (!gpu::gpuUploadNeighborIds(gpuBufs_, nb32.data(), nf)) { gpu::freeGpuBuffers(gpuBufs_); gpuBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDiffusion: GPU mode was selected, but " - "uploading GPU neighbor IDs failed." + - gpuErrorDetail()); - } - if (useIlu0 && !gpu::gpuSetupCSR(gpuBufs_, nb32.data(), - static_cast(n), 2 * D)) { + reportGpuUnavailable("OxidationDiffusion: GPU mode was selected, but " + "uploading GPU neighbor IDs failed."); + } else if (useIlu0 && + !gpu::gpuSetupCSR(gpuBufs_, nb32.data(), + static_cast(n), 2 * D)) { gpu::freeGpuBuffers(gpuBufs_); gpuBufs_ = nullptr; - VIENNACORE_LOG_ERROR("OxidationDiffusion: GPU mode was selected, but " + reportGpuUnavailable("OxidationDiffusion: GPU mode was selected, but " "CUSPARSE setup for the GPU BiCGSTAB solver " - "failed." + - gpuErrorDetail()); + "failed."); + } else { + // Only claim the GPU backend once every setup step has succeeded. + logDiffusionBackend("GPU BiCGSTAB", + "preconditioner=" + + std::string(useIlu0 ? "ILU0" : "Jacobi")); + loggedBackend = true; } - logDiffusionBackend("GPU BiCGSTAB", - "preconditioner=" + - std::string(useIlu0 ? "ILU0" : "Jacobi")); - loggedBackend = true; - } else { - VIENNACORE_LOG_ERROR("OxidationDiffusion: GPU mode was selected, but " - "CUDA buffers could not be " - "allocated or the CUDA context could not be " - "initialized." + - gpuErrorDetail()); } } #endif if (!loggedBackend) { #ifdef VIENNALS_GPU_BICGSTAB - logDiffusionBackend("CPU BiCGSTAB", "GPU mode not selected"); + logDiffusionBackend("CPU BiCGSTAB", + gpuMode_ == GpuMode::Cpu + ? "GPU mode not selected" + : "GPU requested but unavailable"); #else logDiffusionBackend("CPU BiCGSTAB", "ViennaLS was built without GPU BiCGSTAB support"); @@ -945,7 +956,9 @@ class OxidationDiffusion final : public VelocityField, .addTiming(tag + " GPU BiCGSTAB", tSolve) .print(); } - VIENNACORE_LOG_ERROR( + // Under GpuMode::Auto this warns and execution continues into the CPU + // BiCGSTAB below, which recomputes the solution from the initial guess. + reportGpuUnavailable( "OxidationDiffusion: GPU mode was selected, but GPU BiCGSTAB " "failed, did not converge, or produced non-finite concentrations " "(iters=" + @@ -957,6 +970,10 @@ class OxidationDiffusion final : public VelocityField, VIENNACORE_LOG_ERROR("OxidationDiffusion: explicit GPU mode was " "requested, but ViennaLS was built without " "VIENNALS_GPU_BICGSTAB."); + } else if (gpuMode_ == GpuMode::Auto) { + VIENNACORE_LOG_WARNING("OxidationDiffusion: GPU mode Auto was requested, " + "but ViennaLS was built without " + "VIENNALS_GPU_BICGSTAB. Using the CPU solver."); } #endif @@ -1151,6 +1168,19 @@ class OxidationDiffusion final : public VelocityField, return std::string(" Detail: ") + detail; return {}; } + + /// Reports that the GPU solver cannot be used. + /// GpuMode::Gpu means "GPU or nothing", so this raises an error (which + /// aborts). GpuMode::Auto asked for the GPU only if it happens to work, so + /// it warns and lets the caller continue on to the CPU solver. + void reportGpuUnavailable(const std::string &message) const { + if (gpuMode_ == GpuMode::Auto) { + VIENNACORE_LOG_WARNING(message + gpuErrorDetail() + + " Falling back to the CPU solver."); + } else { + VIENNACORE_LOG_ERROR(message + gpuErrorDetail()); + } + } #endif void logDiffusionBackend(const std::string &backend, diff --git a/include/viennals/lsVersion.hpp b/include/viennals/lsVersion.hpp index 2ed5622d..d19776a4 100644 --- a/include/viennals/lsVersion.hpp +++ b/include/viennals/lsVersion.hpp @@ -5,10 +5,10 @@ namespace viennals { // Version information generated by CMake -inline constexpr const char *version = "5.8.4"; +inline constexpr const char *version = "5.8.5"; inline constexpr int versionMajor = 5; inline constexpr int versionMinor = 8; -inline constexpr int versionPatch = 4; +inline constexpr int versionPatch = 5; // Utility functions for version comparison inline constexpr uint32_t versionAsInteger() { diff --git a/pyproject.toml b/pyproject.toml index 259b0158..d3336c7d 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -6,7 +6,7 @@ requires = [ build-backend = "scikit_build_core.build" [project] -version = "5.8.4" +version = "5.8.5" name = "ViennaLS" readme = "README.md" license = {file = "LICENSE"} diff --git a/python/CMakeLists.txt b/python/CMakeLists.txt index 7e05fd33..c67869f4 100644 --- a/python/CMakeLists.txt +++ b/python/CMakeLists.txt @@ -64,7 +64,18 @@ endif() if(VIENNALS_USE_GPU) message(STATUS "[ViennaLS] Adding GPU support for 'viennals' python package") if(TARGET ViennaLS_GPU) + # Built alongside the extension but never linked into it: _core must stay + # importable on machines without a CUDA runtime. It is placed next to + # _core so the dlopen() search in lsOxidationBiCGSTABInterface.hpp finds it + # via the calling module's own directory. add_dependencies(${MOD} ViennaLS_GPU) + set_target_properties( + ViennaLS_GPU PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/${PKG}" + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/${PKG}") + install( + TARGETS ViennaLS_GPU + LIBRARY DESTINATION ${PKG} + RUNTIME DESTINATION ${PKG}) endif() endif() @@ -204,17 +215,6 @@ if((VIENNALS_PACKAGE_PYTHON AND NOT APPLE) OR MSVC) endif() endif() - if(VIENNALS_USE_GPU) - foreach(_CUDA_TARGET CUDA::cudart CUDA::cusparse) - if(TARGET ${_CUDA_TARGET}) - add_custom_command( - TARGET ${MOD} - POST_BUILD - COMMAND ${CMAKE_COMMAND} -E copy_directory $ - ${VIENNALS_LIB_FOLDER}) - endif() - endforeach() - endif() endif() if(NOT VIENNALS_PACKAGE_PYTHON) diff --git a/python/pyWrap.hpp b/python/pyWrap.hpp index 03fb95b4..13aaf197 100644 --- a/python/pyWrap.hpp +++ b/python/pyWrap.hpp @@ -95,6 +95,7 @@ inline void bindOxidationSharedTypes(py::module &module) { py::enum_(module, "GpuMode", py::module_local()) .value("Cpu", GpuMode::Cpu) .value("Gpu", GpuMode::Gpu) + .value("Auto", GpuMode::Auto) .export_values(); py::enum_(module, "GpuPreconditioner", py::module_local()) diff --git a/python/viennals/__init__.pyi b/python/viennals/__init__.pyi index 3f3322e3..6df3496d 100644 --- a/python/viennals/__init__.pyi +++ b/python/viennals/__init__.pyi @@ -133,7 +133,7 @@ ILU0: _core.GpuPreconditioner # value = Jacobi: _core.GpuPreconditioner # value = PROXY_DIM: int = 2 _SHARED_OXIDATION_TYPES: tuple = ('OxidationParameters', 'OxidationPresets', 'OxidationDeformationParameters', 'OxidationMaskParameters', 'OxidationCouplingParameters') -__version__: str = '5.8.4' +__version__: str = '5.8.5' _name: str = 'OxidationCouplingParameters' -version: str = '5.8.4' +version: str = '5.8.5' _C = _core diff --git a/python/viennals/_core.pyi b/python/viennals/_core.pyi index 1da8ac10..f86251b4 100644 --- a/python/viennals/_core.pyi +++ b/python/viennals/_core.pyi @@ -1063,6 +1063,6 @@ Cpu: GpuMode # value = Gpu: GpuMode # value = ILU0: GpuPreconditioner # value = Jacobi: GpuPreconditioner # value = -__version__: str = '5.8.4' -version: str = '5.8.4' +__version__: str = '5.8.5' +version: str = '5.8.5' IntegrationSchemeEnum = SpatialSchemeEnum diff --git a/vcpkg.json b/vcpkg.json index bd5200f7..8789726a 100644 --- a/vcpkg.json +++ b/vcpkg.json @@ -2,7 +2,14 @@ "name": "viennals-ci", "version-string": "1.0.0", "dependencies": [ - "vtk" + "vtk", + { + "name": "hdf5", + "default-features": false, + "features": [ + "zlib" + ] + } ], "builtin-baseline": "72d876b20f06129a1b1e0db80c7075fe3e419d7c" -} \ No newline at end of file +}