diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml index 78e7adc1..8efab491 100644 --- a/.github/workflows/build.yml +++ b/.github/workflows/build.yml @@ -2,7 +2,9 @@ name: Build on: push: - branches: [ "main", "dev" ] + branches: [ "main" ] + pull_request: + branches: [ "main" ] jobs: build: @@ -26,3 +28,26 @@ jobs: cargo build --release cmake -S . -B build -DCMAKE_CXX_COMPILER=clang++-14 -DCMAKE_C_COMPILER=clang-14 -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_EXPORT_COMPILE_COMMANDS=ON -DSHUFFLE_METHOD=g++11 cmake --build build -j 32 + build-gdr: + runs-on: self-hosted # or `ubuntu-22.04` + + steps: + - uses: actions/checkout@v4 + + - name: Prepare + run: | + sudo apt install -y cmake libuv1-dev liblz4-dev liblzma-dev libdouble-conversion-dev libdwarf-dev libunwind-dev + sudo apt install -y libaio-dev libgflags-dev libgoogle-glog-dev libgtest-dev libgmock-dev clang-format-14 clang-14 clang-tidy-14 lld-14 + sudo apt install -y libgoogle-perftools-dev google-perftools libssl-dev gcc-12 g++-12 libboost-all-dev meson rustc cargo + sudo apt install -y nvidia-cuda-toolkit + wget https://github.com/apple/foundationdb/releases/download/7.1.61/foundationdb-clients_7.1.61-1_amd64.deb && sudo dpkg -i foundationdb-clients_7.1.61-1_amd64.deb + git clone https://github.com/libfuse/libfuse.git libfuse -b fuse-3.16.2 --depth 1 && mkdir libfuse/build && cd libfuse/build && meson setup .. && ninja && sudo ninja install && cd ../.. && rm -rf libfuse + git submodule update --init --recursive + ./patches/apply.sh + + - name: Build + run: | + cargo build --release + cmake -S . -B build -DCMAKE_CXX_COMPILER=clang++-14 -DCMAKE_C_COMPILER=clang-14 -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_EXPORT_COMPILE_COMMANDS=ON -DSHUFFLE_METHOD=g++11 -DHF3FS_ENABLE_GDR=ON + cmake --build build -j 32 + ctest --test-dir build -R test_gdr --output-on-failure diff --git a/CMakeLists.txt b/CMakeLists.txt index c21dbb82..afec99f7 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -201,6 +201,7 @@ include(cmake/DumpConfig.cmake) include(cmake/Jemalloc.cmake) include(cmake/ApacheArrow.cmake) include(cmake/AddCrate.cmake) +include(cmake/Cuda.cmake) configure_file(cmake/CTestCustom.cmake ${CMAKE_BINARY_DIR} @ONLY) add_subdirectory(src) diff --git a/cmake/CLangFormat.cmake b/cmake/CLangFormat.cmake index c74b3dd9..31bf7211 100644 --- a/cmake/CLangFormat.cmake +++ b/cmake/CLangFormat.cmake @@ -1,5 +1,5 @@ -set(CLANG_FORMAT "/usr/bin/clang-format-14") -if(EXISTS ${CLANG_FORMAT}) +find_program(CLANG_FORMAT NAMES clang-format-14 clang-format) +if(CLANG_FORMAT) message(STATUS "Found clang-format at ${CLANG_FORMAT}") set(SOURCE_DIRS diff --git a/cmake/Cuda.cmake b/cmake/Cuda.cmake new file mode 100644 index 00000000..49382e33 --- /dev/null +++ b/cmake/Cuda.cmake @@ -0,0 +1,89 @@ +# CUDA/GDR Support for 3FS +# +# This module provides optional CUDA support for GPU Direct RDMA (GDR). +# GDR enables direct data transfers between storage and GPU memory. +# +# Options: +# HF3FS_ENABLE_GDR - Enable GPU Direct RDMA support (default: OFF) +# +# When HF3FS_ENABLE_GDR is ON: +# - CUDA Toolkit will be searched +# - HF3FS_ENABLE_GDR=1 will be defined for opted-in targets +# - CUDA libraries will be linked to relevant targets +# +# Usage in CMakeLists.txt: +# include(cmake/Cuda.cmake) +# target_add_gdr_support(mytarget SCOPE PRIVATE) + +option(HF3FS_ENABLE_GDR "Enable GPU Direct RDMA (GDR) support" OFF) + +set(HF3FS_CUDA_LIBRARIES "") +set(HF3FS_CUDA_INCLUDE_DIRS "") + +if(HF3FS_ENABLE_GDR) + message(STATUS "GDR support requested, searching for CUDA...") + + # Try to find CUDA Toolkit + # CMake 3.17+ has FindCUDAToolkit which is preferred + if(CMAKE_VERSION VERSION_GREATER_EQUAL "3.17") + find_package(CUDAToolkit QUIET) + if(CUDAToolkit_FOUND) + if(NOT TARGET CUDA::cudart OR NOT TARGET CUDA::cuda_driver) + message(FATAL_ERROR + "CUDAToolkit was found, but required CUDA::cudart and CUDA::cuda_driver targets are unavailable") + endif() + set(HF3FS_CUDA_LIBRARIES CUDA::cudart CUDA::cuda_driver) + set(HF3FS_CUDA_INCLUDE_DIRS ${CUDAToolkit_INCLUDE_DIRS}) + message(STATUS "Found CUDA Toolkit ${CUDAToolkit_VERSION}") + message(STATUS " CUDA include: ${CUDAToolkit_INCLUDE_DIRS}") + message(STATUS " CUDA libraries: CUDA::cudart;CUDA::cuda_driver") + endif() + else() + # Fallback for older CMake + find_package(CUDA QUIET) + if(CUDA_FOUND AND CUDA_CUDA_LIBRARY) + set(HF3FS_CUDA_LIBRARIES ${CUDA_LIBRARIES} ${CUDA_CUDA_LIBRARY}) + list(REMOVE_DUPLICATES HF3FS_CUDA_LIBRARIES) + set(HF3FS_CUDA_INCLUDE_DIRS ${CUDA_INCLUDE_DIRS}) + message(STATUS "Found CUDA ${CUDA_VERSION}") + message(STATUS " CUDA include: ${CUDA_INCLUDE_DIRS}") + message(STATUS " CUDA libraries: ${HF3FS_CUDA_LIBRARIES}") + elseif(CUDA_FOUND) + message(STATUS "Legacy CUDA discovery did not provide the required driver library") + endif() + endif() + + if(NOT HF3FS_CUDA_LIBRARIES) + message(FATAL_ERROR + "HF3FS_ENABLE_GDR=ON requires CUDA runtime and driver libraries, but required CUDA components were not found") + endif() + + message(STATUS "GDR support enabled") + + # Check for nvidia_peermem (informational only) + if(EXISTS "/sys/module/nvidia_peermem") + message(STATUS "nvidia_peermem module detected") + else() + message(STATUS "nvidia_peermem module not detected (required at runtime for GDR)") + endif() +else() + message(STATUS "GDR support disabled (use -DHF3FS_ENABLE_GDR=ON to enable)") +endif() + +# Helper function to add GDR support to a target +function(target_add_gdr_support TARGET) + cmake_parse_arguments(GDR "" "SCOPE" "" ${ARGN}) + if(NOT GDR_SCOPE) + set(GDR_SCOPE PRIVATE) + endif() + + if(HF3FS_ENABLE_GDR) + target_compile_definitions(${TARGET} ${GDR_SCOPE} HF3FS_ENABLE_GDR=1) + target_include_directories(${TARGET} PRIVATE ${HF3FS_CUDA_INCLUDE_DIRS}) + # Project targets use the plain target_link_libraries signature. Keeping + # that signature also propagates CUDA link requirements from static + # libraries to their final consumers. + target_link_libraries(${TARGET} ${HF3FS_CUDA_LIBRARIES}) + message(STATUS "Added GDR support to target: ${TARGET} (${GDR_SCOPE})") + endif() +endfunction() diff --git a/docs/gdr-integration.md b/docs/gdr-integration.md new file mode 100644 index 00000000..97553e77 --- /dev/null +++ b/docs/gdr-integration.md @@ -0,0 +1,637 @@ +# GPU Direct RDMA (GDR) Integration Guide + +## Overview + +GPU Direct RDMA (GDR) lets the 3FS FUSE data path transfer data directly +between RDMA and CUDA device memory. The application process does not register +GPU memory with InfiniBand. It allocates or wraps CUDA memory, exports a CUDA +IPC capability, and publishes that capability through the 3FS virtual IOV +namespace. The FUSE process imports the CUDA allocation and owns the +`ibv_mr` objects for the published view. + +This split creates two distinct capabilities: + +1. **Application-side CUDA IPC capability.** `hf3fs_gdr_available()` reports + whether the local process can use CUDA IPC on at least one visible device. + It does not inspect FUSE, `nvidia_peermem`, the selected NIC, or + `ibv_reg_mr`. +2. **FUSE-side GDR capability.** When FUSE resolves a GDR IOV publication, it + imports the CUDA IPC allocation and registers the requested view with every + active IB device. + +A successful application capability query is therefore only a prerequisite. +The return value from the publication operation (`hf3fs_iovcreate_device()` or +`hf3fs_iovwrap_device()`) is the final per-buffer result: success means FUSE +accepted the symlink, imported the CUDA allocation, and registered the +published view with every active IB device. `ibv_reg_mr` can still fail after +`hf3fs_gdr_available()` returned `true`. + +## Prerequisites + +### Hardware + +- A CUDA device that supports unified addressing and is not in prohibited + compute mode. +- An InfiniBand or RoCE device that can register the CUDA allocation. +- A GPU/NIC/driver combination that supports peer-memory registration. + +PCIe locality affects performance. It does not replace the final registration +check. + +### Software + +- A working CUDA runtime and driver. +- RDMA userspace and kernel drivers. +- NVIDIA peer-memory support, normally provided by `nvidia_peermem`: + + ```bash + sudo modprobe nvidia_peermem + lsmod | grep nvidia_peermem + ``` + +### Build + +Build 3FS with GDR support: + +```bash +cmake -DHF3FS_ENABLE_GDR=ON ... +``` + +- `HF3FS_ENABLE_GDR` requests CUDA/GDR support at CMake configure time. +- The same name is defined as `HF3FS_ENABLE_GDR=1` for GDR-enabled C++ targets. +- Configuration fails when GDR is requested but the CUDA runtime or driver + library is unavailable. There is no separate runtime environment switch. + +`hf3fs_iovopen_device()` and `hf3fs_iovwrap_device()` are declared only when +the consuming target is compiled with `HF3FS_ENABLE_GDR`. +`hf3fs_iovcreate_device()` remains declared in all builds for API +compatibility, but returns `-ENOTSUP` when CUDA/GDR support is not compiled in. + +### Runtime + +- Start `IBManager` before publishing a GPU IOV. Per-buffer registration fails + if IB has not started or no active IB device exists. +- Run the application on the same host as the FUSE process that serves the + mount. +- The application and FUSE process must be allowed to share CUDA allocations + through CUDA IPC. +- The FUSE virtual IOV namespace must accept the `.gdr.d{device_id}` + publication. Publication failure is returned to the application as + `-errno`. + +### No Host Fallback + +The device APIs never return a host IOV. A build without CUDA/GDR support, no +locally usable CUDA IPC device, an invalid device pointer, CUDA IPC failure, +FUSE import failure, or RDMA registration failure is returned as an error. +A CUDA device pointer is never passed to the host-memory registration path. +Applications that want host memory must call `hf3fs_iovcreate()` explicitly. + +## API Reference + +All functions return `0` on success and a negative errno value on failure +unless stated otherwise. The caller owns the `struct hf3fs_iov` object itself; +the library fills its fields. + +Include: + +```c +#include +``` + +### Device IOV Creation + +#### `hf3fs_iovcreate_device` + +Allocate CUDA memory, export the whole allocation through CUDA IPC, and +publish a view covering the allocation: + +```c +int hf3fs_iovcreate_device(struct hf3fs_iov *iov, + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); +``` + +Constraints: + +- `iov` and `hf3fs_mount_point` must be non-null. +- `size` must be non-zero. +- `block_size` must be `0`; GDR IOV block partitioning is unsupported. +- `device_id` must name a locally usable CUDA IPC device. + +On the CUDA path, the application owns the allocation and the publication. +The application does **not** create or own an RDMA MR. Publication success is +the final FUSE-side import/registration result. On destroy, the publisher +unlinks the publication and frees its CUDA allocation. + +Representative errors include `-EINVAL` for invalid arguments, +`-ENOTSUP` when CUDA/GDR support or local CUDA IPC capability is unavailable, +`-ENODEV` for an unavailable selected device, `-ENOMEM` for allocation +failure, `-EIO` for CUDA IPC or FUSE registration failures, and the exact +negative errno returned by publication (for example, `-EEXIST` for a +conflicting key). No failure is converted into a host allocation. + +--- + +#### `hf3fs_iovopen_device` + +Open an existing device IOV by UUID: + +```c +int hf3fs_iovopen_device(struct hf3fs_iov *iov, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); +``` + +This API reads the existing GDR v2 publication, imports its CUDA allocation, +and exposes the published view in the importing application. It does not +create a second publication and does not register an application-side MR. + +Constraints: + +- The original publisher and its CUDA allocation must still be alive. +- `size` and `device_id` must match the publication. +- `block_size` must be `0`. +- The imported handle must resolve to an allocation base, and its actual + allocation size must match the v2 `allocation` field. + +The importer borrows the publication. Its `hf3fs_iovunlink()` is a no-op for +that GPU IOV, and its `hf3fs_iovdestroy()` closes only its CUDA IPC mapping and +local metadata. + +Representative errors include `-ENOTSUP` when local CUDA IPC capability is +unavailable, `-ENOENT` when the publication does not exist, `-ENODEV` for an +unavailable device, `-EINVAL` for malformed or mismatched metadata, and +`-EIO` for CUDA import failures. + +**Availability:** declared only with `HF3FS_ENABLE_GDR`. + +--- + +#### `hf3fs_iovwrap_device` + +Publish a view of an existing CUDA allocation: + +```c +int hf3fs_iovwrap_device(struct hf3fs_iov *iov, + void *device_ptr, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); +``` + +`device_ptr` may be the allocation base or a pointer to a view/suballocation, +including a PyTorch tensor view. The implementation uses +`cuMemGetAddressRange()` to discover the underlying allocation base, full +allocation size, and view offset. It verifies that +`[device_ptr, device_ptr + size)` is within that allocation and exports the +allocation base through CUDA IPC. + +The caller retains ownership of the CUDA allocation and must keep the entire +underlying allocation alive until the publication and all importers are gone. +The wrapped IOV owns the publication but not the allocation. +`hf3fs_iovdestroy()` unlinks the publication and releases 3FS metadata; it does +not call `cudaFree()` on the wrapped allocation. + +Constraints: + +- `device_ptr`, `id`, and `hf3fs_mount_point` must be non-null. +- `size` must be non-zero and fit within the discovered allocation. +- `device_id` must match the CUDA pointer's device. +- `id` must be unique in the mount's IOV namespace. +- `block_size` must be `0`. + +Representative errors include `-ENOTSUP`, `-ENODEV`, `-EINVAL`, `-ENOMEM`, +`-EIO`, and publication errors returned as negative errno values. The wrapped +allocation remains caller-owned on every failure path. + +**Availability:** declared only with `HF3FS_ENABLE_GDR`. + +### IOV Lifecycle + +#### `hf3fs_iovdestroy` + +```c +void hf3fs_iovdestroy(struct hf3fs_iov *iov); +``` + +The effect depends on how the IOV was obtained: + +- `hf3fs_iovcreate_device()`: unlink the owned publication, then free the CUDA + allocation. +- `hf3fs_iovwrap_device()`: unlink the owned publication, but leave the + caller-owned CUDA allocation untouched. +- `hf3fs_iovopen_device()`: close the imported CUDA mapping without unlinking + the borrowed publication. + +For cross-process sharing, importers destroy first and the exporter destroys +last. There is no distributed reference count. + +#### `hf3fs_iovunlink` + +```c +void hf3fs_iovunlink(struct hf3fs_iov *iov); +``` + +For a create/wrap publisher, remove the GDR publication while leaving the +local IOV object alive. For an open importer, this is intentionally a no-op: +an importer is not allowed to remove the exporter's publication. + +The publisher must wait for all I/O using the IOV to complete and destroy every +application importer before unlinking or destroying the publication. CUDA +requires the exported allocation to remain alive until every imported mapping +has been closed. + +### Synchronization + +#### `hf3fs_iovsync` + +```c +int hf3fs_iovsync(const struct hf3fs_iov *iov, int direction); +``` + +For a GPU IOV, this selects its CUDA device and calls +`cudaDeviceSynchronize()`. The current implementation accepts the direction +convention but uses the same device-wide synchronization for both values: + +- `direction = 0`: application GPU writes must be visible before submitting a + storage write. +- `direction = 1`: RDMA writes must be visible before the application consumes + data on the GPU. + +The function is a no-op for host IOVs. It returns `-EINVAL` for an invalid GPU +IOV, `-ENODEV` if its device cannot be selected, and `-EIO` if synchronization +fails. + +### Capability Queries + +#### `hf3fs_gdr_available` + +```c +bool hf3fs_gdr_available(void); +``` + +Returns `true` when at least one locally visible CUDA device supports the +application's CUDA IPC requirements (unified addressing and a usable compute +mode). It does not query FUSE or prove that an MR can be created. + +Use the return value of create/wrap publication, which includes FUSE-side +import and registration, as the final per-buffer decision. + +#### `hf3fs_gdr_device_count` + +```c +int hf3fs_gdr_device_count(void); +``` + +Returns the local CUDA device count, or `0` when CUDA enumeration is +unavailable. It does not query FUSE or probe RDMA registration. + +### Additional Utilities + +#### `hf3fs_iov_mem_type` + +```c +enum hf3fs_mem_type hf3fs_iov_mem_type(const struct hf3fs_iov *iov); +``` + +Returns `HF3FS_MEM_DEVICE` for a tracked GPU IOV and `HF3FS_MEM_HOST` +otherwise. The enum also reserves `HF3FS_MEM_MANAGED`, but these device APIs +do not report that value. + +#### `hf3fs_iov_device_id` + +```c +int hf3fs_iov_device_id(const struct hf3fs_iov *iov); +``` + +Returns the CUDA device index for a tracked GPU IOV, or `-1` for a host or +invalid IOV. + +### I/O Submission + +The normal user I/O APIs accept a GPU IOV: + +- `hf3fs_prep_io()` uses `ptr` as an address within the IOV view. +- `hf3fs_submit_ios()` submits prepared operations. +- `hf3fs_wait_for_ios()` returns completion entries. + +For a GPU IOV, `iov->base` and pointers derived from it are CUDA device +pointers. The application passes them as opaque addresses; it must not +dereference them on the CPU. + +## Usage Examples + +### Example A: Require a GDR Allocation + +`hf3fs_gdr_available()` is useful for an early local check, but the create +result is authoritative: + +```c +struct hf3fs_iov iov = {0}; + +if (!hf3fs_gdr_available()) { + /* No local CUDA IPC capability; do not request mandatory GDR. */ + return -ENOTSUP; +} + +int rc = hf3fs_iovcreate_device(&iov, "/mnt/3fs", + 128UL * 1024 * 1024, + /*block_size=*/0, + /*device_id=*/0); +if (rc != 0) { + /* Includes FUSE CUDA import or final MR registration failure. */ + return rc; +} + +/* Prepare and complete I/O using addresses within iov.base. */ + +rc = hf3fs_iovsync(&iov, /*direction=*/1); +hf3fs_iovdestroy(&iov); +return rc; +``` + +No application-side `ibv_reg_mr()` call is required or owned by this code. + +### Example B: Publish a PyTorch View + +The wrapped pointer may be inside a larger allocator-owned CUDA allocation: + +```c +/* tensor_ptr and tensor_bytes come from a live CUDA PyTorch tensor/view. + * uuid is a unique 16-byte identifier supplied by the application. + */ +struct hf3fs_iov iov = {0}; +int rc = hf3fs_iovwrap_device(&iov, + tensor_ptr, + uuid, + "/mnt/3fs", + tensor_bytes, + /*block_size=*/0, + tensor_device); +if (rc != 0) { + /* The tensor remains caller-owned on every failure path. */ + return rc; +} + +/* Use iov.base only as a CUDA/RDMA address. */ + +hf3fs_iovdestroy(&iov); /* Removes the publication; does not free tensor_ptr. */ +``` + +The whole underlying CUDA allocation remains the CUDA IPC sharing unit even +when only a tensor view is published. The application must keep that +allocation alive until all FUSE and application importers have released it. + +### Example C: Cross-Process Open + +The exporter publishes and shares `iov.id` through an application-defined +control channel. An importer borrows that publication: + +```c +struct hf3fs_iov imported = {0}; +int rc = hf3fs_iovopen_device(&imported, + exported_uuid, + "/mnt/3fs", + exported_view_size, + /*block_size=*/0, + exported_device); +if (rc != 0) { + return rc; +} + +/* Submit I/O against imported.base. */ + +hf3fs_iovdestroy(&imported); /* Closes this mapping; does not unlink. */ +``` + +Required shutdown order: + +1. Stop and reap I/O that uses the imported view. +2. Destroy every application importer. +3. Destroy the create/wrap exporter last. + +The publisher owns the publication. Importer-first cleanup cannot unlink it. + +## Runtime Behavior + +### Capability and Ownership Model + +The application process owns CUDA allocation/IPC publication only: + +- create owns both the allocation and publication; +- wrap owns the publication while the caller owns the allocation; +- open owns an imported CUDA mapping and borrows the publication. + +FUSE owns the data-plane import and registration: + +- the v2 URI is parsed and the full CUDA allocation is imported; +- only the published view is exposed to I/O and offered for RDMA registration; +- one shared `RDMABuf::Inner` owns every per-IB `ibv_mr`, the CUDA device + identity, and the imported mapping owner; +- publication succeeds only when every active IB device accepts the MR. + +Removing a FUSE GPU IOV releases its `IOBuffer`. `RDMABuf::Inner` first +deregisters every MR and only then releases the owner that closes the CUDA IPC +mapping. FUSE shutdown clears GPU IOVs before the process-wide shutdown stops +`IBManager`. The lifetime constraints are: + +```text +IBManager outlives every ibv_mr +CUDA IPC mapping outlives every ibv_mr that covers it +exporter allocation outlives every CUDA IPC mapping +``` + +Explicit unlink and shutdown paths preserve these constraints. An exporter +that exits or frees its allocation before importers close violates the CUDA IPC +contract; 3FS cannot repair that ordering after the fact. + +### GDR URI v2 Contract + +The GDR symlink key and target are: + +```text +{uuid}.gdr.d{device_id} -> +gdr://v2/device/{device}/allocation/{allocation_bytes}/offset/{view_offset}/size/{view_bytes}/ipc/{128_hex_chars} +``` + +Fields: + +- `device`: non-negative CUDA device index. +- `allocation`: byte size of the complete CUDA allocation represented by the + IPC handle. +- `offset`: byte offset of the published view from the allocation base. +- `size`: byte size of the published view and IOV. +- `ipc`: exactly 64 bytes of CUDA IPC handle encoded as 128 hexadecimal + characters. + +The parser is strict: + +- the `gdr://v2` prefix, field names, separators, and field order are exact; +- numeric fields contain ASCII decimal digits only; +- `device` must fit in `int`; +- `allocation`, `offset`, and `size` must fit in `size_t`; +- `allocation` and `size` must be non-zero; +- `offset <= allocation` and `size <= allocation - offset`; +- `ipc` must contain exactly 128 valid hexadecimal characters (upper- or + lower-case is accepted; the formatter emits lower-case); +- trailing or missing data is rejected. + +FUSE also requires the device in the key to equal the URI device. GDR keys +reject block-size and I/O-ring attributes. `hf3fs_iovopen_device()` additionally +requires its requested device and view size to match, and CUDA import verifies +that the imported pointer is the allocation base and that the actual allocation +size matches the URI. + +### CUDA IPC Trust Boundary + +CUDA IPC shares the complete underlying allocation, not an independently +protected subrange. `offset` and `size` constrain the 3FS IOV and the RDMA MR, +but the importing FUSE process maps the full allocation before deriving that +view. + +Treat FUSE and any application importer as trusted with respect to the full +allocation. Do not co-locate unrelated secrets in the same CUDA allocation +when that trust is unacceptable. A PyTorch view remains valid only while its +underlying allocator-owned block is alive and has not been reused. + +### Tagged IOV Table and Ring Ordering + +FUSE stores host and GPU buffers as tagged alternatives in one `IovEntry` +table. Namespace-key and UUID I/O lookup resolve the same entry, and its memory +kind selects the host or GPU view. This keeps access checks, descriptor +identity, and lifetime ownership in one table. + +An I/O-ring batch has one absolute logical order even when its storage wraps +at the physical end of the ring. File lookup, buffer lookup, submission, and +completion preserve that order across both physical spans. + +### Device Memory Semantics + +A CUDA device pointer is not CPU-dereferenceable. Do not use `memcpy`, +`memset`, CPU checksum routines, string/formatting reads, or ordinary C/C++ +loads and stores on `iov->base`. Use CUDA APIs or kernels, or pass the pointer +as an opaque I/O address. + +Before an allocation is exported or directly registered, 3FS validates its +CUDA device and allocation bounds and enables +`CU_POINTER_ATTRIBUTE_SYNC_MEMOPS` for the whole allocation. FUSE repeats this +preparation on its imported CUDA mapping before `ibv_reg_mr`; owner-side and +importer-side CUDA contexts must both use conservative memory-operation +semantics. Address offsets inside 3FS use checked integer arithmetic; +converting a GPU-backed `RDMABuf` to a host `span` or `ByteRange` is a fatal +invariant violation. + +For GPU I/O: + +- inline reads are disabled for a batch containing a GPU buffer; +- inline writes are disabled for GPU buffers; +- client-side CPU checksum calculation/comparison is disabled because the CPU + cannot read the buffer; +- failed GPU MR creation is an error; the host fallback path has been removed. + +Skipping the client CPU comparison does not provide independent end-to-end +verification. Applications that require that property must validate the data +with a GPU-capable mechanism. + +### Sparse Read Holes + +Sparse-read zero filling is memory-kind aware. Host ranges use `memset`; CUDA +device ranges use `cudaMemset` on the owning device. Device zeroing is followed +by `cudaDeviceSynchronize()` before the I/O result is finalized and before the +CQE is published. A CUDA zeroing or synchronization error becomes the I/O +error instead of exposing a completion for unfinished zero fill. + +`HF3FS_IOR_FORBID_READ_HOLES` continues to reject holes instead of filling +them. + +### GPU Write Checksums + +`ChecksumType::NONE` means checksum verification is disabled. It is not an +implicit request for storage-side checksum computation. This remains true when +the per-request verification option is enabled: a configured target type of +`NONE` sends a `NONE` checksum and does not set +`FeatureFlags::SERVER_COMPUTE_CHECKSUM`. + +For a GPU write with checksum verification enabled: + +1. The client sends the configured non-`NONE` checksum type with a placeholder + value and sets `FeatureFlags::SERVER_COMPUTE_CHECKSUM`. +2. Storage completes the RDMA read into its host staging buffer. +3. Storage computes the requested checksum over that host staging data and + replaces the placeholder before dispatching the update to the backend. + +Storage rejects the flag for a non-WRITE update, a `NONE` target type, missing +staged data, or an RDMA-bypass request that supplies neither RDMA nor inline +data. + +There is no capability negotiation for this feature. During a rolling upgrade, +upgrade every storage node before enabling GPU writes from new clients. A new +client can send a placeholder checksum to an old storage node; the old node +does not materialize it and can report a checksum mismatch. Do not enable this +path while an old storage node can be the first node handling the write. + +Backend behavior remains backend-specific: + +- The default `ChunkReplica` path never clears the `ChunkMetadata` object merely + because verification is disabled. A server-computed non-`NONE` checksum + follows the existing verification and metadata-update path. With request + type `NONE`, only the checksum fields become the disabled state (`NONE`, + value `0`); version, state, size, identity, and other metadata continue + through the normal update path. +- `ChunkEngine` uses its `without_checksum` request semantic when verification + is disabled, but continues to maintain and return its CRC32C chunk metadata + for non-remove chunks. Thus request `NONE` does not mean that ChunkEngine + lacks internal checksum metadata. + +## Known Limitations + +- **Per-buffer MR registration is authoritative.** Local CUDA IPC capability + does not guarantee that a specific GPU/view/NIC registration will succeed. +- **`nvidia_peermem` is required by the implemented registration path.** + Device API failures are returned to the caller; no host IOV is substituted. +- **CUDA IPC exposes the whole allocation.** View bounds restrict 3FS and RDMA, + not the CUDA IPC mapping's trust boundary. +- **Device-wide synchronization.** `hf3fs_iovsync()` and sparse device zeroing + use `cudaDeviceSynchronize()`, not stream-scoped fencing. +- **No GDR block partitioning.** Device IOV APIs require `block_size == 0`. +- **Publisher-coordinated lifetime.** There is no distributed reference count; + all I/O must complete and the exporter must outlive FUSE use and every + importer. +- **All-HCA registration.** Partial-HCA GDR is not supported. A GPU IOV is + rejected if any active IB device cannot register it; transport selection is + therefore never given a sparse rkey set. +- **IPC-import MR support is platform-dependent.** The target driver, + `nvidia_peermem`, HCA, and CUDA combination must pass the two-process + `IpcImportedPointerRegistersWithEveryHca` hardware test. +- **No registration cache.** An IOV is registered once for its published + lifetime. Add a 64KB-normalized, allocation-ID-aware cache only if measured + registration latency requires it. + +## Target-Machine Validation + +Run the GDR test binary on the NVIDIA/RDMA host used for deployment: + +```bash +ctest --test-dir build -R test_gdr --output-on-failure +``` + +Do not accept a skipped +`TestCudaRDMABuf.IpcImportedPointerRegistersWithEveryHca` test as GDR +validation. That two-process test proves that the exporter can prepare and +publish a CUDA allocation, the importer can enable `SYNC_MEMOPS`, and every +active HCA can register the imported view. The deployment acceptance test must +also submit at least one real storage read and write through a mounted GDR IOV; +the unit test does not emulate an HCA data transfer. + +## References + +- [NVIDIA GPUDirect RDMA: synchronization, registration, and memory ordering](https://docs.nvidia.com/cuda/gpudirect-rdma/) +- [NVIDIA CUDA Driver API: pointer attributes and `SYNC_MEMOPS`](https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__UNIFIED.html) diff --git a/src/client/storage/CMakeLists.txt b/src/client/storage/CMakeLists.txt index 97dc0170..eb811da6 100644 --- a/src/client/storage/CMakeLists.txt +++ b/src/client/storage/CMakeLists.txt @@ -1 +1,5 @@ target_add_lib(storage-client storage-fbs common mgmtd-client) + +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(storage-client SCOPE PRIVATE) +endif() diff --git a/src/client/storage/StorageClient.cc b/src/client/storage/StorageClient.cc index 0763d23a..18cae39d 100644 --- a/src/client/storage/StorageClient.cc +++ b/src/client/storage/StorageClient.cc @@ -1,7 +1,13 @@ #include +#include + +#ifdef HF3FS_ENABLE_GDR +#include +#endif #include "StorageClientImpl.h" #include "StorageClientInMem.h" +#include "common/cuda/CudaMemory.h" #include "common/monitor/ScopedMetricsWriter.h" #include "common/net/ib/RDMABuf.h" @@ -14,6 +20,31 @@ static monitor::DistributionRecorder iobuf_reg_size{"storage_client.iobuf_reg.si const StorageClient::Config StorageClient::kDefaultConfig; +#ifdef HF3FS_ENABLE_GDR +Result detail::zeroCudaDeviceRange(void *devicePtr, size_t length, int deviceId) { + auto guard = cuda::ScopedDevice::create(deviceId); + if (!guard) { + return makeError( + StorageClientCode::kRemoteIOError, + fmt::format("failed to select CUDA device {} while zeroing device memory: {}", deviceId, guard.error())); + } + + auto cudaResult = cudaMemset(devicePtr, 0, length); + if (cudaResult != cudaSuccess) { + return makeError(StorageClientCode::kRemoteIOError, + fmt::format("cudaMemset failed while zeroing device memory: {}", cudaGetErrorString(cudaResult))); + } + + cudaResult = cudaDeviceSynchronize(); + if (cudaResult != cudaSuccess) { + return makeError( + StorageClientCode::kRemoteIOError, + fmt::format("cudaDeviceSynchronize failed after zeroing device memory: {}", cudaGetErrorString(cudaResult))); + } + return Void{}; +} +#endif + std::shared_ptr StorageClient::create(ClientId clientId, const Config &config, hf3fs::client::ICommonMgmtdClient &mgmtdClient) { @@ -94,6 +125,57 @@ TruncateChunkOp StorageClient::createTruncateOp(ChainId chainId, return TruncateChunkOp(requestId, chainId, chunkId, chunkLen, chunkSize, onlyExtendChunk, userCtx); } +Result IOBuffer::zeroRange(size_t offset, size_t length) const { + if (!rdmabuf_.valid()) { + return makeError(StorageClientCode::kInvalidArg, "cannot zero an invalid IOBuffer"); + } + if (offset > size() || length > size() - offset) { + return makeError( + StorageClientCode::kInvalidArg, + fmt::format("IOBuffer zero range offset {} length {} exceeds buffer size {}", offset, length, size())); + } + if (length == 0) { + return Void{}; + } + + auto *ptr = dataAtOffset(offset); + if (ptr == nullptr) { + return makeError(StorageClientCode::kInvalidArg, "IOBuffer zero range has an invalid address"); + } + + if (!isDeviceMemory()) { + std::memset(ptr, 0, length); + return Void{}; + } + +#ifdef HF3FS_ENABLE_GDR + return detail::zeroCudaDeviceRange(ptr, length, cudaDeviceId()); +#else + return makeError(StorageClientCode::kRemoteIOError, + "cannot zero a device IOBuffer because CUDA/GDR support is disabled"); +#endif +} + +Result prepareWriteChecksum(ChecksumType targetType, + bool verifyChecksum, + bool deviceMemory, + const uint8_t *data, + size_t length) { + if (!verifyChecksum) { + return PreparedWriteChecksum{ChecksumInfo{ChecksumType::NONE, 0}, false}; + } + if (targetType == ChecksumType::NONE) { + return PreparedWriteChecksum{ChecksumInfo{ChecksumType::NONE, 0}, false}; + } + if (deviceMemory) { + return PreparedWriteChecksum{ChecksumInfo{targetType, 0}, true}; + } + if (data == nullptr && length != 0) { + return makeError(StorageClientCode::kInvalidArg, "cannot checksum a null host buffer"); + } + return PreparedWriteChecksum{ChecksumInfo::create(targetType, data, length), false}; +} + Result StorageClient::registerIOBuffer(uint8_t *buf, size_t len) { monitor::ScopedLatencyWriter latencyWriter(iobuf_reg_latency); iobuf_reg_size.addSample(len); @@ -109,4 +191,38 @@ Result StorageClient::registerIOBuffer(uint8_t *buf, size_t len) { } } +Result StorageClient::registerGpuIOBuffer(uint8_t *gpuPtr, size_t len) { + (void)gpuPtr; + (void)len; + return makeError(StorageClientCode::kNotAvailable, "GPU IOBuffer registration requires an RDMA storage client"); +} + +Result StorageClientImpl::registerGpuIOBuffer(uint8_t *gpuPtr, size_t len) { + monitor::ScopedLatencyWriter latencyWriter(iobuf_reg_latency); + iobuf_reg_size.addSample(len); + + if (gpuPtr == nullptr || len == 0) { + iobuf_reg_failed_ops.addSample(1); + return makeError(StorageClientCode::kInvalidArg, "GPU IOBuffer pointer and length must be non-zero"); + } + +#ifdef HF3FS_ENABLE_GDR + auto gpuBuf = hf3fs::net::RDMABuf::createFromCudaBuffer(gpuPtr, len, -1); + if (gpuBuf) { + iobuf_reg_success_ops.addSample(1); + return IOBuffer{std::move(*gpuBuf)}; + } + + iobuf_reg_failed_ops.addSample(1); + auto code = gpuBuf.error().code() == StatusCode::kInvalidArg ? StorageClientCode::kInvalidArg + : StorageClientCode::kMemoryError; + return makeError( + code, + fmt::format("failed to register CUDA device pointer {} for RDMA: {}", fmt::ptr(gpuPtr), gpuBuf.error())); +#else + iobuf_reg_failed_ops.addSample(1); + return makeError(StorageClientCode::kNotAvailable, "GPU IOBuffer registration requires CUDA/GDR support"); +#endif +} + } // namespace hf3fs::storage::client diff --git a/src/client/storage/StorageClient.h b/src/client/storage/StorageClient.h index 5f34e850..22796a6a 100644 --- a/src/client/storage/StorageClient.h +++ b/src/client/storage/StorageClient.h @@ -1,5 +1,8 @@ #pragma once +#include +#include +#include #include #include @@ -7,6 +10,7 @@ #include "UpdateChannelAllocator.h" #include "client/mgmtd/ICommonMgmtdClient.h" #include "common/net/Client.h" +#include "common/net/ib/RDMABuf.h" #include "common/utils/Address.h" #include "common/utils/Coroutine.h" #include "common/utils/Result.h" @@ -45,19 +49,63 @@ class RoutingTarget { class IOBuffer : public folly::MoveOnly { public: - uint8_t *data() const { return const_cast(rdmabuf.ptr()); } + /** + * Get the raw data pointer. + * WARNING: For GPU buffers, this returns a device address that is NOT + * CPU-dereferenceable. Use subrangeRemote() for RDMA operations. + */ + uint8_t *data() const { return const_cast(rdmabuf_.ptr()); } - size_t size() const { return rdmabuf.size(); } + size_t size() const { return rdmabuf_.size(); } - bool contains(const uint8_t *data, uint32_t len) const { return rdmabuf.contains(data, len); } + bool contains(const uint8_t *data, uint32_t len) const { return rdmabuf_.contains(data, len); } - net::RDMABuf subrange(size_t offset, size_t length) const { return rdmabuf.subrange(offset, length); } + std::optional offsetOf(const uint8_t *data, uint32_t len) const { + if (!contains(data, len)) { + return std::nullopt; + } + return reinterpret_cast(data) - reinterpret_cast(rdmabuf_.ptr()); + } + + uint8_t *dataAtOffset(size_t offset) const { + auto *basePtr = rdmabuf_.ptr(); + auto size = rdmabuf_.size(); + auto base = reinterpret_cast(basePtr); + if (!basePtr || offset > size || base > std::numeric_limits::max() - offset) { + return nullptr; + } + return reinterpret_cast(base + offset); + } + + /** + * Get an RDMARemoteBuf for a subrange — works for both host and GPU. + * For GPU buffers, the returned RDMARemoteBuf contains the device virtual + * address and per-HCA rkeys, suitable for direct RDMA read/write without + * CPU copies. + */ + net::RDMARemoteBuf subrangeRemote(size_t offset, size_t length) const { + return rdmabuf_.toRemoteBuf().subrange(offset, length); + } IOBuffer(hf3fs::net::RDMABuf rdmabuf) - : rdmabuf(std::move(rdmabuf)) {} + : rdmabuf_(std::move(rdmabuf)) {} + + /** Device buffers must not be accessed by CPU memcpy, memset, or checksum code. */ + bool isDeviceMemory() const { return rdmabuf_.isDeviceMemory(); } + + /** Returns the CUDA device ID for GPU buffers, or -1 for host buffers. */ + int cudaDeviceId() const { return rdmabuf_.cudaDeviceId().value_or(-1); } + + /** + * Zero a byte range relative to data(). + * + * Host memory is zeroed synchronously with memset. Device memory is zeroed + * with CUDA and synchronized before this method returns. + */ + Result zeroRange(size_t offset, size_t length) const; private: - const hf3fs::net::RDMABuf rdmabuf; + net::RDMABuf rdmabuf_; friend class IOBase; friend class StorageClient; @@ -65,6 +113,25 @@ class IOBuffer : public folly::MoveOnly { friend class StorageClientInMem; }; +struct PreparedWriteChecksum { + ChecksumInfo checksum; + bool computeOnServer = false; +}; + +Result prepareWriteChecksum(ChecksumType targetType, + bool verifyChecksum, + bool deviceMemory, + const uint8_t *data, + size_t length); + +#ifdef HF3FS_ENABLE_GDR +namespace detail { + +Result zeroCudaDeviceRange(void *devicePtr, size_t length, int deviceId); + +} // namespace detail +#endif + class IOBase : public folly::MoveOnly { private: IOBase(ChainId chainId, @@ -89,7 +156,16 @@ class IOBase : public folly::MoveOnly { status_code_t statusCode() const { return hf3fs::getStatusCode(result.lengthInfo); } uint32_t resultLen() const { return result.lengthInfo ? *result.lengthInfo : 0; } uint32_t dataLen() const { return length; } - uint8_t *dataEnd() const { return data + length; } + uint8_t *dataEnd() const { + if (!data) { + return nullptr; + } + auto addr = reinterpret_cast(data); + if (addr > std::numeric_limits::max() - length) { + return nullptr; + } + return reinterpret_cast(addr + length); + } ChunkIdRange chunkRange() const { return {chunkId, chunkId, 1}; } uint32_t numProcessedChunks() const { return bool(result.lengthInfo); } void resetResult() { result = IOResult{}; } @@ -516,6 +592,13 @@ class StorageClient : public folly::MoveOnly { // delete the returned IOBuffer object to deregister the buffer virtual Result registerIOBuffer(uint8_t *buf, size_t len); + // Register a GPU memory buffer for RDMA. Non-RDMA implementations reject it. + // The returned IOBuffer has isDeviceMemory() == true, which prevents + // CPU operations (inline memcpy, checksum) on the buffer. + // The caller must keep the CUDA allocation alive until the IOBuffer and all + // I/O operations using it have been destroyed. + virtual Result registerGpuIOBuffer(uint8_t *gpuPtr, size_t len); + virtual CoTryTask batchRead(std::span readIOs, const flat::UserInfo &userInfo, const ReadOptions &options = ReadOptions(), diff --git a/src/client/storage/StorageClientImpl.cc b/src/client/storage/StorageClientImpl.cc index be64ddaf..9a493847 100644 --- a/src/client/storage/StorageClientImpl.cc +++ b/src/client/storage/StorageClientImpl.cc @@ -1,10 +1,12 @@ #include "StorageClientImpl.h" #include +#include #include #include #include #include +#include #include #include @@ -653,6 +655,14 @@ uint32_t buildFeatureFlagsFromOptions(const DebugOptions &debugOptions) { return featureFlags; } +uint32_t buildWriteFeatureFlagsFromOptions(const DebugOptions &debugOptions, bool computeChecksumOnServer) { + auto featureFlags = buildFeatureFlagsFromOptions(debugOptions); + if (computeChecksumOnServer) { + BITFLAGS_SET(featureFlags, hf3fs::storage::FeatureFlags::SERVER_COMPUTE_CHECKSUM); + } + return featureFlags; +} + template BatchReq buildBatchRequest(const ClientRequestContext &requestCtx, const ClientId &clientId, @@ -678,6 +688,7 @@ typename hf3fs::storage::BatchReadReq buildBatchRequest(const ClientRequestConte std::vector payloads; payloads.reserve(ops.size()); size_t requestedBytes = 0; + bool hasGpuBuffer = false; // Track if any IO uses GPU buffer auto requestTagSet = monitor::instanceTagSet("batchRead"); auto tagged_bytes_per_operation = bytes_per_operation.getRecorderWithTag(requestTagSet); @@ -687,14 +698,19 @@ typename hf3fs::storage::BatchReadReq buildBatchRequest(const ClientRequestConte for (auto &op : ops) { hf3fs::storage::GlobalKey key{op->routingTarget.getVersionedChainId(), op->chunkId}; - size_t offset = op->data - op->buffer->data(); - auto iobuf = op->buffer->subrange(offset, op->length); + auto offset = op->buffer->offsetOf(op->data, op->length); + XLOGF_IF(DFATAL, !offset, "Read IO data range was not validated before request build"); + auto remoteBuf = op->buffer->subrangeRemote(offset.value_or(0), op->length); requestedBytes += op->length; tagged_bytes_per_operation->addSample(op->length); + if (op->buffer && op->buffer->isDeviceMemory()) { + hasGpuBuffer = true; + } + op->requestId = requestId; - payloads.push_back({op->offset, op->length, std::move(key), iobuf.toRemoteBuf()}); + payloads.push_back({op->offset, op->length, std::move(key), remoteBuf}); } bytes_per_request.addSample(requestedBytes, requestTagSet); @@ -703,7 +719,8 @@ typename hf3fs::storage::BatchReadReq buildBatchRequest(const ClientRequestConte auto checksumType = options.verifyChecksum() ? config.chunk_checksum_type() : ChecksumType::NONE; uint32_t featureFlags = buildFeatureFlagsFromOptions(options.debug()); - if (requestedBytes < requestCtx.clientConfig.max_inline_read_bytes()) { + // Do not use inline data for GPU buffers - cannot memcpy to GPU memory + if (requestedBytes < requestCtx.clientConfig.max_inline_read_bytes() && !hasGpuBuffer) { BITFLAGS_SET(featureFlags, hf3fs::storage::FeatureFlags::SEND_DATA_INLINE); } @@ -889,9 +906,12 @@ CoTryTask StorageClientImpl::callMessengerMethod(StorageMessenger &messenge template std::vector validateDataRange(const std::vector &ios, bool checkOverlappingBuffers) { std::vector sortedIOs(begin(ios), end(ios)); + auto addrOf = [](const uint8_t *ptr) { return reinterpret_cast(ptr); }; if (checkOverlappingBuffers) { - std::sort(begin(sortedIOs), end(sortedIOs), [](const IO *a, const IO *b) { return a->data < b->data; }); + std::sort(begin(sortedIOs), end(sortedIOs), [addrOf](const IO *a, const IO *b) { + return addrOf(a->data) < addrOf(b->data); + }); } std::vector validIOs; @@ -928,7 +948,17 @@ std::vector validateDataRange(const std::vector &ios, bool checkOver return {}; } - if (checkOverlappingBuffers && lastIO != nullptr && io->data < lastIO->dataEnd()) { + if (checkOverlappingBuffers && lastIO != nullptr) { + auto lastStart = addrOf(lastIO->data); + auto lastEnd = lastStart > std::numeric_limits::max() - lastIO->length + ? std::numeric_limits::max() + : lastStart + lastIO->length; + if (addrOf(io->data) >= lastEnd) { + validIOs.push_back(io); + lastIO = io; + continue; + } + XLOGF(ERR, "#{}/{} User data overlaps: current data starts at {}, length {}; " "last data starts at {}, ends at {}, length {}", @@ -937,7 +967,7 @@ std::vector validateDataRange(const std::vector &ios, bool checkOver fmt::ptr(io->data), io->length, fmt::ptr(lastIO->data), - fmt::ptr(lastIO->dataEnd()), + fmt::ptr(reinterpret_cast(lastEnd)), lastIO->length); setErrorCodeOfOps(sortedIOs, StorageClientCode::kInvalidArg); return {}; @@ -965,11 +995,22 @@ std::vector splitReadIOs(StorageClientImpl &client, uint32_t ioLen = std::min(ioEnd - offset, length); assert(ioLen > 0); + auto parentDataOffset = parentIO->buffer->offsetOf(parentIO->data, parentIO->length); + XLOGF_IF(DFATAL, !parentDataOffset, "Parent read IO data range was not validated before split"); + auto childDelta = static_cast(offset - parentIO->offset); + auto childOffset = parentDataOffset.value_or(0); + XLOGF_IF(DFATAL, + childDelta > std::numeric_limits::max() - childOffset, + "Split read IO data offset overflow"); + childOffset = childDelta > std::numeric_limits::max() - childOffset ? std::numeric_limits::max() + : childOffset + childDelta; + auto childData = parentIO->buffer->dataAtOffset(childOffset); + parentIO->splittedIOs.push_back(client.createReadIO(parentIO->routingTarget.chainId, parentIO->chunkId, offset, ioLen, - parentIO->data + (offset - parentIO->offset), + childData, parentIO->buffer)); XLOGF(DBG5, @@ -1583,7 +1624,7 @@ CoTryTask StorageClientImpl::batchReadWithRetry(ClientRequestContext &requ const bool splitLargeIOs = maxIOBytes > 0; std::vector splittedIOs; - auto sendOps = [ this, &requestCtx, &userInfo, &options ](const std::vector &ops) -> auto{ + auto sendOps = [this, &requestCtx, &userInfo, &options](const std::vector &ops) -> auto { return batchReadWithoutRetry(requestCtx, ops, userInfo, options); }; @@ -1712,6 +1753,13 @@ CoTryTask StorageClientImpl::batchReadWithoutRetry(ClientRequestContext &r auto inlinebuf = &response->inlinebuf.data[0]; for (auto readIO : batchIOs) { + // Skip CPU memcpy for GPU memory buffers - this should not happen + // because inline data is disabled for GPU buffers, but check anyway + if (readIO->buffer && readIO->buffer->isDeviceMemory()) { + XLOGF(ERR, "BUG: Inline data received for GPU buffer - cannot memcpy to GPU memory"); + setErrorCodeOfOp(readIO, StorageClientCode::kInvalidArg); + continue; + } std::memcpy(readIO->data, inlinebuf, readIO->resultLen()); inlinebuf += readIO->resultLen(); } @@ -1720,6 +1768,11 @@ CoTryTask StorageClientImpl::batchReadWithoutRetry(ClientRequestContext &r if (options.verifyChecksum()) { for (auto readIO : batchIOs) { if (readIO->result.lengthInfo && *readIO->result.lengthInfo > 0) { + // Skip CPU checksum verification for GPU memory - server checksum is trusted + if (readIO->buffer && readIO->buffer->isDeviceMemory()) { + XLOGF(DBG, "Skipping CPU checksum verification for GPU buffer"); + continue; + } auto checksum = ChecksumInfo::create(readIO->result.checksum.type, readIO->data, *readIO->result.lengthInfo); if (FAULT_INJECTION_POINT(requestCtx.debugFlags.injectClientError(), true, @@ -1773,7 +1826,7 @@ CoTryTask StorageClientImpl::batchWriteWithRetry(ClientRequestContext &req const flat::UserInfo &userInfo, const WriteOptions &options, std::vector &failedIOs) { - auto sendOps = [ this, &requestCtx, userInfo, options ](const std::vector &ops) -> auto{ + auto sendOps = [this, &requestCtx, userInfo, options](const std::vector &ops) -> auto { return batchWriteWithoutRetry(requestCtx, ops, userInfo, options); }; @@ -1868,34 +1921,51 @@ CoTryTask StorageClientImpl::sendWriteRequest(ClientRequestContext &reques hf3fs::storage::ChainVer(writeIO->routingTarget.chainVer)}; hf3fs::storage::GlobalKey key{vChainId, hf3fs::storage::ChunkId(writeIO->chunkId)}; - size_t offset = writeIO->data - writeIO->buffer->data(); - auto iobuf = writeIO->buffer->subrange(offset, writeIO->length); + auto offset = writeIO->buffer->offsetOf(writeIO->data, writeIO->length); + XLOGF_IF(DFATAL, !offset, "Write IO data range was not validated before request build"); + auto remoteBuf = writeIO->buffer->subrangeRemote(offset.value_or(0), writeIO->length); bytes_per_operation.addSample(writeIO->length, requestCtx.requestTagSet); bytes_per_request.addSample(writeIO->length, requestCtx.requestTagSet); ops_per_request.addSample(1, requestCtx.requestTagSet); - if (options.verifyChecksum()) { + // Check if buffer is GPU memory - skip CPU operations if so + bool isGpuBuffer = writeIO->buffer && writeIO->buffer->isDeviceMemory(); + + auto preparedChecksum = prepareWriteChecksum(config_.chunk_checksum_type(), + options.verifyChecksum(), + isGpuBuffer, + writeIO->data, + writeIO->length); + if (UNLIKELY(!preparedChecksum)) { + XLOGF(ERR, "Failed to prepare write checksum: {}", preparedChecksum.error()); + setErrorCodeOfOp(writeIO, preparedChecksum.error().code()); + co_return Void{}; + } + writeIO->checksum = preparedChecksum->checksum; + + if (options.verifyChecksum() && !isGpuBuffer && writeIO->length != 0) { writeIO->checksum = FAULT_INJECTION_POINT( requestCtx.debugFlags.injectClientError(), ChecksumInfo::create(config_.chunk_checksum_type(), writeIO->data, std::max(writeIO->length / 2, 1U)), - ChecksumInfo::create(config_.chunk_checksum_type(), writeIO->data, writeIO->length)); + writeIO->checksum); } hf3fs::storage::RequestId requestId(writeIO->requestId); hf3fs::storage::MessageTag tag{clientId_, requestId, writeIO->routingTarget.channel}; - uint32_t featureFlags = buildFeatureFlagsFromOptions(options.debug()); + uint32_t featureFlags = buildWriteFeatureFlagsFromOptions(options.debug(), preparedChecksum->computeOnServer); hf3fs::storage::UpdateIO payload{writeIO->offset, writeIO->length, writeIO->chunkSize, key, - iobuf.toRemoteBuf(), + remoteBuf, hf3fs::storage::ChunkVer(0) /*updateVer*/, UpdateType::WRITE, writeIO->checksum}; - if (writeIO->length <= requestCtx.clientConfig.max_inline_write_bytes()) { + // Do not use inline data for GPU buffers - cannot memcpy from GPU memory + if (writeIO->length <= requestCtx.clientConfig.max_inline_write_bytes() && !isGpuBuffer) { payload.inlinebuf.data.assign(writeIO->data, writeIO->data + writeIO->length); BITFLAGS_SET(featureFlags, hf3fs::storage::FeatureFlags::SEND_DATA_INLINE); } @@ -2035,7 +2105,7 @@ CoTryTask StorageClientImpl::queryLastChunk(std::span op std::vector failedIOVec; if (failedOps == nullptr) failedOps = &failedIOVec; - auto sendOps = [ this, &requestCtx, userInfo, options ](const std::vector &ops) -> auto{ + auto sendOps = [this, &requestCtx, userInfo, options](const std::vector &ops) -> auto { return queryLastChunkWithoutRetry(requestCtx, ops, userInfo, options); }; @@ -2138,7 +2208,7 @@ CoTryTask StorageClientImpl::removeChunks(std::span ops, std::vector failedIOVec; if (failedOps == nullptr) failedOps = &failedIOVec; - auto sendOps = [ this, &requestCtx, userInfo, options ](const std::vector &ops) -> auto{ + auto sendOps = [this, &requestCtx, userInfo, options](const std::vector &ops) -> auto { return removeChunksWithoutRetry(requestCtx, ops, userInfo, options); }; @@ -2269,7 +2339,7 @@ CoTryTask StorageClientImpl::truncateChunks(std::span ops std::vector failedIOVec; if (failedOps == nullptr) failedOps = &failedIOVec; - auto sendOps = [ this, &requestCtx, userInfo, options ](const std::vector &ops) -> auto{ + auto sendOps = [this, &requestCtx, userInfo, options](const std::vector &ops) -> auto { return truncateChunksWithoutRetry(requestCtx, ops, userInfo, options); }; diff --git a/src/client/storage/StorageClientImpl.h b/src/client/storage/StorageClientImpl.h index 6fa9e02c..195d60a0 100644 --- a/src/client/storage/StorageClientImpl.h +++ b/src/client/storage/StorageClientImpl.h @@ -16,6 +16,8 @@ namespace hf3fs::storage::client { class ClientRequestContext; +uint32_t buildWriteFeatureFlagsFromOptions(const DebugOptions &debugOptions, bool computeChecksumOnServer); + class StorageClientImpl : public StorageClient { public: StorageClientImpl(const ClientId &clientId, const Config &config, hf3fs::client::ICommonMgmtdClient &mgmtdClient); @@ -27,6 +29,8 @@ class StorageClientImpl : public StorageClient { hf3fs::client::ICommonMgmtdClient &getMgmtdClient() override { return mgmtdClient_; } void stop() override; + Result registerGpuIOBuffer(uint8_t *gpuPtr, size_t len) override; + CoTryTask batchRead(std::span readIOs, const flat::UserInfo &userInfo, const ReadOptions &options = ReadOptions(), diff --git a/src/common/CMakeLists.txt b/src/common/CMakeLists.txt index 39b7d6e9..262e3fee 100644 --- a/src/common/CMakeLists.txt +++ b/src/common/CMakeLists.txt @@ -7,5 +7,15 @@ target_add_lib(common memory-common version-info folly ibverbs scn::scn clickhou add_dependencies(common MonitorCollectorService-fbs) target_sources(common PRIVATE utils/Linenoise.c) +# Add CUDA/GDR support if enabled +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(common SCOPE PRIVATE) +endif() + target_add_shared_lib(hf3fs_common_shared memory-common version-info folly ibverbs scn::scn clickhouse-cpp-lib-static toml11 libzstd_static 3fs_liburing) add_dependencies(hf3fs_common_shared MonitorCollectorService-fbs) + +# Add CUDA/GDR support to shared library if enabled +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(hf3fs_common_shared SCOPE PRIVATE) +endif() diff --git a/src/common/cuda/CudaMemory.cc b/src/common/cuda/CudaMemory.cc new file mode 100644 index 00000000..6d8f4da0 --- /dev/null +++ b/src/common/cuda/CudaMemory.cc @@ -0,0 +1,254 @@ +#include "common/cuda/CudaMemory.h" + +#include +#include +#include +#include +#include + +#ifdef HF3FS_ENABLE_GDR +#include +#include +#endif + +namespace hf3fs::cuda { +namespace { + +#ifdef HF3FS_ENABLE_GDR +std::string driverError(CUresult result) { + const char *name = nullptr; + const char *message = nullptr; + cuGetErrorName(result, &name); + cuGetErrorString(result, &message); + return fmt::format("{}: {}", name ? name : "CUDA_ERROR_UNKNOWN", message ? message : "unknown error"); +} +#endif + +} // namespace + +ScopedDevice::ScopedDevice(ScopedDevice &&other) noexcept + : previousDevice_(std::exchange(other.previousDevice_, -1)), + restore_(std::exchange(other.restore_, false)) {} + +ScopedDevice &ScopedDevice::operator=(ScopedDevice &&other) noexcept { + if (this != &other) { + restore(); + previousDevice_ = std::exchange(other.previousDevice_, -1); + restore_ = std::exchange(other.restore_, false); + } + return *this; +} + +ScopedDevice::~ScopedDevice() { restore(); } + +void ScopedDevice::restore() noexcept { +#ifdef HF3FS_ENABLE_GDR + if (restore_) { + restore_ = false; + auto error = cudaSetDevice(previousDevice_); + XLOGF_IF(WARN, + error != cudaSuccess, + "cudaSetDevice({}) failed while restoring CUDA device: {}", + previousDevice_, + cudaGetErrorString(error)); + } +#endif +} + +Result ScopedDevice::create(int deviceId) { +#ifdef HF3FS_ENABLE_GDR + if (deviceId < 0) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, "invalid CUDA device {}", deviceId); + } + + int previousDevice = -1; + auto error = cudaGetDevice(&previousDevice); + if (error != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, "cudaGetDevice failed: {}", cudaGetErrorString(error)); + } + if (previousDevice == deviceId) { + return ScopedDevice(previousDevice, false); + } + + error = cudaSetDevice(deviceId); + if (error != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, "cudaSetDevice({}) failed: {}", deviceId, cudaGetErrorString(error)); + } + return ScopedDevice(previousDevice, true); +#else + (void)deviceId; + return makeError(StatusCode::kNotImplemented, "CUDA/GDR is disabled in this build"); +#endif +} + +Result deviceCount() { +#ifdef HF3FS_ENABLE_GDR + int count = 0; + auto error = cudaGetDeviceCount(&count); + if (error == cudaErrorNoDevice) { + cudaGetLastError(); + return 0; + } + if (error != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, "cudaGetDeviceCount failed: {}", cudaGetErrorString(error)); + } + return count; +#else + return 0; +#endif +} + +Result supportsIpc(int deviceId) { +#ifdef HF3FS_ENABLE_GDR + auto count = deviceCount(); + RETURN_ON_ERROR(count); + if (deviceId < 0 || deviceId >= *count) { + return false; + } + + int unifiedAddressing = 0; + auto error = cudaDeviceGetAttribute(&unifiedAddressing, cudaDevAttrUnifiedAddressing, deviceId); + if (error != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, + "failed to query unified addressing for CUDA device {}: {}", + deviceId, + cudaGetErrorString(error)); + } + + int computeMode = cudaComputeModeProhibited; + error = cudaDeviceGetAttribute(&computeMode, cudaDevAttrComputeMode, deviceId); + if (error != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, + "failed to query compute mode for CUDA device {}: {}", + deviceId, + cudaGetErrorString(error)); + } + + return unifiedAddressing != 0 && computeMode != cudaComputeModeProhibited; +#else + (void)deviceId; + return false; +#endif +} + +Result inspectDeviceMemory(void *devicePtr, size_t viewSize, int expectedDeviceId) { +#ifdef HF3FS_ENABLE_GDR + if (!devicePtr || viewSize == 0 || expectedDeviceId < -1) { + return makeError(StatusCode::kInvalidArg, "invalid CUDA device-memory view"); + } + + int deviceId = expectedDeviceId; + if (deviceId < 0) { + cudaPointerAttributes attributes; + auto runtimeResult = cudaPointerGetAttributes(&attributes, devicePtr); + if (runtimeResult != cudaSuccess) { + auto message = std::string(cudaGetErrorString(runtimeResult)); + cudaGetLastError(); + return MAKE_ERROR_F(StatusCode::kInvalidArg, "failed to inspect CUDA pointer: {}", message); + } + if (attributes.type != cudaMemoryTypeDevice || attributes.device < 0) { + return makeError(StatusCode::kInvalidArg, "pointer is not CUDA device memory"); + } + deviceId = attributes.device; + } + + auto guard = ScopedDevice::create(deviceId); + RETURN_ON_ERROR(guard); + + auto address = reinterpret_cast(devicePtr); + CUmemorytype memoryType = CU_MEMORYTYPE_HOST; + auto driverResult = cuPointerGetAttribute(&memoryType, CU_POINTER_ATTRIBUTE_MEMORY_TYPE, address); + if (driverResult != CUDA_SUCCESS || memoryType != CU_MEMORYTYPE_DEVICE) { + if (driverResult != CUDA_SUCCESS) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "failed to inspect CUDA pointer {}: {}", + fmt::ptr(devicePtr), + driverError(driverResult)); + } + return MAKE_ERROR_F(StatusCode::kInvalidArg, "pointer {} is not CUDA device memory", fmt::ptr(devicePtr)); + } + + CUdevice pointerDevice = -1; + driverResult = cuPointerGetAttribute(&pointerDevice, CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL, address); + if (driverResult != CUDA_SUCCESS) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, "failed to query CUDA pointer device: {}", driverError(driverResult)); + } + if (pointerDevice != deviceId) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "CUDA pointer belongs to device {}, expected {}", + pointerDevice, + deviceId); + } + + CUdeviceptr allocationBase = 0; + size_t allocationSize = 0; + driverResult = cuMemGetAddressRange(&allocationBase, &allocationSize, address); + if (driverResult != CUDA_SUCCESS) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "cuMemGetAddressRange failed for CUDA pointer: {}", + driverError(driverResult)); + } + if (address < allocationBase) { + return makeError(StatusCode::kInvalidArg, "CUDA allocation range does not contain the pointer"); + } + + auto offset = static_cast(address - allocationBase); + if (allocationSize == 0 || offset > allocationSize || viewSize > allocationSize - offset) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "CUDA view [{}, {}) exceeds allocation size {}", + offset, + offset + viewSize, + allocationSize); + } + + return DeviceMemoryView{ + .allocationBase = reinterpret_cast(allocationBase), + .allocationSize = allocationSize, + .offset = offset, + .deviceId = pointerDevice, + }; +#else + (void)devicePtr; + (void)viewSize; + (void)expectedDeviceId; + return makeError(StatusCode::kNotImplemented, "CUDA/GDR is disabled in this build"); +#endif +} + +Result prepareGdrMemory(void *devicePtr, size_t viewSize, int expectedDeviceId) { +#ifdef HF3FS_ENABLE_GDR + auto view = inspectDeviceMemory(devicePtr, viewSize, expectedDeviceId); + RETURN_ON_ERROR(view); + + auto guard = ScopedDevice::create(view->deviceId); + RETURN_ON_ERROR(guard); + + unsigned int enabled = 0; + auto allocationBase = reinterpret_cast(view->allocationBase); + auto driverResult = cuPointerGetAttribute(&enabled, CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, allocationBase); + if (driverResult != CUDA_SUCCESS) { + return MAKE_ERROR_F(StatusCode::kIOError, + "failed to query CU_POINTER_ATTRIBUTE_SYNC_MEMOPS: {}", + driverError(driverResult)); + } + if (enabled != 0) { + return *view; + } + + enabled = 1; + driverResult = cuPointerSetAttribute(&enabled, CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, allocationBase); + if (driverResult != CUDA_SUCCESS) { + return MAKE_ERROR_F(StatusCode::kIOError, + "failed to enable CU_POINTER_ATTRIBUTE_SYNC_MEMOPS: {}", + driverError(driverResult)); + } + return *view; +#else + (void)devicePtr; + (void)viewSize; + (void)expectedDeviceId; + return makeError(StatusCode::kNotImplemented, "CUDA/GDR is disabled in this build"); +#endif +} + +} // namespace hf3fs::cuda diff --git a/src/common/cuda/CudaMemory.h b/src/common/cuda/CudaMemory.h new file mode 100644 index 00000000..9a5a2aea --- /dev/null +++ b/src/common/cuda/CudaMemory.h @@ -0,0 +1,50 @@ +#pragma once + +#include + +#include "common/utils/Result.h" + +namespace hf3fs::cuda { + +struct DeviceMemoryView { + void *allocationBase = nullptr; + size_t allocationSize = 0; + size_t offset = 0; + int deviceId = -1; +}; + +class ScopedDevice { + public: + ScopedDevice(const ScopedDevice &) = delete; + ScopedDevice &operator=(const ScopedDevice &) = delete; + ScopedDevice(ScopedDevice &&other) noexcept; + ScopedDevice &operator=(ScopedDevice &&other) noexcept; + ~ScopedDevice(); + + static Result create(int deviceId); + + private: + explicit ScopedDevice(int previousDevice, bool restore) + : previousDevice_(previousDevice), + restore_(restore) {} + + void restore() noexcept; + + int previousDevice_ = -1; + bool restore_ = false; +}; + +Result deviceCount(); +/** Local CUDA IPC prerequisites only; this does not probe an HCA or ibv_reg_mr. */ +Result supportsIpc(int deviceId); + +/** Validate a CUDA device-memory view; pass -1 to discover its device. */ +Result inspectDeviceMemory(void *devicePtr, size_t viewSize, int expectedDeviceId); + +/** + * Validate a CUDA device-memory view and enable the allocation-wide synchronization + * behavior required before publishing or directly registering it for GPUDirect RDMA. + */ +Result prepareGdrMemory(void *devicePtr, size_t viewSize, int expectedDeviceId); + +} // namespace hf3fs::cuda diff --git a/src/common/net/ib/RDMABuf.cc b/src/common/net/ib/RDMABuf.cc index 73e7e518..5d0b07f8 100644 --- a/src/common/net/ib/RDMABuf.cc +++ b/src/common/net/ib/RDMABuf.cc @@ -9,6 +9,7 @@ #include #include +#include "common/cuda/CudaMemory.h" #include "common/monitor/Recorder.h" #include "common/net/ib/IBDevice.h" #include "memory/common/GlobalMemoryAllocator.h" @@ -47,6 +48,44 @@ RDMABuf RDMABuf::createFromUserBuffer(uint8_t *buf, size_t len) { return RDMABuf(std::shared_ptr(new RDMABuf::Inner(std::move(inner)), RDMABuf::Inner::deallocate)); } +Result RDMABuf::createFromCudaBuffer(uint8_t *ptr, + size_t len, + int expectedCudaDeviceId, + BackingOwner backingOwner) { +#ifdef HF3FS_ENABLE_GDR + // NVIDIA requires SYNC_MEMOPS to be enabled before a CUDA allocation is + // pinned for GPUDirect RDMA. Do this in the registering process as well as + // in the exporter, because an IPC mapping has its own CUDA context. + auto view = cuda::prepareGdrMemory(ptr, len, expectedCudaDeviceId); + RETURN_ON_ERROR(view); + + if (!IBManager::initialized()) { + return makeError(RPCCode::kIBDeviceNotInitialized, "IB is not initialized"); + } + if (IBDevice::all().empty()) { + return makeError(RPCCode::kIBDeviceNotFound, "no active IB device is available"); + } + + // Keep the imported/allocation device current while nvidia_peermem pins the + // view through ibv_reg_mr; the caller's previous device is restored after + // registration completes. + auto deviceGuard = cuda::ScopedDevice::create(view->deviceId); + RETURN_ON_ERROR(deviceGuard); + + RDMABuf::Inner inner(ptr, len, view->deviceId, std::move(backingOwner)); + if (inner.registerMemory() != 0) { + return makeError(StatusCode::kIOError, "failed to register CUDA memory with every active IB device"); + } + return RDMABuf(std::shared_ptr(new RDMABuf::Inner(std::move(inner)), RDMABuf::Inner::deallocate)); +#else + (void)ptr; + (void)len; + (void)expectedCudaDeviceId; + (void)backingOwner; + return makeError(StatusCode::kNotImplemented, "CUDA/GDR is disabled in this build"); +#endif +} + RDMABuf::Inner::~Inner() { XLOGF(DBG, "RDMABuf free and deregister, ptr {}", (void *)ptr_); for (auto &dev : IBDevice::all()) { @@ -56,7 +95,7 @@ RDMABuf::Inner::~Inner() { dev->deregMemory(mr); } } - if (ptr_ && !userBuffer_) { + if (ptr_ && ownsAllocation_) { rdmaBufMem.addSample(-capacity_); hf3fs::memory::deallocate(ptr_); } @@ -174,4 +213,4 @@ void RDMABufPool::deallocate(RDMABuf::Inner *buf) { sem_.signal(); } -} // namespace hf3fs::net \ No newline at end of file +} // namespace hf3fs::net diff --git a/src/common/net/ib/RDMABuf.h b/src/common/net/ib/RDMABuf.h index f6fce297..9b50d472 100644 --- a/src/common/net/ib/RDMABuf.h +++ b/src/common/net/ib/RDMABuf.h @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -38,6 +39,7 @@ namespace hf3fs::net { using RDMABufMR = ibv_mr *; class RDMARemoteBuf { + public: struct Rkey { uint32_t rkey = 0; int devId = -1; @@ -45,7 +47,6 @@ class RDMARemoteBuf { bool operator==(const Rkey &) const = default; }; - public: RDMARemoteBuf() : addr_(0), length_(0), @@ -87,7 +88,7 @@ class RDMARemoteBuf { } RDMARemoteBuf subrange(size_t offset, size_t len) const { - if (UNLIKELY(offset + len > length_)) { + if (UNLIKELY(offset > length_ || len > length_ - offset)) { return RDMARemoteBuf(); } RDMARemoteBuf remote = *this; @@ -137,6 +138,8 @@ class RDMABufPool; class RDMABuf { public: + using BackingOwner = std::shared_ptr; + RDMABuf() : RDMABuf(nullptr, nullptr, 0) {} @@ -159,7 +162,14 @@ class RDMABuf { size_t size() const { return length_; } bool empty() const { return size() == 0; } - bool contains(const uint8_t *data, uint32_t len) const { return ptr() <= data && data + len <= ptr() + capacity(); } + bool contains(const uint8_t *data, uint32_t len) const { + if (!buf_ || !begin_ || !data) { + return false; + } + auto base = reinterpret_cast(begin_); + auto address = reinterpret_cast(data); + return address >= base && address - base <= length_ && len <= length_ - (address - base); + } void resetRange() { if (LIKELY(buf_ != nullptr)) { @@ -172,7 +182,11 @@ class RDMABuf { if (UNLIKELY(n > size())) { return false; } - begin_ += n; + auto next = addAddress(begin_, n); + if (UNLIKELY(!next)) { + return false; + } + begin_ = next; length_ -= n; return true; } @@ -185,10 +199,11 @@ class RDMABuf { } RDMABuf subrange(size_t offset, size_t length) const { - if (UNLIKELY(offset + length > size())) { + if (UNLIKELY(offset > size() || length > size() - offset)) { return RDMABuf(); } - return RDMABuf(buf_, begin_ + offset, length); + auto begin = addAddress(begin_, offset); + return begin ? RDMABuf(buf_, begin, length) : RDMABuf(); } RDMABuf first(size_t length) const { return subrange(0, length); } @@ -218,20 +233,35 @@ class RDMABuf { } RDMARemoteBuf toRemoteBuf() const { - std::array rkeys; + std::array rkeys{}; if (UNLIKELY(!buf_ || !buf_->getRkeys(rkeys))) { return RDMARemoteBuf(); } - return RDMARemoteBuf((uint64_t)begin_, length_, rkeys); + return RDMARemoteBuf(reinterpret_cast(begin_), length_, rkeys); } operator RDMARemoteBuf() { return toRemoteBuf(); } - operator std::span() const { return {ptr(), size()}; } - operator std::span() { return {ptr(), size()}; } + operator std::span() const { + assertHostAccessible(); + return {ptr(), size()}; + } + operator std::span() { + assertHostAccessible(); + return {ptr(), size()}; + } - operator folly::MutableByteRange() { return {ptr(), size()}; } - operator folly::ByteRange() const { return {ptr(), size()}; } + operator folly::MutableByteRange() { + assertHostAccessible(); + return {ptr(), size()}; + } + operator folly::ByteRange() const { + assertHostAccessible(); + return {ptr(), size()}; + } + + bool isDeviceMemory() const { return buf_ && buf_->cudaDeviceId().has_value(); } + std::optional cudaDeviceId() const { return buf_ ? buf_->cudaDeviceId() : std::nullopt; } bool operator==(const RDMABuf &o) const = default; @@ -251,21 +281,36 @@ class RDMABuf { ptr_(nullptr), capacity_(capacity), mrs_(), - userBuffer_(false) {} + ownsAllocation_(true), + cudaDeviceId_(), + backingOwner_() {} Inner(uint8_t *buf, size_t len) : pool_(std::weak_ptr()), ptr_(buf), capacity_(len), mrs_(), - userBuffer_(true) {} + ownsAllocation_(false), + cudaDeviceId_(), + backingOwner_() {} + + Inner(uint8_t *buf, size_t len, int cudaDeviceId, BackingOwner backingOwner) + : pool_(), + ptr_(buf), + capacity_(len), + mrs_(), + ownsAllocation_(false), + cudaDeviceId_(cudaDeviceId), + backingOwner_(std::move(backingOwner)) {} Inner(Inner &&o) : pool_(std::move(o.pool_)), ptr_(std::exchange(o.ptr_, nullptr)), capacity_(std::exchange(o.capacity_, 0)), mrs_(std::exchange(o.mrs_, std::array())), - userBuffer_(std::exchange(o.userBuffer_, false)) {} + ownsAllocation_(std::exchange(o.ownsAllocation_, false)), + cudaDeviceId_(std::exchange(o.cudaDeviceId_, std::nullopt)), + backingOwner_(std::move(o.backingOwner_)) {} ~Inner(); @@ -281,13 +326,17 @@ class RDMABuf { const uint8_t *ptr() const { return ptr_; } uint8_t *ptr() { return ptr_; } size_t capacity() const { return capacity_; } + std::optional cudaDeviceId() const { return cudaDeviceId_; } private: std::weak_ptr pool_; uint8_t *ptr_; size_t capacity_; std::array mrs_; - bool userBuffer_; + bool ownsAllocation_; + std::optional cudaDeviceId_; + // Destroyed after ~Inner() deregisters every MR and releases any owned allocation. + BackingOwner backingOwner_; }; static RDMABuf allocate(size_t size, std::weak_ptr pool); @@ -309,7 +358,30 @@ class RDMABuf { static RDMABuf createFromUserBuffer(uint8_t *buf, size_t len); + /** + * Prepare a CUDA device-memory view for GPUDirect RDMA and register it with + * every active IB device. The optional owner is released after all MRs have + * been deregistered. Pass -1 as the expected device to discover it from the + * pointer. + */ + static Result createFromCudaBuffer(uint8_t *ptr, + size_t len, + int expectedCudaDeviceId, + BackingOwner backingOwner = {}); + private: + static uint8_t *addAddress(uint8_t *ptr, size_t offset) { + auto address = reinterpret_cast(ptr); + if (!ptr || address > std::numeric_limits::max() - offset) { + return nullptr; + } + return reinterpret_cast(address + offset); + } + + void assertHostAccessible() const { + XLOGF_IF(FATAL, isDeviceMemory(), "CUDA device memory cannot be accessed as a host byte range"); + } + std::shared_ptr buf_; uint8_t *begin_; size_t length_; diff --git a/src/fbs/storage/Common.h b/src/fbs/storage/Common.h index b5d17726..b11693a4 100644 --- a/src/fbs/storage/Common.h +++ b/src/fbs/storage/Common.h @@ -75,6 +75,7 @@ enum class FeatureFlags : uint32_t { BYPASS_RDMAXMIT = 2, SEND_DATA_INLINE = 4, ALLOW_READ_UNCOMMITTED = 8, + SERVER_COMPUTE_CHECKSUM = 16, }; constexpr auto kAIOAlignSize = 4096ul; diff --git a/src/fuse/CMakeLists.txt b/src/fuse/CMakeLists.txt index eaaeb9cd..e27e1330 100644 --- a/src/fuse/CMakeLists.txt +++ b/src/fuse/CMakeLists.txt @@ -5,6 +5,9 @@ elseif(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64") endif() target_add_lib(hf3fs_fuse common core-app meta-client storage-client fuse3 client-lib-common) +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(hf3fs_fuse SCOPE PUBLIC) +endif() target_add_bin(hf3fs_fuse_main hf3fs_fuse.cpp hf3fs_fuse) if (ENABLE_FUSE_APPLICATION) diff --git a/src/fuse/FuseClients.cc b/src/fuse/FuseClients.cc index 3bd676c4..3b1d914d 100644 --- a/src/fuse/FuseClients.cc +++ b/src/fuse/FuseClients.cc @@ -1,5 +1,6 @@ #include "FuseClients.h" +#include #include #include #include @@ -45,6 +46,23 @@ Result establishClientSession(client::IMgmtdClientForClient &mgmtdClient) } } // namespace +void detail::lookupIovBuffers(const IovTable &iovs, + std::vector> &output, + meta::Uid requester, + const IoArgs *args, + const IoSqe *sqes, + int count) { + std::vector requests; + requests.reserve(count); + for (int i = 0; i < count; ++i) { + const auto &arg = args[sqes[i].index]; + Uuid id; + memcpy(id.data, arg.bufId, sizeof(id.data)); + requests.push_back(IovLookupRequest{id, arg.bufOff, arg.ioLen}); + } + iovs.lookupBufs(output, requests, requester); +} + FuseClients::~FuseClients() { stop(); } Result FuseClients::init(const flat::AppInfo &appInfo, @@ -213,6 +231,9 @@ void FuseClients::stop() { client->stopAndJoin(); client.reset(); } +#ifdef HF3FS_ENABLE_GDR + iovs.clearGpuIovs(); +#endif } CoTask FuseClients::ioRingWorker(int i, int ths) { @@ -293,44 +314,12 @@ CoTask FuseClients::ioRingWorker(int i, int ths) { ins.push_back(it == inodes.end() ? (std::shared_ptr()) : it->second); } }; - auto lookupBufs = - [this](std::vector> &bufs, const IoArgs *args, const IoSqe *sqe, int sqec) { - auto lastId = Uuid::zero(); - std::shared_ptr lastShm; - - std::lock_guard lock(iovs.shmLock); - for (int i = 0; i < sqec; ++i) { - auto &arg = args[sqe[i].index]; - Uuid id; - memcpy(id.data, arg.bufId, sizeof(id.data)); - - std::shared_ptr shm; - if (i && id == lastId) { - shm = lastShm; - } else { - auto it = iovs.shmsById.find(id); - if (it == iovs.shmsById.end()) { - bufs.emplace_back(makeError(StatusCode::kInvalidArg, "buf id not found")); - continue; - } - - auto iovd = it->second; - shm = iovs.iovs->table[iovd].load(); - if (!shm) { - bufs.emplace_back(makeError(StatusCode::kInvalidArg, "buf id not found")); - continue; - } else if (shm->size < arg.bufOff + arg.ioLen) { - bufs.emplace_back(makeError(StatusCode::kInvalidArg, "invalid buf off and/or io len")); - continue; - } - - lastId = id; - lastShm = shm; - } - - bufs.emplace_back(lib::ShmBufForIO(std::move(shm), arg.bufOff)); - } - }; + auto lookupBufs = [this, requester = job.ior->userInfo().uid](std::vector> &bufs, + const IoArgs *args, + const IoSqe *sqe, + int sqec) { + detail::lookupIovBuffers(iovs, bufs, requester, args, sqe, sqec); + }; co_await job.ior->process(job.sqeProcTail, job.toProc, diff --git a/src/fuse/FuseClients.h b/src/fuse/FuseClients.h index 8cb0e799..3b8fb6d4 100644 --- a/src/fuse/FuseClients.h +++ b/src/fuse/FuseClients.h @@ -176,6 +176,17 @@ struct DirEntryInodeVector { inodes(std::move(inodes)) {} }; +namespace detail { + +void lookupIovBuffers(const IovTable &iovs, + std::vector> &output, + meta::Uid requester, + const IoArgs *args, + const IoSqe *sqes, + int count); + +} // namespace detail + struct FuseClients { FuseClients() = default; ~FuseClients(); diff --git a/src/fuse/FuseOps.cc b/src/fuse/FuseOps.cc index 6ff77883..8d8b8b90 100644 --- a/src/fuse/FuseOps.cc +++ b/src/fuse/FuseOps.cc @@ -598,7 +598,7 @@ bool flushAndSync(fuse_req_t req, struct fuse_entry_param *e = nullptr) { auto ino = real_ino(fino); - struct fuse_file_info fi2 {}; + struct fuse_file_info fi2{}; if (!fi) { fi = &fi2; } @@ -815,7 +815,7 @@ void hf3fs_setattr(fuse_req_t req, fuse_ino_t fino, struct stat *attr, int to_se } auto ino = real_ino(fino); - struct fuse_file_info fi2 {}; + struct fuse_file_info fi2{}; if (!fi) { fi = &fi2; } @@ -1234,7 +1234,14 @@ void hf3fs_symlink(fuse_req_t req, const char *link, fuse_ino_t fparent, const c ior->ioDepth, *ior->iora); if (!res2) { - handle_error(req, res); + auto rollback = d.iovs.rmIov(name, userInfo, ior); + XLOGF_IF(ERR, + !rollback, + "failed to roll back iov {} after io-ring creation failed: {}", + name, + rollback.error()); + handle_error(req, res2); + return; } // record the ior index for later removal res->second->iorIndex = *res2; @@ -1779,7 +1786,7 @@ void hf3fs_opendir(fuse_req_t req, fuse_ino_t fino, struct fuse_file_info *fi) { XLOGF(OP_LOG_LEVEL, "hf3fs_opendir(ino={}, pid={})", ino, fuse_req_ctx(req)->pid); record("opendir", fuse_req_ctx(req)->uid); - fi->fh = (uintptr_t) new DirHandle{d.dirHandle.fetch_add(1), fuse_req_ctx(req)->pid, false}; + fi->fh = (uintptr_t)new DirHandle{d.dirHandle.fetch_add(1), fuse_req_ctx(req)->pid, false}; auto dname = checkVirtDir(ino); if (dname && *dname == "iovs") { @@ -1815,17 +1822,9 @@ void hf3fs_releasedir(fuse_req_t req, fuse_ino_t fino, struct fuse_file_info *fi if (dh->iovDir) { // releasedir() is called only the last process with the inherited fd closes it or exits - auto &iovs = *d.iovs.iovs; - auto n = iovs.slots.nextAvail.load(); - for (int i = 0; i < n; ++i) { - auto iov = iovs.table[i].load(); - if (iov && iov->pid == dh->pid) { - XLOGF(INFO, "unlinking iov {} symlink from dead pid {}", iov->key, dh->pid); - d.iovs.rmIov(iov->key.c_str(), meta::UserInfo{iov->user, meta::Gid{iov->user.toUnderType()}}); - if (iov->isIoRing) { - d.iors.rmIoRing(iov->iorIndex); - } - } + auto ioRingIndexes = d.iovs.removeIovsByPid(dh->pid); + for (auto ioRingIndex : ioRingIndexes) { + d.iors.rmIoRing(ioRingIndex); } } diff --git a/src/fuse/IoRing.cc b/src/fuse/IoRing.cc index 7009198d..5be294df 100644 --- a/src/fuse/IoRing.cc +++ b/src/fuse/IoRing.cc @@ -71,7 +71,7 @@ CoTask IoRing::process( const storage::client::IoOptions &storageIo, UserConfig &userConfig, std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupFiles, - std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupBufs) { + std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupBufs) { static monitor::LatencyRecorder overallLatency("usrbio.piov.overall", monitor::TagSet{{"mount_name", mountName}}); static monitor::LatencyRecorder prepareLatency("usrbio.piov.prepare", monitor::TagSet{{"mount_name", mountName}}); static monitor::LatencyRecorder submitLatency("usrbio.piov.submit", monitor::TagSet{{"mount_name", mountName}}); @@ -109,7 +109,7 @@ CoTask IoRing::process( lookupFiles(inodes, ringSection, sqeSection, toProc - (int)inodes.size()); } - std::vector> bufs; + std::vector> bufs; bufs.reserve(toProc); lookupBufs(bufs, ringSection, sqeSection + spt, std::min(toProc, entries - spt)); if ((int)bufs.size() < toProc) { @@ -148,11 +148,19 @@ CoTask IoRing::process( continue; } +#ifdef HF3FS_ENABLE_GDR + auto bufPtr = ioBufPtr(bufs[i].value()); + auto memh = + co_await std::visit([&](auto &b) -> CoTryTask { return b.memh(args.ioLen); }, + bufs[i].value()); +#else + auto bufPtr = bufs[i]->ptr(); auto memh = co_await bufs[i]->memh(args.ioLen); +#endif if (!memh) { res[i] = -static_cast(memh.error().code()); continue; - } else if (!bufs[i]->ptr() || !*memh) { + } else if (!bufPtr || !*memh) { XLOGF(ERR, "{} is null when doing usrbio", *memh ? "buf ptr" : "memh"); res[i] = -static_cast(ClientAgentCode::kIovShmFail); continue; @@ -168,9 +176,8 @@ CoTask IoRing::process( truncateVers[i] = *beginWrite; } - auto addRes = forRead_ - ? ioExec.addRead(i, inodes[i]->inode, 0, args.fileOff, args.ioLen, bufs[i]->ptr(), **memh) - : ioExec.addWrite(i, inodes[i]->inode, 0, args.fileOff, args.ioLen, bufs[i]->ptr(), **memh); + auto addRes = forRead_ ? ioExec.addRead(i, inodes[i]->inode, 0, args.fileOff, args.ioLen, bufPtr, **memh) + : ioExec.addWrite(i, inodes[i]->inode, 0, args.fileOff, args.ioLen, bufPtr, **memh); if (!addRes) { res[i] = -static_cast(addRes.error().code()); } diff --git a/src/fuse/IoRing.h b/src/fuse/IoRing.h index b7d022ce..d114f98d 100644 --- a/src/fuse/IoRing.h +++ b/src/fuse/IoRing.h @@ -3,7 +3,7 @@ #include #include -#include "IovTable.h" +#include "IovTypes.h" #include "UserConfig.h" #include "client/storage/StorageClient.h" #include "common/utils/AtomicSharedPtrTable.h" @@ -13,6 +13,7 @@ #include "lib/common/Shm.h" namespace hf3fs::fuse { + struct RcInode; struct IoArgs { uint8_t bufId[16]; @@ -120,6 +121,7 @@ class IoRing : public std::enable_shared_from_this { } std::vector jobsToProc(int maxJobs); int cqeCount() const { return (cqeHead.load() + entries - cqeTail.load()) % entries; } + const meta::UserInfo &userInfo() const { return userInfo_; } CoTask process( int spt, int toProc, @@ -127,7 +129,7 @@ class IoRing : public std::enable_shared_from_this { const storage::client::IoOptions &storageIo, UserConfig &userConfig, std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupFiles, - std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupBufs); + std::function> &, const IoArgs *, const IoSqe *, int)> &&lookupBufs); public: bool addSqe(int idx, const void *userdata) { diff --git a/src/fuse/IovTable.cc b/src/fuse/IovTable.cc index f0cca1ff..5a68cde2 100644 --- a/src/fuse/IovTable.cc +++ b/src/fuse/IovTable.cc @@ -1,9 +1,19 @@ #include "IovTable.h" +#include +#include +#include +#include #include +#include +#include +#include +#include +#include #include "IoRing.h" #include "fbs/meta/Common.h" +#include "lib/common/GdrUri.h" namespace hf3fs::fuse { @@ -11,9 +21,80 @@ using hf3fs::lib::IorAttrs; const Path linkPref = "/dev/shm"; +namespace { + +std::optional parsePositiveSize(std::string_view text) { + if (text.empty()) return std::nullopt; + size_t value = 0; + auto [ptr, ec] = std::from_chars(text.data(), text.data() + text.size(), value); + if (ec != std::errc{} || ptr != text.data() + text.size() || value == 0) { + return std::nullopt; + } + return value; +} + +std::optional parseNonNegativeInt(std::string_view text) { + if (text.empty()) return std::nullopt; + int value = 0; + auto [ptr, ec] = std::from_chars(text.data(), text.data() + text.size(), value); + if (ec != std::errc{} || ptr != text.data() + text.size() || value < 0) { + return std::nullopt; + } + return value; +} + +std::optional parseBinaryFlags(std::string_view text) { + if (text.empty()) return std::nullopt; + uint64_t value = 0; + auto [ptr, ec] = std::from_chars(text.data(), text.data() + text.size(), value, 2); + if (ec != std::errc{} || ptr != text.data() + text.size()) { + return std::nullopt; + } + return value; +} + +} // namespace + void IovTable::init(const Path &mount, int cap) { mountName = mount.native(); - iovs = std::make_unique>(cap); + entries_ = std::make_unique>(cap); +} + +Result IovTable::publishEntry(int iovd, std::shared_ptr entry) { + std::unique_lock lock(mutex_); + if (!entries_ || iovd < 0 || iovd >= (int)entries_->table.size() || entries_->table[iovd].load()) { + return makeError(StatusCode::kInvalidArg, "invalid iov descriptor"); + } + if (iovdsByKey_.find(entry->key) != iovdsByKey_.end() || iovdsById_.find(entry->id) != iovdsById_.end()) { + return makeError(MetaCode::kExists, "iov key or UUID already exists"); + } + + iovdsByKey_.emplace(entry->key, iovd); + iovdsById_.emplace(entry->id, iovd); + entries_->table[iovd].store(std::move(entry)); + return Void{}; +} + +bool IovTable::removeEntryLocked(int iovd, const std::shared_ptr &expected) { + if (!entries_ || iovd < 0 || iovd >= (int)entries_->table.size()) { + return false; + } + + auto current = entries_->table[iovd].load(); + if (!current || current != expected) { + return false; + } + + auto keyIt = iovdsByKey_.find(current->key); + if (keyIt != iovdsByKey_.end() && keyIt->second == iovd) { + iovdsByKey_.erase(keyIt); + } + auto idIt = iovdsById_.find(current->id); + if (idIt != iovdsById_.end() && idIt->second == iovd) { + iovdsById_.erase(idIt); + } + entries_->remove(iovd); + return true; } struct IovAttrs { @@ -23,6 +104,8 @@ struct IovAttrs { bool forRead = true; int ioDepth = 0; std::optional iora; + bool isGdr = false; + int gpuDeviceId = -1; }; static Result parseKey(const char *key) { @@ -30,6 +113,9 @@ static Result parseKey(const char *key) { std::vector fnParts; folly::split('.', key, fnParts); + if (fnParts.empty() || fnParts[0].empty()) { + return makeError(StatusCode::kInvalidArg, "invalid shm key"); + } auto idRes = Uuid::fromHexString(fnParts[0]); RETURN_ON_ERROR(idRes); @@ -37,22 +123,28 @@ static Result parseKey(const char *key) { for (size_t i = 1; i < fnParts.size(); ++i) { auto dec = fnParts[i]; + if (dec.empty()) { + return makeError(StatusCode::kInvalidArg, "empty attr in shm key"); + } switch (dec[0]) { case 'b': { // block size - auto i = atoll(dec.c_str() + 1); - if (i <= 0) { + auto blockSize = parsePositiveSize(std::string_view(dec).substr(1)); + if (!blockSize) { return makeError(StatusCode::kInvalidArg, "invalid block size set in shm key"); } - iova.blockSize = (size_t)i; + iova.blockSize = *blockSize; break; } case 'r': case 'w': { // is io ring - auto i = atoll(dec.c_str() + 1); + auto depth = parseNonNegativeInt(std::string_view(dec).substr(1)); + if (!depth) { + return makeError(StatusCode::kInvalidArg, "invalid io batch size set in shm key"); + } iova.isIoRing = true; iova.forRead = dec[0] == 'r'; - iova.ioDepth = i; + iova.ioDepth = *depth; break; } @@ -60,11 +152,11 @@ static Result parseKey(const char *key) { if (!iova.iora) { iova.iora = IorAttrs{}; } - auto i = atoi(dec.c_str() + 1); - if (i < 0) { + auto timeoutMs = parseNonNegativeInt(std::string_view(dec).substr(1)); + if (!timeoutMs) { return makeError(StatusCode::kInvalidArg, "invalid io job check timeout {}", dec.c_str() + 1); } - iova.iora->timeout = Duration(std::chrono::nanoseconds((uint64_t)i * 1000000)); + iova.iora->timeout = Duration(std::chrono::nanoseconds((uint64_t)*timeoutMs * 1000000)); break; } @@ -72,12 +164,11 @@ static Result parseKey(const char *key) { if (!iova.iora) { iova.iora = IorAttrs{}; } - char *ep; - auto i = strtoull(dec.c_str() + 1, &ep, 2); - if (*ep != 0 || i < 0) { + auto flags = parseBinaryFlags(std::string_view(dec).substr(1)); + if (!flags) { return makeError(StatusCode::kInvalidArg, "invalid io exec flags {}", dec.c_str() + 1); } - iova.iora->flags = i; + iova.iora->flags = *flags; break; } @@ -85,6 +176,9 @@ static Result parseKey(const char *key) { if (!iova.iora) { iova.iora = IorAttrs{}; } + if (dec.size() > 2) { + return makeError(StatusCode::kInvalidArg, "invalid priority set in shm key"); + } switch (dec.c_str()[1]) { case 'l': iova.iora->priority = 2; @@ -100,11 +194,42 @@ static Result parseKey(const char *key) { return makeError(StatusCode::kInvalidArg, "invalid priority set in shm key"); } break; + + case 'g': // gdr marker (e.g. ".gdr") + if (dec == "gdr") { + iova.isGdr = true; + } else { + return makeError(StatusCode::kInvalidArg, "invalid gdr attr in shm key"); + } + break; + + case 'd': { // gpu device id (e.g. ".d0", ".d1") + auto devId = parseNonNegativeInt(std::string_view(dec).substr(1)); + if (!devId) { + return makeError(StatusCode::kInvalidArg, "invalid gpu device id in key"); + } + iova.gpuDeviceId = *devId; + break; + } + default: + return makeError(StatusCode::kInvalidArg, "unknown attr in shm key"); } } - if (!iova.isIoRing && iova.iora) { - return makeError(StatusCode::kInvalidArg, "ioring attrs set for non-ioring"); + if (iova.isGdr) { + if (iova.gpuDeviceId < 0) { + return makeError(StatusCode::kInvalidArg, "gdr key missing gpu device id"); + } + if (iova.blockSize != 0 || iova.isIoRing || iova.iora) { + return makeError(StatusCode::kInvalidArg, "gdr key does not support block or io-ring attrs"); + } + } else { + if (!iova.isIoRing && iova.iora) { + return makeError(StatusCode::kInvalidArg, "ioring attrs set for non-ioring"); + } + if (iova.gpuDeviceId >= 0) { + return makeError(StatusCode::kInvalidArg, "gpu device id set for non-gdr key"); + } } return iova; @@ -138,6 +263,70 @@ Result>> IovTable::addIov(co auto iovaRes = parseKey(key); RETURN_ON_ERROR(iovaRes); + { + std::shared_lock lock(mutex_); + if (iovdsByKey_.find(key) != iovdsByKey_.end() || iovdsById_.find(iovaRes->id) != iovdsById_.end()) { + return makeError(MetaCode::kExists, "iov key or UUID already exists"); + } + } + +#ifndef HF3FS_ENABLE_GDR + if (iovaRes->isGdr) { + return makeError(StatusCode::kInvalidArg, "GDR not enabled in this build"); + } +#endif + +#ifdef HF3FS_ENABLE_GDR + // GDR path: shmPath is a gdr:// URI, not a filesystem path + if (iovaRes->isGdr) { + auto gdrTarget = lib::parseGdrUri(shmPath.native()); + if (!gdrTarget) { + return makeError(StatusCode::kInvalidArg, "failed to parse GDR target URI"); + } + if (iovaRes->gpuDeviceId != gdrTarget->deviceId) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "gdr key device {} does not match URI device {}", + iovaRes->gpuDeviceId, + gdrTarget->deviceId); + } + + lib::CudaIpcHandle ipcHandle; + std::copy(gdrTarget->ipcHandle.begin(), gdrTarget->ipcHandle.end(), ipcHandle.begin()); + auto gpuShmResult = lib::GpuShmBuf::create(ipcHandle, + gdrTarget->allocationSize, + gdrTarget->offset, + gdrTarget->size, + gdrTarget->deviceId); + RETURN_ON_ERROR(gpuShmResult); + auto gpuShm = std::move(*gpuShmResult); + + // Allocate iov descriptor slot + auto iovdRes = entries_->alloc(); + if (!iovdRes) { + return makeError(ClientAgentCode::kTooManyOpenFiles, "too many iovs allocated"); + } + auto iovd = *iovdRes; + bool dealloc = true; + SCOPE_EXIT { + if (dealloc) { + entries_->dealloc(iovd); + } + }; + + auto entry = std::make_shared( + IovEntry{std::string(key), iovaRes->id, shmPath, ui.uid, ui.gid, pid, IovBuffer{gpuShm}}); + RETURN_ON_ERROR(publishEntry(iovd, entry)); + + // For GPU iovs, we return the GDR URI as the symlink target + auto inode = + meta::Inode{meta::InodeId::iov(iovd), + meta::InodeData{meta::Symlink{shmPath}, meta::Acl{ui.uid, ui.gid, meta::Permission(0400)}}}; + + dealloc = false; + return std::make_pair(inode, std::shared_ptr()); + } +#endif + Path shmOpenPath("/"); shmOpenPath /= shmPath.lexically_relative(linkPref); @@ -153,7 +342,7 @@ Result>> IovTable::addIov(co } while (true) { - auto iovdRes = iovs->alloc(); + auto iovdRes = entries_->alloc(); if (!iovdRes) { return makeError(ClientAgentCode::kTooManyOpenFiles, "too many iovs allocated"); } @@ -161,7 +350,7 @@ Result>> IovTable::addIov(co bool dealloc = true; SCOPE_EXIT { if (dealloc) { - iovs->dealloc(iovd); + entries_->dealloc(iovd); } }; @@ -170,29 +359,29 @@ Result>> IovTable::addIov(co std::shared_ptr shm; try { - shm.reset( - new lib::ShmBuf(shmOpenPath, 0, st.st_size, iovaRes->blockSize, iovaRes->id), - [uids, - &shmSizeCount = shmSizeCount, - &mapTimesCount = mapTimesCount, - &mapBytesDist = mapBytesDist, - &allocLatency = allocLatency, - &ibRegLatency = ibRegLatency](auto p) { - auto start = SteadyClock::now(); - folly::coro::blockingWait(p->deregisterForIO()); - auto now = SteadyClock::now(); - ibRegLatency.addSample(now - start, monitor::TagSet{{"instance", "dereg"}, {"uid", uids}}); - - start = now; - p->unmapBuf(); - allocLatency.addSample(SteadyClock::now() - start, monitor::TagSet{{"instance", "free"}, {"uid", uids}}); - - mapTimesCount.addSample(1, monitor::TagSet{{"instance", "free"}, {"uid", uids}}); - mapBytesDist.addSample(p->size, monitor::TagSet{{"instance", "free"}, {"uid", uids}}); - shmSizeCount.addSample(-p->size); - - delete p; - }); + shm.reset(new lib::ShmBuf(shmOpenPath, 0, st.st_size, iovaRes->blockSize, iovaRes->id), + [uids, + &shmSizeCount = shmSizeCount, + &mapTimesCount = mapTimesCount, + &mapBytesDist = mapBytesDist, + &allocLatency = allocLatency, + &ibRegLatency = ibRegLatency](auto p) { + auto start = SteadyClock::now(); + folly::coro::blockingWait(p->deregisterForIO()); + auto now = SteadyClock::now(); + ibRegLatency.addSample(now - start, monitor::TagSet{{"instance", "dereg"}, {"uid", uids}}); + + start = now; + p->unmapBuf(); + allocLatency.addSample(SteadyClock::now() - start, + monitor::TagSet{{"instance", "free"}, {"uid", uids}}); + + mapTimesCount.addSample(1, monitor::TagSet{{"instance", "free"}, {"uid", uids}}); + mapBytesDist.addSample(p->size, monitor::TagSet{{"instance", "free"}, {"uid", uids}}); + shmSizeCount.addSample(-p->size); + + delete p; + }); } catch (const std::runtime_error &e) { return makeError(ClientAgentCode::kIovShmFail, std::string("failed to open/map shm for iov ") + e.what()); } @@ -210,9 +399,6 @@ Result>> IovTable::addIov(co shm->ioDepth = iovaRes->ioDepth; shm->iora = iovaRes->iora; - // the idx should be reserved by us - iovs->table[iovd].store(shm); - start = SteadyClock::now(); auto recordMetrics = [blockSize = shm->blockSize, start, uids]() mutable { ibRegBytesDist.addSample(blockSize, monitor::TagSet{{"instance", "reg"}, {"uid", uids}}); @@ -223,15 +409,9 @@ Result>> IovTable::addIov(co folly::coro::blockingWait(shm->registerForIO(exec, sc, recordMetrics)); } - { - std::unique_lock lock(iovdLock_); - iovds_[key] = iovd; - } - - { - std::unique_lock lock(shmLock); - shmsById[iovaRes->id] = iovd; - } + auto entry = std::make_shared( + IovEntry{std::string(key), iovaRes->id, linkPref / shm->path, ui.uid, ui.gid, pid, IovBuffer{shm}}); + RETURN_ON_ERROR(publishEntry(iovd, entry)); auto statRes = statIov(iovd, ui); RETURN_ON_ERROR(statRes); @@ -241,70 +421,225 @@ Result>> IovTable::addIov(co } } -Result> IovTable::rmIov(const char *key, const meta::UserInfo &ui) { - auto res = lookupIov(key, ui); - RETURN_ON_ERROR(res); - - { - std::unique_lock lock(iovdLock_); - iovds_.erase(key); - } - +Result> IovTable::rmIov(const char *key, + const meta::UserInfo &ui, + const std::shared_ptr &expectedBuffer) { + std::shared_ptr entry; + int iovd = -1; { - auto res = parseKey(key); + std::unique_lock lock(mutex_); + auto it = iovdsByKey_.find(key); + if (it == iovdsByKey_.end() || !entries_) { + return makeError(MetaCode::kNotFound, std::string("iov key not found ") + key); + } - std::unique_lock lock(shmLock); - shmsById.erase(res->id); + iovd = it->second; + entry = entries_->table[iovd].load(); + if (!entry || entry->key != key) { + return makeError(MetaCode::kNotFound, std::string("iov key not found ") + key); + } + if (entry->user != ui.uid) { + XLOGF(ERR, "removing user {} iov belongs to {}", ui.uid, entry->user); + return makeError(MetaCode::kNoPermission, "iov not for user"); + } + if (expectedBuffer) { + auto host = std::get_if>(&entry->buffer); + if (!host || *host != expectedBuffer) { + return makeError(MetaCode::kNotFound, "iov changed before rollback"); + } + } + XLOGF_IF(FATAL, !removeEntryLocked(iovd, entry), "iov entry changed while holding the table lock"); } - auto iovd = iovDesc(res->id); - auto shm = iovs->table[*iovd].load(); - iovs->remove(*iovd); +#ifdef HF3FS_ENABLE_GDR + if (entry->isGpu()) { + // RDMABuf deregisters every MR before releasing its CUDA IPC mapping owner. + entry.reset(); + return std::shared_ptr(); + } +#endif - return shm; + return std::get>(entry->buffer); } Result IovTable::statIov(int iovd, const meta::UserInfo &ui) { - if (iovd < 0 || iovd >= (int)iovs->table.size()) { + std::shared_lock lock(mutex_); + if (!entries_ || iovd < 0 || iovd >= (int)entries_->table.size()) { return makeError(MetaCode::kNotFound, "invalid iov desc"); } - auto shm = iovs->table[iovd].load(); - if (!shm) { + auto entry = entries_->table[iovd].load(); + if (!entry) { return makeError(MetaCode::kNotFound, - fmt::format("iov desc {} not found, next avail {}", iovd, iovs->slots.nextAvail.load())); + fmt::format("iov desc {} not found, next avail {}", iovd, entries_->slots.nextAvail.load())); } - if (shm->user != ui.uid) { - XLOGF(ERR, "statting user {} iov belongs to {}", ui.uid, shm->user); + if (entry->user != ui.uid) { + XLOGF(ERR, "statting user {} iov belongs to {}", ui.uid, entry->user); return makeError(MetaCode::kNoPermission, "iov not for user"); } return meta::Inode{ meta::InodeId::iov(iovd), - meta::InodeData{meta::Symlink{linkPref / shm->path}, meta::Acl{ui.uid, ui.gid, meta::Permission(0400)}}}; + meta::InodeData{meta::Symlink{entry->target}, meta::Acl{entry->user, entry->gid, meta::Permission(0400)}}}; } Result IovTable::lookupIov(const char *key, const meta::UserInfo &ui) { - int iovd = -1; + std::shared_lock lock(mutex_); + auto it = iovdsByKey_.find(key); + if (it == iovdsByKey_.end() || !entries_) { + return makeError(MetaCode::kNotFound, std::string("iov key not found ") + key); + } + + auto entry = entries_->table[it->second].load(); + if (!entry || entry->key != key) { + return makeError(MetaCode::kNotFound, std::string("iov key not found ") + key); + } + if (entry->user != ui.uid) { + XLOGF(ERR, "looking up user {} iov belongs to {}", ui.uid, entry->user); + return makeError(MetaCode::kNoPermission, "iov not for user"); + } + + return meta::Inode{ + meta::InodeId::iov(it->second), + meta::InodeData{meta::Symlink{entry->target}, meta::Acl{entry->user, entry->gid, meta::Permission(0400)}}}; +} + +namespace { + +Result makeIoBufForIO(const std::shared_ptr &entry, const IovLookupRequest &request) { + return std::visit( + [&](const auto &buffer) -> Result { + using Buffer = typename std::decay_t::element_type; + if (!buffer || request.offset > buffer->size || request.length > buffer->size - request.offset) { + return makeError(StatusCode::kInvalidArg, "invalid buf off and/or io len"); + } + + if constexpr (std::is_same_v) { + return IoBufForIO{lib::ShmBufForIO(buffer, request.offset)}; + } +#ifdef HF3FS_ENABLE_GDR + else { + return IoBufForIO{lib::GpuShmBufForIO(buffer, request.offset)}; + } +#endif + }, + entry->buffer); +} + +} // namespace + +void IovTable::lookupBufs(std::vector> &output, + std::span requests, + meta::Uid requester) const { + output.reserve(output.size() + requests.size()); + std::shared_lock lock(mutex_); + + bool hasLast = false; + bool lastDenied = false; + Uuid lastId = Uuid::zero(); + std::shared_ptr lastEntry; + for (const auto &request : requests) { + if (!hasLast || request.id != lastId) { + hasLast = true; + lastDenied = false; + lastId = request.id; + lastEntry.reset(); + + if (entries_) { + auto it = iovdsById_.find(request.id); + if (it != iovdsById_.end()) { + auto entry = entries_->table[it->second].load(); + if (entry && entry->id == request.id && entry->user == requester) { + lastEntry = std::move(entry); + } else if (entry && entry->id == request.id) { + lastDenied = true; + } + } + } + } + + if (!lastEntry) { + output.emplace_back(lastDenied ? makeError(MetaCode::kNoPermission, "iov not for user") + : makeError(StatusCode::kInvalidArg, "buf id not found")); + continue; + } + output.emplace_back(makeIoBufForIO(lastEntry, request)); + } +} + +Result IovTable::lookupBuf(const IovLookupRequest &request, meta::Uid requester) const { + std::vector> result; + lookupBufs(result, std::span(&request, 1), requester); + return std::move(result.front()); +} + +std::shared_ptr IovTable::entryAt(int iovd) const { + std::shared_lock lock(mutex_); + if (!entries_ || iovd < 0 || iovd >= (int)entries_->table.size()) { + return nullptr; + } + return entries_->table[iovd].load(); +} + +std::vector IovTable::removeIovsByPid(pid_t pid) { + struct IovToRemove { + int iovd; + std::shared_ptr entry; + std::optional ioRingIndex; + }; + + std::vector targets; { - std::shared_lock lock(iovdLock_); - auto it = iovds_.find(key); - if (it == iovds_.end()) { - return makeError(MetaCode::kNotFound, std::string("iov key not found ") + key); - } else { - iovd = it->second; + std::shared_lock lock(mutex_); + if (!entries_) { + return {}; + } + + auto n = entries_->slots.nextAvail.load(); + targets.reserve(n); + for (int i = 0; i < n; ++i) { + auto entry = entries_->table[i].load(); + if (!entry || entry->pid != pid) { + continue; + } + + std::optional ioRingIndex; + if (auto host = std::get_if>(&entry->buffer); host && *host && (*host)->isIoRing) { + ioRingIndex = (*host)->iorIndex; + } + targets.push_back(IovToRemove{i, std::move(entry), ioRingIndex}); } } - return statIov(iovd, ui); + std::vector ioRingIndexes; + ioRingIndexes.reserve(targets.size()); + for (const auto &target : targets) { + bool removed = false; + { + std::unique_lock lock(mutex_); + if (target.entry->pid == pid) { + removed = removeEntryLocked(target.iovd, target.entry); + } + } + if (!removed) { + XLOGF(DBG, "iov {} changed before dead pid {} cleanup", target.entry->key, pid); + continue; + } + XLOGF(INFO, "unlinked iov {} symlink from dead pid {}", target.entry->key, pid); + if (target.ioRingIndex) { + ioRingIndexes.push_back(*target.ioRingIndex); + } + } + return ioRingIndexes; } std::pair>, std::shared_ptr>>> IovTable::listIovs(const meta::UserInfo &ui) { meta::DirEntry de{meta::InodeId::iovDir(), ""}; - auto n = iovs->slots.nextAvail.load(); + std::shared_lock lock(mutex_); + auto n = entries_ ? entries_->slots.nextAvail.load() : 0; std::vector des; std::vector> ins; des.reserve(n + 3); @@ -318,20 +653,60 @@ IovTable::listIovs(const meta::UserInfo &ui) { ins.emplace_back(std::move(inode)); } - meta::Acl acl{meta::Uid{ui.uid}, meta::Gid{ui.gid}, meta::Permission{0400}}; for (int i = 0; i < n; ++i) { - auto iov = iovs->table[i].load(); - if (!iov || iov->user != ui.uid) { + auto entry = entries_->table[i].load(); + if (!entry || entry->user != ui.uid) { continue; } - de.name = iov->key; + de.name = entry->key; des.emplace_back(de); - ins.emplace_back( - meta::Inode{meta::InodeId{meta::InodeId::iov(i)}, meta::InodeData{meta::Symlink{linkPref / iov->path}, acl}}); + ins.emplace_back(meta::Inode{ + meta::InodeId{meta::InodeId::iov(i)}, + meta::InodeData{meta::Symlink{entry->target}, meta::Acl{entry->user, entry->gid, meta::Permission{0400}}}}); } return std::make_pair(std::make_shared>(std::move(des)), std::make_shared>>(std::move(ins))); } + +#ifdef HF3FS_ENABLE_GDR +void IovTable::clearGpuIovs() { + std::vector> gpuBuffers; + { + std::unique_lock lock(mutex_); + if (!entries_) { + return; + } + + auto n = entries_->slots.nextAvail.load(); + gpuBuffers.reserve(n); + for (int i = 0; i < n; ++i) { + auto entry = entries_->table[i].load(); + if (!entry) { + continue; + } + auto gpu = std::get_if>(&entry->buffer); + if (!gpu) { + continue; + } + + auto keyIt = iovdsByKey_.find(entry->key); + if (keyIt != iovdsByKey_.end() && keyIt->second == i) { + iovdsByKey_.erase(keyIt); + } + auto idIt = iovdsById_.find(entry->id); + if (idIt != iovdsById_.end() && idIt->second == i) { + iovdsById_.erase(idIt); + } + if (*gpu) { + gpuBuffers.push_back(*gpu); + } + entries_->remove(i); + } + } + + gpuBuffers.clear(); +} +#endif } // namespace hf3fs::fuse diff --git a/src/fuse/IovTable.h b/src/fuse/IovTable.h index 3a2085a2..3f8f8179 100644 --- a/src/fuse/IovTable.h +++ b/src/fuse/IovTable.h @@ -1,12 +1,50 @@ #pragma once +#include +#include +#include +#include #include +#include +#include +#include #include "common/utils/AtomicSharedPtrTable.h" #include "fbs/meta/Schema.h" +#include "fuse/IovTypes.h" #include "lib/common/Shm.h" +#ifdef HF3FS_ENABLE_GDR +#include "lib/common/GpuShm.h" +#endif namespace hf3fs::fuse { + +class IovTableTestHelper; + +#ifdef HF3FS_ENABLE_GDR +using IovBuffer = std::variant, std::shared_ptr>; +#else +using IovBuffer = std::variant>; +#endif + +struct IovEntry { + std::string key; + Uuid id; + Path target; + meta::Uid user{0}; + meta::Gid gid{0}; + pid_t pid = 0; + IovBuffer buffer; + + bool isGpu() const { +#ifdef HF3FS_ENABLE_GDR + return std::holds_alternative>(buffer); +#else + return false; +#endif + } +}; + class IovTable { public: IovTable() = default; @@ -17,23 +55,36 @@ class IovTable { const meta::UserInfo &ui, folly::Executor::KeepAlive<> exec, storage::client::StorageClient &sc); - Result> rmIov(const char *key, const meta::UserInfo &ui); + Result> rmIov(const char *key, + const meta::UserInfo &ui, + const std::shared_ptr &expectedBuffer = {}); Result lookupIov(const char *key, const meta::UserInfo &ui); std::optional iovDesc(meta::InodeId iid); Result statIov(int key, const meta::UserInfo &ui); + std::vector removeIovsByPid(pid_t pid); + Result lookupBuf(const IovLookupRequest &request, meta::Uid requester) const; + // Appends exactly one final result per request; existing output is untouched. + void lookupBufs(std::vector> &output, + std::span requests, + meta::Uid requester) const; + std::shared_ptr entryAt(int iovd) const; +#ifdef HF3FS_ENABLE_GDR + void clearGpuIovs(); +#endif - public: std::pair>, std::shared_ptr>>> listIovs(const meta::UserInfo &ui); - public: - std::string mountName; - std::shared_mutex shmLock; - robin_hood::unordered_map shmsById; - std::unique_ptr> iovs; - private: - mutable std::shared_mutex iovdLock_; - robin_hood::unordered_map iovds_; + friend class IovTableTestHelper; + + Result publishEntry(int iovd, std::shared_ptr entry); + bool removeEntryLocked(int iovd, const std::shared_ptr &expected); + + std::string mountName; + mutable std::shared_mutex mutex_; + robin_hood::unordered_map iovdsByKey_; + robin_hood::unordered_map iovdsById_; + std::unique_ptr> entries_; }; } // namespace hf3fs::fuse diff --git a/src/fuse/IovTypes.h b/src/fuse/IovTypes.h new file mode 100644 index 00000000..d5ffef1c --- /dev/null +++ b/src/fuse/IovTypes.h @@ -0,0 +1,33 @@ +#pragma once + +#include +#include +#include + +#include "common/utils/Uuid.h" +#include "lib/common/Shm.h" +#ifdef HF3FS_ENABLE_GDR +#include "lib/common/GpuShm.h" +#endif + +namespace hf3fs::fuse { + +#ifdef HF3FS_ENABLE_GDR +using IoBufForIO = std::variant; + +inline uint8_t *ioBufPtr(const IoBufForIO &buf) { + return std::visit([](const auto &b) -> uint8_t * { return b.ptr(); }, buf); +} +#else +using IoBufForIO = lib::ShmBufForIO; + +inline uint8_t *ioBufPtr(const IoBufForIO &buf) { return buf.ptr(); } +#endif + +struct IovLookupRequest { + Uuid id; + size_t offset; + size_t length; +}; + +} // namespace hf3fs::fuse diff --git a/src/fuse/PioV.cc b/src/fuse/PioV.cc index 2eb23815..87f57874 100644 --- a/src/fuse/PioV.cc +++ b/src/fuse/PioV.cc @@ -1,5 +1,7 @@ #include "PioV.h" +#include + namespace hf3fs::lib::agent { PioV::PioV(storage::client::StorageClient &storageClient, int chunkSizeLim, std::vector &res) : storageClient_(storageClient), @@ -183,16 +185,44 @@ CoTryTask PioV::executeWrite(const UserInfo &userInfo, const storage::clie } template -void concatIoRes(bool read, std::vector &res, const Io &ios, bool allowHoles) { +Result zeroReadRange(const Io &io, size_t offset, size_t length) { + if (io.buffer == nullptr) { + return makeError(StorageClientCode::kInvalidArg, "read hole has no registered IOBuffer"); + } + if (offset > io.length || length > io.length - offset) { + return makeError(StorageClientCode::kInvalidArg, + fmt::format("read hole range [{}, {}) exceeds IO length {}", offset, offset + length, io.length)); + } + + auto ioBufferOffset = io.buffer->offsetOf(io.data, io.length); + if (!ioBufferOffset || *ioBufferOffset > std::numeric_limits::max() - offset) { + return makeError(StorageClientCode::kInvalidArg, "read hole is outside its registered IOBuffer"); + } + return io.buffer->zeroRange(*ioBufferOffset + offset, length); +} + +template +void concatIoRes(bool read, + std::vector &res, + const Io &ios, + bool allowHoles, + const std::function(const typename Io::value_type &, size_t, size_t)> &zeroRange) { ssize_t lastIovIdx = -1; bool inHole = false; - std::optional holeIo = 0; + std::optional holeIo; size_t holeOff = 0; size_t holeSize = 0; ssize_t iovIdx = 0; for (size_t i = 0; i < ios.size(); ++i, lastIovIdx = iovIdx) { const auto &io = ios[i]; iovIdx = reinterpret_cast(io.userCtx); + if (lastIovIdx != iovIdx) { + inHole = false; + holeIo = std::nullopt; + holeOff = 0; + holeSize = 0; + } + uint32_t iolen = 0; if (io.result.lengthInfo) { iolen = *io.result.lengthInfo; @@ -224,21 +254,23 @@ void concatIoRes(bool read, std::vector &res, const Io &ios, bool allow if (read && allowHoles) { // zerofill the hole we found auto &hio = ios[*holeIo]; - memset(hio.data + holeOff, 0, hio.length - holeOff); - for (size_t j = *holeIo + 1; j < i; ++j) { - memset(ios[j].data, 0, ios[j].length); + auto zeroResult = zeroRange(hio, holeOff, hio.length - holeOff); + for (size_t j = *holeIo + 1; zeroResult && j < i; ++j) { + zeroResult = zeroRange(ios[j], 0, ios[j].length); } - res[iovIdx] += holeSize; + if (zeroResult) { + res[iovIdx] += holeSize; + } else { + XLOGF(ERR, "Failed to zero-fill read hole for iov index {}: {}", iovIdx, zeroResult.error()); + res[iovIdx] = -static_cast(zeroResult.error().code()); + } inHole = false; // out of hole now, but we may begin a new hole holeIo = std::nullopt; } else { res[iovIdx] = -static_cast(ClientAgentCode::kHoleInIoOutcome); } - } else if (lastIovIdx != iovIdx) { - inHole = false; - holeIo = std::nullopt; } } else if (read && io.result.lengthInfo.error().code() == StorageClientCode::kChunkNotFound) { // ignore @@ -265,11 +297,20 @@ void concatIoRes(bool read, std::vector &res, const Io &ios, bool allow } } +void detail::finishReadResults(std::vector &res, + std::span ios, + bool allowHoles, + const ReadHoleZeroer &zeroRange) { + concatIoRes(true, res, ios, allowHoles, zeroRange); +} + void PioV::finishIo(bool allowHoles) { if (wios_.empty()) { - concatIoRes(true, res_, rios_, allowHoles); + detail::finishReadResults(res_, rios_, allowHoles, zeroReadRange); } else { - concatIoRes(false, res_, wios_, false); + concatIoRes(false, res_, wios_, false, [](const storage::client::WriteIO &, size_t, size_t) -> Result { + return makeError(StorageClientCode::kInvalidArg, "write result cannot contain a readable hole"); + }); } } } // namespace hf3fs::lib::agent diff --git a/src/fuse/PioV.h b/src/fuse/PioV.h index 65696566..855b86dc 100644 --- a/src/fuse/PioV.h +++ b/src/fuse/PioV.h @@ -1,6 +1,7 @@ #pragma once #include +#include #include "client/meta/MetaClient.h" #include "client/storage/StorageClient.h" @@ -8,6 +9,18 @@ namespace hf3fs::lib::agent { using flat::UserInfo; + +namespace detail { + +using ReadHoleZeroer = std::function(const storage::client::ReadIO &, size_t offset, size_t length)>; + +void finishReadResults(std::vector &res, + std::span ios, + bool allowHoles, + const ReadHoleZeroer &zeroRange); + +} // namespace detail + class PioV { public: PioV(storage::client::StorageClient &storageClient, int chunkSizeLim, std::vector &res); diff --git a/src/lib/api/CMakeLists.txt b/src/lib/api/CMakeLists.txt index a3857365..815aebed 100644 --- a/src/lib/api/CMakeLists.txt +++ b/src/lib/api/CMakeLists.txt @@ -1,6 +1,12 @@ target_add_lib(hf3fs_api client-lib-common storage-client numa rt) target_add_shared_lib(hf3fs_api_shared client-lib-common storage-client numa rt) +# Add CUDA/GDR support if enabled +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(hf3fs_api SCOPE PUBLIC) + target_add_gdr_support(hf3fs_api_shared SCOPE PUBLIC) +endif() + add_custom_command( TARGET hf3fs_api_shared POST_BUILD diff --git a/src/lib/api/UsrbIo.cc b/src/lib/api/UsrbIo.cc index dbcd3ebc..86679f08 100644 --- a/src/lib/api/UsrbIo.cc +++ b/src/lib/api/UsrbIo.cc @@ -1,9 +1,10 @@ #include #include #include -#include #include +#include #include +#include #include #include @@ -11,6 +12,7 @@ #include "common/utils/Duration.h" #include "common/utils/Path.h" #include "fuse/IoRing.h" +#include "lib/api/UsrbIoGdrInternal.h" #include "lib/api/fuse.h" #include "lib/api/hf3fs.h" #include "lib/api/hf3fs_usrbio.h" @@ -31,6 +33,25 @@ struct Hf3fsLibAliveness { static Hf3fsLibAliveness alive; +extern "C" int hf3fs_ensure_iov_mount_fd_internal(const char *hf3fs_mount_point) { + if (!hf3fs_mount_point || !*hf3fs_mount_point) { + return -EINVAL; + } + + std::lock_guard lock(alive.mtx); + if (alive.mountFds.contains(hf3fs_mount_point)) { + return 0; + } + + auto fd = open(fmt::format("{}/3fs-virt/iovs", hf3fs_mount_point).c_str(), O_DIRECTORY); + if (fd < 0) { + return -errno; + } + alive.mountFds.emplace(hf3fs_mount_point, fd); + XLOGF(INFO, "fd {} for mount {}", fd, hf3fs_mount_point); + return 0; +} + bool hf3fs_is_hf3fs(int fd) { uint32_t magic = 0; @@ -117,7 +138,7 @@ int hf3fs_iovcreate_general(struct hf3fs_iov *iov, int priority = 0, int timeout = 0, uint64_t flags = 0) { - if (!iov) { + if (!iov || !hf3fs_mount_point || !*hf3fs_mount_point) { return -EINVAL; } @@ -152,6 +173,11 @@ int hf3fs_iovcreate_general(struct hf3fs_iov *iov, is_io_ring && priority != 0 ? fmt::format(".p{}", priority < 0 ? 'h' : 'l') : std::string(), is_io_ring ? fmt::format(".t{}", timeout) : std::string(), is_io_ring && flags != 0 ? fmt::format(".f{:b}", flags) : std::string()); + auto aliveResult = hf3fs_ensure_iov_mount_fd_internal(hf3fs_mount_point); + if (aliveResult != 0) { + XLOGF(ERR, "failed to hold iovs directory for mount '{}': {}", hf3fs_mount_point, strerror(-aliveResult)); + return aliveResult; + } auto lres = symlink(target.c_str(), link.c_str()); if (lres < 0) { XLOGF(ERR, "failed to register iov '{}' to hf3fs '{}'", target, link); @@ -172,13 +198,6 @@ int hf3fs_iovcreate_general(struct hf3fs_iov *iov, succ = true; - std::lock_guard lock(alive.mtx); - if (alive.mountFds.find(hf3fs_mount_point) == alive.mountFds.end()) { - auto fd = open(fmt::format("{}/3fs-virt/iovs", hf3fs_mount_point).c_str(), O_DIRECTORY); - alive.mountFds[hf3fs_mount_point] = fd; - XLOGF(INFO, "fd {} for mount {}", fd, hf3fs_mount_point); - } - return 0; } @@ -222,15 +241,40 @@ void hf3fs_iovdestroy_general(struct hf3fs_iov *iov, } int hf3fs_iovcreate(struct hf3fs_iov *iov, const char *hf3fs_mount_point, size_t size, size_t block_size, int numa) { + // numa < 0 means "don't bind to any NUMA node" (original semantics) + // For device memory, use hf3fs_iovcreate_device() instead return hf3fs_iovcreate_general(iov, hf3fs_mount_point, size, block_size, numa, false, true, 0); } +int hf3fs_iovcreate_device(struct hf3fs_iov *iov, + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id) { + if (block_size != 0) { + XLOGF(ERR, "device iov does not support block_size: {}", block_size); + return -EINVAL; + } +#ifdef HF3FS_ENABLE_GDR + if (hf3fs_gdr_available()) { + XLOGF(DBG, "Using GDR path for device {}", device_id); + return hf3fs_iovcreate_gpu_internal(iov, hf3fs_mount_point, size, block_size, device_id); + } +#endif + XLOGF(DBG, "CUDA/GDR is unavailable for device {}", device_id); + return -ENOTSUP; +} + int hf3fs_iovopen(struct hf3fs_iov *iov, const uint8_t id[16], const char *hf3fs_mount_point, size_t size, size_t block_size, int numa) { + // numa < 0 means "don't bind to any NUMA node" (original semantics) + // For device memory, use hf3fs_iovopen_device() instead + + // Host memory path hf3fs::Uuid uuid; memcpy(uuid.data, id, sizeof(uuid.data)); @@ -282,12 +326,55 @@ int hf3fs_iovopen(struct hf3fs_iov *iov, return 0; } +// Only compiled when HF3FS_ENABLE_GDR — matches header guard. +#ifdef HF3FS_ENABLE_GDR +int hf3fs_iovopen_device(struct hf3fs_iov *iov, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id) { + if (block_size != 0) { + XLOGF(ERR, "device iov does not support block_size: {}", block_size); + return -EINVAL; + } + if (hf3fs_gdr_available()) { + XLOGF(DBG, "Using GDR path for iovopen_device, device {}", device_id); + return hf3fs_iovopen_gpu_internal(iov, id, hf3fs_mount_point, size, block_size, device_id); + } + return -ENOTSUP; +} +#endif + void hf3fs_iovunlink(struct hf3fs_iov *iov) { + if (!iov || !iov->iovh) { + return; + } + +#ifdef HF3FS_ENABLE_GDR + if (hf3fs_iov_is_gpu_internal(iov)) { + auto result = hf3fs_iovunlink_gpu_internal(iov); + if (result != 0) { + XLOGF(ERR, "failed to unlink GPU iov publication: {}", strerror(-result)); + } + return; + } +#endif + auto *shm = static_cast(iov->iovh); shm->maybeUnlinkShm(); } -void hf3fs_iovdestroy(struct hf3fs_iov *iov) { hf3fs_iovdestroy_general(iov, false, true, 0); } +void hf3fs_iovdestroy(struct hf3fs_iov *iov) { +#ifdef HF3FS_ENABLE_GDR + if (iov && hf3fs_iov_is_gpu_internal(iov)) { + hf3fs_iovdestroy_gpu_internal(iov); + return; + } +#endif + + hf3fs_iovdestroy_general(iov, false, true, 0); +} size_t hf3fs_ior_size(int entries) { return hf3fs::fuse::IoRing::bytesRequired(entries); } @@ -302,6 +389,10 @@ int hf3fs_iovwrap(struct hf3fs_iov *iov, return -EINVAL; } + // numa < 0 means "don't bind to any NUMA node" (original semantics) + // For device memory, use hf3fs_iovwrap_device() instead + + // Host memory path if (strlen(hf3fs_mount_point) >= sizeof(iov->mount_point)) { XLOGF(ERR, "mount point too long '{}'", hf3fs_mount_point); return -EINVAL; @@ -324,6 +415,27 @@ int hf3fs_iovwrap(struct hf3fs_iov *iov, return 0; } +// Only compiled when HF3FS_ENABLE_GDR — matches header guard. +#ifdef HF3FS_ENABLE_GDR +int hf3fs_iovwrap_device(struct hf3fs_iov *iov, + void *device_ptr, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id) { + if (block_size != 0) { + XLOGF(ERR, "device iov does not support block_size: {}", block_size); + return -EINVAL; + } + if (hf3fs_gdr_available()) { + XLOGF(DBG, "Using GDR path for iovwrap_device, device {}", device_id); + return hf3fs_iovwrap_gpu_internal(iov, device_ptr, id, hf3fs_mount_point, size, block_size, device_id); + } + return -ENOTSUP; +} +#endif + struct Hf3fsIorHandle { std::unique_ptr ior; sem_t *submitSem; @@ -621,10 +733,23 @@ int hf3fs_prep_io(const struct hf3fs_ior *ior, size_t off, uint64_t len, const void *userdata) { - auto p = (uint8_t *)ptr; - auto afd = abs(fd); - if (!ior || !ior->iorh || read != ior->for_read || !iov || len <= 0 || !iov->base || p < iov->base || - p + len > iov->base + iov->size || afd >= (int)regfds.size()) { + if (!ior || !ior->iorh || read != ior->for_read || !iov || !ptr || !iov->base || len == 0 || + fd == std::numeric_limits::min()) { + return -EINVAL; + } + + auto base = reinterpret_cast(iov->base); + auto address = reinterpret_cast(ptr); + if (iov->size > std::numeric_limits::max() - base || address < base) { + return -EINVAL; + } + auto bufferOffset = address - base; + if (bufferOffset > iov->size || len > iov->size - bufferOffset) { + return -EINVAL; + } + + auto afd = fd < 0 ? -fd : fd; + if (afd >= (int)regfds.size()) { return -EINVAL; } @@ -648,7 +773,7 @@ int hf3fs_prep_io(const struct hf3fs_ior *ior, auto &args = ring.ringSection[*idx]; memcpy(args.bufId, iov->id, sizeof(iov->id)); - args.bufOff = p - iov->base; + args.bufOff = static_cast(bufferOffset); args.fileIid = regfd->iid.u64(); args.fileOff = off; args.ioLen = len; @@ -805,3 +930,40 @@ int hf3fs_punchhole(int fd, int n, const size_t *start, const size_t *end, size_ } return 0; } + +enum hf3fs_mem_type hf3fs_iov_mem_type(const struct hf3fs_iov *iov) { + if (!iov) { + return HF3FS_MEM_HOST; + } +#ifdef HF3FS_ENABLE_GDR + if (hf3fs_iov_is_gpu_internal(iov)) { + return HF3FS_MEM_DEVICE; + } +#endif + return HF3FS_MEM_HOST; +} + +int hf3fs_iov_device_id(const struct hf3fs_iov *iov) { + if (!iov) { + return -1; + } +#ifdef HF3FS_ENABLE_GDR + if (hf3fs_iov_is_gpu_internal(iov)) { + return hf3fs_iov_gpu_device_internal(iov); + } +#endif + return -1; +} + +int hf3fs_iovsync(const struct hf3fs_iov *iov, int direction) { + if (!iov) { + return -EINVAL; + } +#ifdef HF3FS_ENABLE_GDR + if (hf3fs_iov_is_gpu_internal(iov)) { + return hf3fs_iovsync_gpu_internal(iov, direction); + } +#endif + (void)direction; + return 0; +} diff --git a/src/lib/api/UsrbIo.md b/src/lib/api/UsrbIo.md index 3ecf72b4..39487d28 100644 --- a/src/lib/api/UsrbIo.md +++ b/src/lib/api/UsrbIo.md @@ -36,7 +36,7 @@ int hf3fs_iorcreate4(struct hf3fs_ior *ior, - **for_read**: `true` if this Ior handles read requests, `false` if this Ior handles write requests. An Ior cannot handle read requests and write requests simultaneously. - **io_depth**: `0` for no control with I/O depth. If greater than 0, then only when `io_depth` I/O requests are in queue, they will be issued to server as a batch. If smaller than 0, then USRBIO will wait for at most `-io_depth` I/O requests are in queue and issue them in one batch. - **timeout**: Maximum wait time for batching when `io_depth` < 0. -- **numa**: Numa ID for Ior shared memory. `-1` for current process numa ID. +- **numa**: Numa ID for Ior shared memory. `-1` for no NUMA binding. - **flags**: A flag composed of OR-ed bits to specify special behaviors. #### Return Value @@ -82,7 +82,7 @@ int hf3fs_iovcreate(struct hf3fs_iov *iov, - **hf3fs_mount_point**: Mount point for 3FS. This parameter is used to distinguish 3FS clusters, enabling a single machine to mount multiple 3FS instances. - **size**: Shared memory size for this Iov. - **block_size**: If not `0`, this function will allocate multiple shared memory blocks, each sized no larger than `block_size`. `0` for allocate a single large shared memory. All IOs on this Iov should not span across the block margin. This parameter is for optimization on IB register time. -- **numa**: Numa ID for Ior shared memory. `-1` for current process numa ID. +- **numa**: Numa ID for Iov shared memory. `-1` for no NUMA binding. For device memory, use `hf3fs_iovcreate_device()`. #### Return Value - If success, return 0. @@ -250,7 +250,7 @@ constexpr uint64_t BLOCK_SIZE = (32 << 20); int main() { struct hf3fs_ior ior; hf3fs_iorcreate4(&ior, "/hf3fs/mount/point", NUM_IOS, true, 0, 0, -1, 0); - + struct hf3fs_iov iov; hf3fs_iovcreate(&iov, "/hf3fs/mount/point", NUM_IOS * BLOCK_SIZE, 0, -1); diff --git a/src/lib/api/UsrbIoGdr.cc b/src/lib/api/UsrbIoGdr.cc new file mode 100644 index 00000000..237c4ed7 --- /dev/null +++ b/src/lib/api/UsrbIoGdr.cc @@ -0,0 +1,506 @@ +/** + * GPU Direct RDMA (GDR) user API implementation. + * + * The application process owns only CUDA allocation/IPC publication. GPU + * memory registration is performed by the FUSE process when the publication + * symlink is resolved. + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "UsrbIoGdrInternal.h" +#include "hf3fs_usrbio.h" + +#ifdef HF3FS_ENABLE_GDR +#include +#endif + +#include "common/cuda/CudaMemory.h" +#include "common/utils/Uuid.h" +#include "lib/common/CudaIpcMemory.h" +#include "lib/common/GdrUri.h" + +namespace { + +static_assert(hf3fs::lib::kCudaIpcHandleBytes == hf3fs::lib::kGdrIpcHandleBytes); + +struct GpuIovHandle { + int deviceId = -1; + + void *allocationBase = nullptr; + size_t allocationSize = 0; + void *viewPtr = nullptr; + size_t viewOffset = 0; + size_t viewSize = 0; + + bool ownsMemory = false; + bool ownsPublication = false; + std::unique_ptr importedMapping; + hf3fs::lib::CudaIpcHandle ipcHandle{}; + + ~GpuIovHandle() { + importedMapping.reset(); + if (ownsMemory && allocationBase) { +#ifdef HF3FS_ENABLE_GDR + auto guard = hf3fs::cuda::ScopedDevice::create(deviceId); + if (!guard) { + XLOGF(WARN, "Failed to select CUDA device {} before freeing GPU iov: {}", deviceId, guard.error()); + return; + } + auto error = cudaFree(allocationBase); + if (error != cudaSuccess) { + XLOGF(WARN, "cudaFree failed for GPU iov: {}", cudaGetErrorString(error)); + } +#endif + } + } +}; + +std::mutex gGpuIovMutex; +std::unordered_map> gGpuIovHandles; + +int resultToErrno(const hf3fs::Status &status) { + switch (status.code()) { + case hf3fs::StatusCode::kInvalidArg: + return -EINVAL; + case hf3fs::StatusCode::kNotImplemented: + return -ENOTSUP; + case hf3fs::StatusCode::kNotEnoughMemory: + return -ENOMEM; + default: + return -EIO; + } +} + +std::string gpuIovLink(const struct hf3fs_iov &iov, int deviceId) { + hf3fs::Uuid uuid; + std::memcpy(uuid.data, iov.id, sizeof(uuid.data)); + return fmt::format("{}/3fs-virt/iovs/{}.gdr.d{}", iov.mount_point, uuid.toHexString(), deviceId); +} + +int createGpuIovSymlink(const struct hf3fs_iov &iov, const GpuIovHandle &handle) { + auto target = hf3fs::lib::formatGdrUri(handle.deviceId, + handle.allocationSize, + handle.viewOffset, + handle.viewSize, + handle.ipcHandle.data(), + handle.ipcHandle.size()); + if (target.empty()) { + return -EINVAL; + } + + auto link = gpuIovLink(iov, handle.deviceId); + if (symlink(target.c_str(), link.c_str()) < 0) { + auto error = errno; + XLOGF(WARN, "Failed to publish GPU iov {} -> {}: {}", link, target, strerror(error)); + return -error; + } + + XLOGF(DBG, "Published GPU iov {} -> {}", link, target); + return 0; +} + +int removeGpuIovSymlink(const struct hf3fs_iov &iov, int deviceId) { + auto link = gpuIovLink(iov, deviceId); + if (unlink(link.c_str()) == 0) { + return 0; + } + + auto error = errno; + if (error != ENOENT) { + XLOGF(WARN, "Failed to unlink GPU iov {}: {}", link, strerror(error)); + } + return -error; +} + +struct GpuIovSnapshot { + int deviceId; + bool ownsPublication; +}; + +std::optional getGpuHandleSnapshot(const struct hf3fs_iov *iov) { + if (!iov || !iov->iovh) { + return std::nullopt; + } + std::lock_guard lock(gGpuIovMutex); + auto it = gGpuIovHandles.find(iov->iovh); + if (it == gGpuIovHandles.end()) { + return std::nullopt; + } + return GpuIovSnapshot{it->second->deviceId, it->second->ownsPublication}; +} + +std::unique_ptr unregisterGpuIov(const struct hf3fs_iov *iov) { + if (!iov || !iov->iovh) { + return nullptr; + } + std::lock_guard lock(gGpuIovMutex); + auto it = gGpuIovHandles.find(iov->iovh); + if (it == gGpuIovHandles.end()) { + return nullptr; + } + auto handle = std::move(it->second); + gGpuIovHandles.erase(it); + return handle; +} + +int finalizeGpuIov(struct hf3fs_iov *iov, + std::unique_ptr handle, + const uint8_t id[16], + const char *mountPoint, + size_t blockSize) { + if (!iov || !handle || !id || !mountPoint || std::strlen(mountPoint) >= sizeof(iov->mount_point)) { + return -EINVAL; + } + + struct hf3fs_iov result{}; + result.base = static_cast(handle->viewPtr); + result.iovh = handle.get(); + std::memcpy(result.id, id, sizeof(result.id)); + std::strcpy(result.mount_point, mountPoint); + result.size = handle->viewSize; + result.block_size = blockSize; + result.numa = -1; + + if (handle->ownsPublication) { + auto aliveResult = hf3fs_ensure_iov_mount_fd_internal(mountPoint); + if (aliveResult != 0) { + XLOGF(ERR, "Failed to hold iovs directory for mount {}: {}", mountPoint, strerror(-aliveResult)); + return aliveResult; + } + } + + auto *key = handle.get(); + try { + std::lock_guard lock(gGpuIovMutex); + auto [it, inserted] = gGpuIovHandles.try_emplace(key); + if (!inserted) { + return -EEXIST; + } + it->second = std::move(handle); + } catch (const std::bad_alloc &) { + return -ENOMEM; + } + + if (key->ownsPublication) { + auto publishResult = createGpuIovSymlink(result, *key); + if (publishResult != 0) { + auto rollback = unregisterGpuIov(&result); + XLOGF_IF(FATAL, !rollback, "GPU iov handle disappeared during publication rollback"); + return publishResult; + } + } + + *iov = result; + return 0; +} + +int registerAndPublishGpuIov(struct hf3fs_iov *iov, + std::unique_ptr handle, + const uint8_t id[16], + const char *mountPoint, + size_t blockSize) { + handle->ownsPublication = true; + return finalizeGpuIov(iov, std::move(handle), id, mountPoint, blockSize); +} + +int validateGpuDevice(int deviceId) { + auto available = hf3fs::cuda::supportsIpc(deviceId); + if (!available) { + XLOGF(ERR, "Failed to query CUDA device {}: {}", deviceId, available.error()); + return resultToErrno(available.error()); + } + return *available ? 0 : -ENODEV; +} + +int allocateGpuMemory(size_t size, int deviceId, void **devicePtr) { + if (!devicePtr || size == 0) { + return -EINVAL; + } +#ifdef HF3FS_ENABLE_GDR + auto guard = hf3fs::cuda::ScopedDevice::create(deviceId); + if (!guard) { + XLOGF(ERR, "Failed to select CUDA device {}: {}", deviceId, guard.error()); + return -ENODEV; + } + auto error = cudaMalloc(devicePtr, size); + if (error != cudaSuccess) { + XLOGF(ERR, "cudaMalloc({}) failed: {}", size, cudaGetErrorString(error)); + return error == cudaErrorMemoryAllocation ? -ENOMEM : -EIO; + } + return 0; +#else + (void)deviceId; + *devicePtr = nullptr; + return -ENOTSUP; +#endif +} + +} // namespace + +extern "C" { + +bool hf3fs_gdr_available(void) { + auto count = hf3fs::cuda::deviceCount(); + if (!count) { + return false; + } + for (int deviceId = 0; deviceId < *count; ++deviceId) { + auto available = hf3fs::cuda::supportsIpc(deviceId); + if (available && *available) { + return true; + } + } + return false; +} + +int hf3fs_gdr_device_count(void) { + auto count = hf3fs::cuda::deviceCount(); + return count ? *count : 0; +} + +int hf3fs_iovcreate_gpu_internal(struct hf3fs_iov *iov, + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id) { + if (!iov || !hf3fs_mount_point || size == 0) { + return -EINVAL; + } + if (block_size != 0) { + return -EINVAL; + } + auto deviceResult = validateGpuDevice(gpu_device_id); + if (deviceResult != 0) { + return deviceResult; + } + + void *devicePtr = nullptr; + auto allocationResult = allocateGpuMemory(size, gpu_device_id, &devicePtr); + if (allocationResult != 0) { + return allocationResult; + } + + auto handle = std::make_unique(); + handle->deviceId = gpu_device_id; + handle->allocationBase = devicePtr; + handle->allocationSize = size; + handle->viewPtr = devicePtr; + handle->viewSize = size; + handle->ownsMemory = true; + + auto ipcExport = hf3fs::lib::exportCudaIpcMemory(devicePtr, size, gpu_device_id); + if (!ipcExport) { + XLOGF(ERR, "Failed to export allocated GPU iov: {}", ipcExport.error()); + return resultToErrno(ipcExport.error()); + } + if (ipcExport->allocationBase != devicePtr || ipcExport->offset != 0) { + XLOGF(ERR, "cudaMalloc returned a pointer that is not the CUDA allocation base"); + return -EIO; + } + handle->allocationBase = ipcExport->allocationBase; + handle->allocationSize = ipcExport->allocationSize; + handle->viewOffset = ipcExport->offset; + handle->ipcHandle = ipcExport->ipcHandle; + + auto uuid = hf3fs::Uuid::random(); + return registerAndPublishGpuIov(iov, + std::move(handle), + reinterpret_cast(uuid.data), + hf3fs_mount_point, + block_size); +} + +int hf3fs_iovopen_gpu_internal(struct hf3fs_iov *iov, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id) { + if (!iov || !id || !hf3fs_mount_point || size == 0 || std::strlen(hf3fs_mount_point) >= sizeof(iov->mount_point)) { + return -EINVAL; + } + if (block_size != 0) { + return -EINVAL; + } + auto deviceResult = validateGpuDevice(gpu_device_id); + if (deviceResult != 0) { + return deviceResult; + } + + hf3fs::Uuid uuid; + std::memcpy(uuid.data, id, sizeof(uuid.data)); + auto link = fmt::format("{}/3fs-virt/iovs/{}.gdr.d{}", hf3fs_mount_point, uuid.toHexString(), gpu_device_id); + + char target[512]; + auto length = readlink(link.c_str(), target, sizeof(target) - 1); + if (length < 0) { + return -errno; + } + target[length] = '\0'; + + auto parsed = hf3fs::lib::parseGdrUri(target); + if (!parsed || parsed->deviceId != gpu_device_id || parsed->size != size) { + return -EINVAL; + } + + hf3fs::lib::CudaIpcHandle ipcHandle; + std::copy(parsed->ipcHandle.begin(), parsed->ipcHandle.end(), ipcHandle.begin()); + auto mapping = hf3fs::lib::CudaIpcMapping::open(gpu_device_id, ipcHandle, parsed->allocationSize); + if (!mapping) { + XLOGF(ERR, "Failed to import GPU iov {}: {}", uuid.toHexString(), mapping.error()); + return resultToErrno(mapping.error()); + } + auto view = mapping->view(parsed->offset, parsed->size); + if (!view) { + return resultToErrno(view.error()); + } + + auto handle = std::make_unique(); + handle->deviceId = gpu_device_id; + handle->allocationBase = mapping->allocationBase(); + handle->allocationSize = parsed->allocationSize; + handle->viewPtr = *view; + handle->viewOffset = parsed->offset; + handle->viewSize = parsed->size; + handle->ipcHandle = ipcHandle; + handle->importedMapping = std::make_unique(std::move(*mapping)); + + return finalizeGpuIov(iov, std::move(handle), id, hf3fs_mount_point, block_size); +} + +int hf3fs_iovwrap_gpu_internal(struct hf3fs_iov *iov, + void *gpu_ptr, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id) { + if (!iov || !gpu_ptr || !id || !hf3fs_mount_point || size == 0) { + return -EINVAL; + } + if (block_size != 0) { + return -EINVAL; + } + auto deviceResult = validateGpuDevice(gpu_device_id); + if (deviceResult != 0) { + return deviceResult; + } + + auto ipcExport = hf3fs::lib::exportCudaIpcMemory(gpu_ptr, size, gpu_device_id); + if (!ipcExport) { + XLOGF(ERR, "Failed to export wrapped GPU iov: {}", ipcExport.error()); + return resultToErrno(ipcExport.error()); + } + + auto handle = std::make_unique(); + handle->deviceId = gpu_device_id; + handle->allocationBase = ipcExport->allocationBase; + handle->allocationSize = ipcExport->allocationSize; + handle->viewPtr = gpu_ptr; + handle->viewOffset = ipcExport->offset; + handle->viewSize = size; + handle->ipcHandle = ipcExport->ipcHandle; + + return registerAndPublishGpuIov(iov, std::move(handle), id, hf3fs_mount_point, block_size); +} + +int hf3fs_iovunlink_gpu_internal(struct hf3fs_iov *iov) { + if (!iov || !iov->iovh) { + return 0; + } + + std::lock_guard lock(gGpuIovMutex); + auto it = gGpuIovHandles.find(iov->iovh); + if (it == gGpuIovHandles.end() || !it->second->ownsPublication) { + return 0; + } + + auto result = removeGpuIovSymlink(*iov, it->second->deviceId); + if (result == 0 || result == -ENOENT) { + it->second->ownsPublication = false; + return 0; + } + return result; +} + +int hf3fs_iovunlink_gpu_publication_internal(const struct hf3fs_iov *iov, int device_id, bool owns_publication) { + if (!iov || !owns_publication) { + return 0; + } + return removeGpuIovSymlink(*iov, device_id); +} + +void hf3fs_iovdestroy_gpu_internal(struct hf3fs_iov *iov) { + if (!iov || !iov->iovh) { + return; + } + + std::unique_ptr handle; + { + std::lock_guard lock(gGpuIovMutex); + auto it = gGpuIovHandles.find(iov->iovh); + if (it == gGpuIovHandles.end()) { + return; + } + + if (it->second->ownsPublication) { + auto unlinkResult = removeGpuIovSymlink(*iov, it->second->deviceId); + if (unlinkResult != 0 && unlinkResult != -ENOENT) { + XLOGF(ERR, "GPU iov destroy retained handle after publication unlink failed: {}", strerror(-unlinkResult)); + return; + } + it->second->ownsPublication = false; + } + + handle = std::move(it->second); + gGpuIovHandles.erase(it); + } + + handle.reset(); + std::memset(iov, 0, sizeof(*iov)); +} + +bool hf3fs_iov_is_gpu_internal(const struct hf3fs_iov *iov) { return getGpuHandleSnapshot(iov).has_value(); } + +int hf3fs_iov_gpu_device_internal(const struct hf3fs_iov *iov) { + auto handle = getGpuHandleSnapshot(iov); + return handle ? handle->deviceId : -1; +} + +int hf3fs_iovsync_gpu_internal(const struct hf3fs_iov *iov, int direction) { + auto handle = getGpuHandleSnapshot(iov); + if (!handle) { + return -EINVAL; + } +#ifdef HF3FS_ENABLE_GDR + auto guard = hf3fs::cuda::ScopedDevice::create(handle->deviceId); + if (!guard) { + return -ENODEV; + } + auto error = cudaDeviceSynchronize(); + if (error != cudaSuccess) { + XLOGF(ERR, "cudaDeviceSynchronize failed: {}", cudaGetErrorString(error)); + return -EIO; + } +#else + (void)direction; + return -ENOTSUP; +#endif + (void)direction; + return 0; +} + +} // extern "C" diff --git a/src/lib/api/UsrbIoGdrInternal.h b/src/lib/api/UsrbIoGdrInternal.h new file mode 100644 index 00000000..9b750c5e --- /dev/null +++ b/src/lib/api/UsrbIoGdrInternal.h @@ -0,0 +1,45 @@ +#pragma once + +#include +#include + +#include "lib/api/hf3fs_usrbio.h" + +#ifdef __cplusplus +extern "C" { +#endif + +bool hf3fs_gdr_available(void); +int hf3fs_ensure_iov_mount_fd_internal(const char *hf3fs_mount_point); + +int hf3fs_iovcreate_gpu_internal(struct hf3fs_iov *iov, + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id); + +int hf3fs_iovopen_gpu_internal(struct hf3fs_iov *iov, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id); + +int hf3fs_iovwrap_gpu_internal(struct hf3fs_iov *iov, + void *gpu_ptr, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int gpu_device_id); + +int hf3fs_iovunlink_gpu_internal(struct hf3fs_iov *iov); +int hf3fs_iovunlink_gpu_publication_internal(const struct hf3fs_iov *iov, int device_id, bool owns_publication); +void hf3fs_iovdestroy_gpu_internal(struct hf3fs_iov *iov); +bool hf3fs_iov_is_gpu_internal(const struct hf3fs_iov *iov); +int hf3fs_iov_gpu_device_internal(const struct hf3fs_iov *iov); +int hf3fs_iovsync_gpu_internal(const struct hf3fs_iov *iov, int direction); + +#ifdef __cplusplus +} +#endif diff --git a/src/lib/api/hf3fs_usrbio.h b/src/lib/api/hf3fs_usrbio.h index c8c237e5..15362f41 100644 --- a/src/lib/api/hf3fs_usrbio.h +++ b/src/lib/api/hf3fs_usrbio.h @@ -68,6 +68,10 @@ int hf3fs_extract_mount_point(char *hf3fs_mount_point, int size, const char *pat // iov ptr itself should be allocated by caller, it could be on either stack or heap // the pointer hf3fs_mount_point will be copied into the corresponding field in hf3fs_iov // it should not be too long +// numa: NUMA node hint for host memory allocation +// >= 0: allocate on specified NUMA node +// < 0: no NUMA binding (e.g., -1) +// For device memory (GPU etc.), use hf3fs_iovcreate_device() instead. int hf3fs_iovcreate(struct hf3fs_iov *iov, const char *hf3fs_mount_point, size_t size, size_t block_size, int numa); int hf3fs_iovopen(struct hf3fs_iov *iov, const uint8_t id[16], @@ -75,6 +79,8 @@ int hf3fs_iovopen(struct hf3fs_iov *iov, size_t size, size_t block_size, int numa); +// GPU publication unlink failures are logged; the handle remains registered so +// hf3fs_iovunlink() or hf3fs_iovdestroy() can retry cleanup. void hf3fs_iovunlink(struct hf3fs_iov *iov); // iov ptr itself will not be freed void hf3fs_iovdestroy(struct hf3fs_iov *iov); @@ -84,6 +90,7 @@ void hf3fs_iovdestroy(struct hf3fs_iov *iov); // iov ptr itself should be allocated by caller // the wrapped iov cannot be destroyed by hf3fs_iovdestroy() // the underlying base ptr should not be unmapped when the iov (or its corresponding ior) is still being used +// For device memory, use hf3fs_iovwrap_device() instead. int hf3fs_iovwrap(struct hf3fs_iov *iov, void *base, const uint8_t id[16], @@ -92,6 +99,49 @@ int hf3fs_iovwrap(struct hf3fs_iov *iov, size_t block_size, int numa); +// Device memory IOV creation (e.g., GPU via GDR) +// device_id: accelerator device index (0, 1, 2, ...) +// Returns -ENOTSUP if CUDA/GDR support or local CUDA IPC is unavailable; +// this API never falls back to host memory. +// A successful CUDA capability check does not guarantee RDMA registration; +// successful FUSE publication is the final per-IOV registration result. +// GDR v2 does not support block partitioning; block_size must be 0. +int hf3fs_iovcreate_device(struct hf3fs_iov *iov, + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); + +#ifdef HF3FS_ENABLE_GDR +// Open an existing device memory IOV by UUID (cross-process reopen) +// The importer borrows the publisher's publication; destroy closes only the +// importer's CUDA IPC mapping and does not unlink the publication. +// Returns -ENOTSUP when local CUDA IPC is not available. +// GDR v2 does not support block partitioning; block_size must be 0. +// Only declared when HF3FS_ENABLE_GDR is defined at compile time. +int hf3fs_iovopen_device(struct hf3fs_iov *iov, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); + +// Wrap externally-allocated device memory as IOV +// device_ptr may point to a subrange of a CUDA allocation and must remain valid +// until the publication and all importers are gone. CUDA IPC shares the entire +// underlying allocation even when only this subrange is published. +// Returns -ENOTSUP when local CUDA IPC is not available. +// GDR v2 does not support block partitioning; block_size must be 0. +// Only declared when HF3FS_ENABLE_GDR is defined at compile time. +int hf3fs_iovwrap_device(struct hf3fs_iov *iov, + void *device_ptr, + const uint8_t id[16], + const char *hf3fs_mount_point, + size_t size, + size_t block_size, + int device_id); +#endif + // calculate required memory size with wanted entries // the calculated size can be used to create the underlying iov size_t hf3fs_ior_size(int entries); @@ -168,6 +218,23 @@ int hf3fs_wait_for_ios(const struct hf3fs_ior *ior, int hf3fs_hardlink(const char *target, const char *link_name); int hf3fs_punchhole(int fd, int n, const size_t *start, const size_t *end, size_t flags); + +enum hf3fs_mem_type { + HF3FS_MEM_HOST = 0, + HF3FS_MEM_DEVICE = 1, + HF3FS_MEM_MANAGED = 2, +}; + +// Reports application-local CUDA IPC/device capability only. It does not prove +// that FUSE can import the allocation or register an RDMA MR; successful +// create/wrap publication is the final per-IOV result. +bool hf3fs_gdr_available(void); +// Returns the number of locally visible CUDA devices; it does not probe FUSE RDMA registration. +int hf3fs_gdr_device_count(void); +enum hf3fs_mem_type hf3fs_iov_mem_type(const struct hf3fs_iov *iov); +int hf3fs_iov_device_id(const struct hf3fs_iov *iov); +int hf3fs_iovsync(const struct hf3fs_iov *iov, int direction); + #ifdef __cplusplus } #endif diff --git a/src/lib/common/CMakeLists.txt b/src/lib/common/CMakeLists.txt index 064e2124..ed459abe 100644 --- a/src/lib/common/CMakeLists.txt +++ b/src/lib/common/CMakeLists.txt @@ -1 +1,6 @@ target_add_lib(client-lib-common common numa rt) + +# Add CUDA/GDR support if enabled +if(HF3FS_ENABLE_GDR) + target_add_gdr_support(client-lib-common SCOPE PRIVATE) +endif() diff --git a/src/lib/common/CudaIpcMemory.cc b/src/lib/common/CudaIpcMemory.cc new file mode 100644 index 00000000..5be79e46 --- /dev/null +++ b/src/lib/common/CudaIpcMemory.cc @@ -0,0 +1,144 @@ +#include "lib/common/CudaIpcMemory.h" + +#include +#include +#include +#include +#include + +#ifdef HF3FS_ENABLE_GDR +#include +#endif + +#include "common/cuda/CudaMemory.h" + +namespace hf3fs::lib { + +Result exportCudaIpcMemory(void *devicePtr, size_t viewSize, int deviceId) { +#ifdef HF3FS_ENABLE_GDR + static_assert(sizeof(cudaIpcMemHandle_t) == kCudaIpcHandleBytes); + + if (!devicePtr || viewSize == 0 || deviceId < 0) { + return makeError(StatusCode::kInvalidArg, "invalid CUDA IPC export parameters"); + } + + auto available = cuda::supportsIpc(deviceId); + RETURN_ON_ERROR(available); + if (!*available) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, "CUDA device {} does not support IPC", deviceId); + } + + auto view = cuda::prepareGdrMemory(devicePtr, viewSize, deviceId); + RETURN_ON_ERROR(view); + auto guard = cuda::ScopedDevice::create(deviceId); + RETURN_ON_ERROR(guard); + + cudaIpcMemHandle_t runtimeHandle; + auto runtimeResult = cudaIpcGetMemHandle(&runtimeHandle, view->allocationBase); + if (runtimeResult != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, "cudaIpcGetMemHandle failed: {}", cudaGetErrorString(runtimeResult)); + } + + CudaIpcExport result; + result.allocationBase = view->allocationBase; + result.allocationSize = view->allocationSize; + result.offset = view->offset; + std::memcpy(result.ipcHandle.data(), &runtimeHandle, sizeof(runtimeHandle)); + return result; +#else + (void)devicePtr; + (void)viewSize; + (void)deviceId; + return makeError(StatusCode::kNotImplemented, "CUDA IPC is disabled in this build"); +#endif +} + +Result CudaIpcMapping::open(int deviceId, const CudaIpcHandle &ipcHandle, size_t allocationSize) { +#ifdef HF3FS_ENABLE_GDR + static_assert(sizeof(cudaIpcMemHandle_t) == kCudaIpcHandleBytes); + + if (deviceId < 0 || allocationSize == 0) { + return makeError(StatusCode::kInvalidArg, "invalid CUDA IPC import parameters"); + } + auto guard = cuda::ScopedDevice::create(deviceId); + RETURN_ON_ERROR(guard); + + cudaIpcMemHandle_t runtimeHandle; + std::memcpy(&runtimeHandle, ipcHandle.data(), sizeof(runtimeHandle)); + + void *importedBase = nullptr; + auto runtimeResult = cudaIpcOpenMemHandle(&importedBase, runtimeHandle, cudaIpcMemLazyEnablePeerAccess); + if (runtimeResult != cudaSuccess) { + return MAKE_ERROR_F(StatusCode::kIOError, "cudaIpcOpenMemHandle failed: {}", cudaGetErrorString(runtimeResult)); + } + + CudaIpcMapping mapping(deviceId, importedBase, allocationSize); + auto view = cuda::inspectDeviceMemory(importedBase, allocationSize, deviceId); + RETURN_ON_ERROR(view); + if (view->allocationBase != importedBase || view->allocationSize != allocationSize || view->offset != 0) { + return MAKE_ERROR_F(StatusCode::kInvalidArg, + "imported CUDA allocation range mismatch: URI size {}, actual size {}", + allocationSize, + view->allocationSize); + } + + return Result(std::move(mapping)); +#else + (void)deviceId; + (void)ipcHandle; + (void)allocationSize; + return makeError(StatusCode::kNotImplemented, "CUDA IPC is disabled in this build"); +#endif +} + +CudaIpcMapping::CudaIpcMapping(CudaIpcMapping &&other) noexcept + : deviceId_(std::exchange(other.deviceId_, -1)), + allocationBase_(std::exchange(other.allocationBase_, nullptr)), + allocationSize_(std::exchange(other.allocationSize_, 0)) {} + +CudaIpcMapping &CudaIpcMapping::operator=(CudaIpcMapping &&other) noexcept { + if (this != &other) { + reset(); + deviceId_ = std::exchange(other.deviceId_, -1); + allocationBase_ = std::exchange(other.allocationBase_, nullptr); + allocationSize_ = std::exchange(other.allocationSize_, 0); + } + return *this; +} + +CudaIpcMapping::~CudaIpcMapping() { reset(); } + +Result CudaIpcMapping::view(size_t offset, size_t size) const { + if (!allocationBase_ || size == 0 || offset > allocationSize_ || size > allocationSize_ - offset) { + return makeError(StatusCode::kInvalidArg, "CUDA IPC view is outside the imported allocation"); + } + auto address = reinterpret_cast(allocationBase_); + if (address > std::numeric_limits::max() - offset) { + return makeError(StatusCode::kInvalidArg, "CUDA IPC view address overflows uintptr_t"); + } + return reinterpret_cast(address + offset); +} + +void CudaIpcMapping::reset() noexcept { + if (!allocationBase_) { + return; + } + + void *base = std::exchange(allocationBase_, nullptr); + allocationSize_ = 0; +#ifdef HF3FS_ENABLE_GDR + auto guard = cuda::ScopedDevice::create(deviceId_); + if (!guard) { + XLOGF(WARN, "Failed to select CUDA device {} before closing IPC mapping: {}", deviceId_, guard.error()); + } + auto error = cudaIpcCloseMemHandle(base); + if (error != cudaSuccess) { + XLOGF(WARN, "cudaIpcCloseMemHandle failed: {}", cudaGetErrorString(error)); + } +#else + (void)base; +#endif + deviceId_ = -1; +} + +} // namespace hf3fs::lib diff --git a/src/lib/common/CudaIpcMemory.h b/src/lib/common/CudaIpcMemory.h new file mode 100644 index 00000000..d0c8d549 --- /dev/null +++ b/src/lib/common/CudaIpcMemory.h @@ -0,0 +1,50 @@ +#pragma once + +#include +#include +#include + +#include "common/utils/Result.h" + +namespace hf3fs::lib { + +constexpr size_t kCudaIpcHandleBytes = 64; +using CudaIpcHandle = std::array; + +struct CudaIpcExport { + void *allocationBase = nullptr; + size_t allocationSize = 0; + size_t offset = 0; + CudaIpcHandle ipcHandle{}; +}; + +Result exportCudaIpcMemory(void *devicePtr, size_t viewSize, int deviceId); + +class CudaIpcMapping { + public: + static Result open(int deviceId, const CudaIpcHandle &ipcHandle, size_t allocationSize); + + CudaIpcMapping(const CudaIpcMapping &) = delete; + CudaIpcMapping &operator=(const CudaIpcMapping &) = delete; + CudaIpcMapping(CudaIpcMapping &&other) noexcept; + CudaIpcMapping &operator=(CudaIpcMapping &&other) noexcept; + ~CudaIpcMapping(); + + void *allocationBase() const { return allocationBase_; } + size_t allocationSize() const { return allocationSize_; } + Result view(size_t offset, size_t size) const; + + private: + CudaIpcMapping(int deviceId, void *allocationBase, size_t allocationSize) + : deviceId_(deviceId), + allocationBase_(allocationBase), + allocationSize_(allocationSize) {} + + void reset() noexcept; + + int deviceId_ = -1; + void *allocationBase_ = nullptr; + size_t allocationSize_ = 0; +}; + +} // namespace hf3fs::lib diff --git a/src/lib/common/GdrUri.cc b/src/lib/common/GdrUri.cc new file mode 100644 index 00000000..88619d78 --- /dev/null +++ b/src/lib/common/GdrUri.cc @@ -0,0 +1,151 @@ +#include "lib/common/GdrUri.h" + +#include +#include +#include +#include + +namespace hf3fs::lib { +namespace { + +constexpr std::string_view kPrefix = "gdr://v2/device/"; +constexpr std::string_view kAllocationSep = "/allocation/"; +constexpr std::string_view kOffsetSep = "/offset/"; +constexpr std::string_view kSizeSep = "/size/"; +constexpr std::string_view kIpcSep = "/ipc/"; + +std::optional parseUnsigned(std::string_view text) { + if (text.empty() || !std::all_of(text.begin(), text.end(), [](char c) { return c >= '0' && c <= '9'; })) { + return std::nullopt; + } + + size_t value = 0; + auto *begin = text.data(); + auto *end = begin + text.size(); + auto [ptr, ec] = std::from_chars(begin, end, value); + if (ec != std::errc{} || ptr != end) { + return std::nullopt; + } + return value; +} + +int hexNibble(char c) { + if (c >= '0' && c <= '9') return c - '0'; + if (c >= 'a' && c <= 'f') return c - 'a' + 10; + if (c >= 'A' && c <= 'F') return c - 'A' + 10; + return -1; +} + +bool decodeHex(std::string_view encoded, std::array &out) { + if (encoded.size() != kGdrIpcHandleBytes * 2) { + return false; + } + + for (size_t i = 0; i < kGdrIpcHandleBytes; ++i) { + int hi = hexNibble(encoded[2 * i]); + int lo = hexNibble(encoded[2 * i + 1]); + if (hi < 0 || lo < 0) { + return false; + } + out[i] = static_cast((hi << 4) | lo); + } + return true; +} + +std::string encodeHex(const uint8_t *data, size_t size) { + static constexpr char kHex[] = "0123456789abcdef"; + std::string out; + out.resize(size * 2); + for (size_t i = 0; i < size; ++i) { + out[2 * i] = kHex[(data[i] >> 4) & 0x0F]; + out[2 * i + 1] = kHex[data[i] & 0x0F]; + } + return out; +} + +} // namespace + +std::optional parseGdrUri(std::string_view uri) { + if (!uri.starts_with(kPrefix)) { + return std::nullopt; + } + + uri.remove_prefix(kPrefix.size()); + auto allocationPos = uri.find(kAllocationSep); + if (allocationPos == std::string_view::npos) { + return std::nullopt; + } + auto deviceText = uri.substr(0, allocationPos); + uri.remove_prefix(allocationPos + kAllocationSep.size()); + + auto offsetPos = uri.find(kOffsetSep); + if (offsetPos == std::string_view::npos) { + return std::nullopt; + } + auto allocationText = uri.substr(0, offsetPos); + uri.remove_prefix(offsetPos + kOffsetSep.size()); + + auto sizePos = uri.find(kSizeSep); + if (sizePos == std::string_view::npos) { + return std::nullopt; + } + auto offsetText = uri.substr(0, sizePos); + uri.remove_prefix(sizePos + kSizeSep.size()); + + auto ipcPos = uri.find(kIpcSep); + if (ipcPos == std::string_view::npos) { + return std::nullopt; + } + auto sizeText = uri.substr(0, ipcPos); + auto ipcHex = uri.substr(ipcPos + kIpcSep.size()); + + auto device = parseUnsigned(deviceText); + auto allocationSize = parseUnsigned(allocationText); + auto offset = parseUnsigned(offsetText); + auto size = parseUnsigned(sizeText); + if (!device || *device > static_cast(std::numeric_limits::max()) || !allocationSize || + *allocationSize == 0 || !offset || !size || *size == 0 || *offset > *allocationSize || + *size > *allocationSize - *offset) { + return std::nullopt; + } + + GdrUri parsed; + parsed.deviceId = static_cast(*device); + parsed.allocationSize = *allocationSize; + parsed.offset = *offset; + parsed.size = *size; + if (!decodeHex(ipcHex, parsed.ipcHandle)) { + return std::nullopt; + } + return parsed; +} + +std::string formatGdrUri(int deviceId, + size_t allocationSize, + size_t offset, + size_t size, + const uint8_t *ipcHandle, + size_t ipcHandleSize) { + if (deviceId < 0 || allocationSize == 0 || size == 0 || offset > allocationSize || size > allocationSize - offset || + ipcHandle == nullptr || ipcHandleSize != kGdrIpcHandleBytes) { + return {}; + } + + auto ipcHex = encodeHex(ipcHandle, ipcHandleSize); + std::string uri; + uri.reserve(kPrefix.size() + 20 + kAllocationSep.size() + 20 + kOffsetSep.size() + 20 + kSizeSep.size() + 20 + + kIpcSep.size() + ipcHex.size()); + uri += kPrefix; + uri += std::to_string(deviceId); + uri += kAllocationSep; + uri += std::to_string(allocationSize); + uri += kOffsetSep; + uri += std::to_string(offset); + uri += kSizeSep; + uri += std::to_string(size); + uri += kIpcSep; + uri += ipcHex; + return uri; +} + +} // namespace hf3fs::lib diff --git a/src/lib/common/GdrUri.h b/src/lib/common/GdrUri.h new file mode 100644 index 00000000..803fc51d --- /dev/null +++ b/src/lib/common/GdrUri.h @@ -0,0 +1,31 @@ +#pragma once + +#include +#include +#include +#include +#include +#include + +namespace hf3fs::lib { + +constexpr size_t kGdrIpcHandleBytes = 64; + +struct GdrUri { + int deviceId = -1; + size_t allocationSize = 0; + size_t offset = 0; + size_t size = 0; + std::array ipcHandle{}; +}; + +std::optional parseGdrUri(std::string_view uri); + +std::string formatGdrUri(int deviceId, + size_t allocationSize, + size_t offset, + size_t size, + const uint8_t *ipcHandle, + size_t ipcHandleSize); + +} // namespace hf3fs::lib diff --git a/src/lib/common/GpuShm.cc b/src/lib/common/GpuShm.cc new file mode 100644 index 00000000..14dbe1c5 --- /dev/null +++ b/src/lib/common/GpuShm.cc @@ -0,0 +1,35 @@ +#include "lib/common/GpuShm.h" + +#include "common/net/ib/RDMABuf.h" + +namespace hf3fs::lib { + +Result> GpuShmBuf::create(const CudaIpcHandle &ipcHandle, + size_t allocationSize, + size_t offset, + size_t size, + int deviceId) { + auto imported = CudaIpcMapping::open(deviceId, ipcHandle, allocationSize); + RETURN_ON_ERROR(imported); + + auto mapping = std::make_shared(std::move(*imported)); + auto view = mapping->view(offset, size); + RETURN_ON_ERROR(view); + + std::shared_ptr backingOwner = mapping; + auto rdmaBuf = + net::RDMABuf::createFromCudaBuffer(static_cast(*view), size, deviceId, std::move(backingOwner)); + RETURN_ON_ERROR(rdmaBuf); + + auto ioBuffer = storage::client::IOBuffer(std::move(*rdmaBuf)); + return std::shared_ptr(new GpuShmBuf(size, std::move(ioBuffer))); +} + +CoTryTask GpuShmBufForIO::memh(size_t len) const { + if (!buf_ || off_ > buf_->size || len > buf_->size - off_) { + co_return makeError(StatusCode::kInvalidArg, "invalid GPU buf off and/or io len"); + } + co_return buf_->ioBuffer(); +} + +} // namespace hf3fs::lib diff --git a/src/lib/common/GpuShm.h b/src/lib/common/GpuShm.h new file mode 100644 index 00000000..3d59aa25 --- /dev/null +++ b/src/lib/common/GpuShm.h @@ -0,0 +1,52 @@ +#pragma once + +#include +#include +#include +#include + +#include "client/storage/StorageClient.h" +#include "common/utils/Coroutine.h" +#include "common/utils/Result.h" +#include "lib/common/CudaIpcMemory.h" + +namespace hf3fs::lib { + +/** A CUDA IPC view imported and registered once for the lifetime of an IOV entry. */ +class GpuShmBuf { + public: + static Result> create(const CudaIpcHandle &ipcHandle, + size_t allocationSize, + size_t offset, + size_t size, + int deviceId); + + uint8_t *dataAtOffset(size_t offset) const { return ioBuffer_.dataAtOffset(offset); } + storage::client::IOBuffer *ioBuffer() { return &ioBuffer_; } + + const size_t size; + + private: + GpuShmBuf(size_t size, storage::client::IOBuffer ioBuffer) + : size(size), + ioBuffer_(std::move(ioBuffer)) {} + + storage::client::IOBuffer ioBuffer_; +}; + +/** Offset view used by the existing host/GPU IOV dispatch in the FUSE worker. */ +class GpuShmBufForIO { + public: + GpuShmBufForIO(std::shared_ptr buf, size_t off) + : buf_(std::move(buf)), + off_(off) {} + + uint8_t *ptr() const { return buf_ ? buf_->dataAtOffset(off_) : nullptr; } + CoTryTask memh(size_t len) const; + + private: + std::shared_ptr buf_; + size_t off_; +}; + +} // namespace hf3fs::lib diff --git a/src/storage/service/StorageOperator.cc b/src/storage/service/StorageOperator.cc index 36a77d29..70e2d7de 100644 --- a/src/storage/service/StorageOperator.cc +++ b/src/storage/service/StorageOperator.cc @@ -57,6 +57,24 @@ monitor::OperationRecorder syncDoneRecorder{"storage.sync_done"}; monitor::OperationRecorder storageReqRemoveChunksRecorder{"storage.req_remove_chunks"}; monitor::OperationRecorder storageRemoveRangeRecorder{"storage.remove_range"}; +Result materializeServerComputedChecksum(UpdateIO &updateIO, uint32_t featureFlags, const uint8_t *data) { + if (!BITFLAGS_CONTAIN(featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)) { + return Void{}; + } + if (!updateIO.isWrite()) { + return makeError(StatusCode::kInvalidArg, "server-computed checksum is only valid for WRITE updates"); + } + if (updateIO.checksum.type == ChecksumType::NONE) { + return makeError(StatusCode::kInvalidArg, "server-computed checksum requires a non-NONE target checksum type"); + } + if (data == nullptr && updateIO.length != 0) { + return makeError(StatusCode::kInvalidArg, "server-computed checksum requires staged write data"); + } + + updateIO.checksum = ChecksumInfo::create(updateIO.checksum.type, data, updateIO.length); + return Void{}; +} + Result StorageOperator::init(uint32_t numberOfDisks) { storageReadAvgBytes.setLambda([&] { auto totalReadBytes = totalReadBytes_.exchange(0); @@ -514,7 +532,7 @@ CoTask StorageOperator::handleUpdate(ServiceRequestContext &requestCtx } CoTask StorageOperator::doUpdate(ServiceRequestContext &requestCtx, - const UpdateIO &updateIO, + UpdateIO &updateIO, const UpdateOptions &updateOptions, uint32_t featureFlags, const std::shared_ptr &target, @@ -591,6 +609,19 @@ CoTask StorageOperator::doUpdate(ServiceRequestContext &requestCtx, } } + if (BITFLAGS_CONTAIN(featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM) && + BITFLAGS_CONTAIN(featureFlags, FeatureFlags::BYPASS_RDMAXMIT) && + !BITFLAGS_CONTAIN(featureFlags, FeatureFlags::SEND_DATA_INLINE)) { + co_return makeError(StatusCode::kInvalidArg, "cannot compute a server checksum when RDMA transfer is bypassed"); + } + + auto checksumResult = materializeServerComputedChecksum(updateIO, featureFlags, job.state().data); + if (UNLIKELY(!checksumResult)) { + XLOGF(ERR, "Failed to materialize server-computed checksum for {}: {}", updateIO, checksumResult.error()); + co_return makeError(std::move(checksumResult.error())); + } + job.updateIO().checksum = updateIO.checksum; + if (BITFLAGS_CONTAIN(featureFlags, FeatureFlags::BYPASS_DISKIO)) { job.setResult(updateIO.length); } else { diff --git a/src/storage/service/StorageOperator.h b/src/storage/service/StorageOperator.h index f5421d7e..c1a0ba99 100644 --- a/src/storage/service/StorageOperator.h +++ b/src/storage/service/StorageOperator.h @@ -27,6 +27,8 @@ namespace hf3fs::storage { struct Components; +Result materializeServerComputedChecksum(UpdateIO &updateIO, uint32_t featureFlags, const uint8_t *data); + class StorageOperator { public: class Config : public ConfigBase { @@ -106,7 +108,7 @@ class StorageOperator { TargetPtr &target); CoTask doUpdate(ServiceRequestContext &requestCtx, - const UpdateIO &updateIO, + UpdateIO &updateIO, const UpdateOptions &updateOptions, uint32_t featureFlags, const std::shared_ptr &target, diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index ee064690..a30eaeb6 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -9,3 +9,4 @@ add_subdirectory(kv) add_subdirectory(storage) add_subdirectory(mgmtd) add_subdirectory(migration) +add_subdirectory(gdr) diff --git a/tests/common/CMakeLists.txt b/tests/common/CMakeLists.txt index 83a47217..7b3180cb 100644 --- a/tests/common/CMakeLists.txt +++ b/tests/common/CMakeLists.txt @@ -1 +1,6 @@ -target_add_test(test_common common fdb mgmtd-fbs) +target_add_test(test_common common client-lib-common fdb mgmtd-fbs) + +if(NOT HF3FS_ENABLE_GDR) + target_sources(test_common PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../gdr/TestGdrOffCompile.cc) + target_link_libraries(test_common hf3fs_api) +endif() diff --git a/tests/common/net/ib/TestRDMABuf.cc b/tests/common/net/ib/TestRDMABuf.cc index 61bcd958..6227928a 100644 --- a/tests/common/net/ib/TestRDMABuf.cc +++ b/tests/common/net/ib/TestRDMABuf.cc @@ -4,6 +4,7 @@ #include #include #include +#include #include #include #include @@ -81,6 +82,15 @@ TEST_F(TestRDMARemoteBuf, Subrange) { ASSERT_FALSE(buf8.subtract(buf.size() + 1)); } +TEST(TestRDMARemoteBufPure, SubrangeRejectsOverflow) { + RDMARemoteBuf buf(0x1000, 16, {}); + + ASSERT_TRUE(buf); + EXPECT_FALSE(buf.subrange(std::numeric_limits::max(), 1)); + EXPECT_FALSE(buf.subrange(8, std::numeric_limits::max())); + EXPECT_TRUE(buf.subrange(16, 0)); +} + class TestRDMABuf : public test::SetupIB {}; TEST_F(TestRDMABuf, Default) { diff --git a/tests/common/utils/TestGdrUri.cc b/tests/common/utils/TestGdrUri.cc new file mode 100644 index 00000000..94b889ed --- /dev/null +++ b/tests/common/utils/TestGdrUri.cc @@ -0,0 +1,123 @@ +#include +#include +#include +#include +#include + +#include "lib/common/GdrUri.h" + +namespace hf3fs::lib { +namespace { + +std::array makeHandle() { + std::array handle{}; + for (size_t i = 0; i < handle.size(); ++i) { + handle[i] = static_cast(i * 3 + 7); + } + return handle; +} + +std::string ipcHex(const std::string &uri) { + auto position = uri.find("/ipc/"); + return uri.substr(position + 5); +} + +} // namespace + +TEST(TestGdrUri, FormatAndParseV2RoundTripWithOffset) { + auto handle = makeHandle(); + + auto uri = formatGdrUri(2, 16384, 4096, 8192, handle.data(), handle.size()); + ASSERT_FALSE(uri.empty()); + EXPECT_EQ(uri, "gdr://v2/device/2/allocation/16384/offset/4096/size/8192/ipc/" + ipcHex(uri)); + + auto parsed = parseGdrUri(uri); + ASSERT_TRUE(parsed); + EXPECT_EQ(parsed->deviceId, 2); + EXPECT_EQ(parsed->allocationSize, 16384u); + EXPECT_EQ(parsed->offset, 4096u); + EXPECT_EQ(parsed->size, 8192u); + EXPECT_EQ(parsed->ipcHandle, handle); +} + +TEST(TestGdrUri, FormatRejectsInvalidInputAndBounds) { + auto handle = makeHandle(); + + EXPECT_TRUE(formatGdrUri(-1, 4096, 0, 4096, handle.data(), handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 0, 0, 1, handle.data(), handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 4096, 0, 0, handle.data(), handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 4096, 4097, 1, handle.data(), handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 4096, 4096, 1, handle.data(), handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 4096, 0, 4096, nullptr, handle.size()).empty()); + EXPECT_TRUE(formatGdrUri(0, 4096, 0, 4096, handle.data(), handle.size() - 1).empty()); +} + +TEST(TestGdrUri, AcceptsMaximumInBoundsViewWithoutOverflow) { + auto handle = makeHandle(); + constexpr auto max = std::numeric_limits::max(); + + auto uri = formatGdrUri(std::numeric_limits::max(), max, max - 1, 1, handle.data(), handle.size()); + ASSERT_FALSE(uri.empty()); + + auto parsed = parseGdrUri(uri); + ASSERT_TRUE(parsed); + EXPECT_EQ(parsed->deviceId, std::numeric_limits::max()); + EXPECT_EQ(parsed->allocationSize, max); + EXPECT_EQ(parsed->offset, max - 1); + EXPECT_EQ(parsed->size, 1u); +} + +TEST(TestGdrUri, ParseRejectsMalformedFieldsAndTrailingData) { + auto handle = makeHandle(); + auto valid = formatGdrUri(1, 8192, 1024, 4096, handle.data(), handle.size()); + ASSERT_FALSE(valid.empty()); + auto hex = ipcHex(valid); + + EXPECT_FALSE(parseGdrUri("")); + EXPECT_FALSE(parseGdrUri("gdr://v1/device/1/allocation/8192/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device//allocation/8192/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/-1/allocation/8192/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/+1/allocation/8192/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation//offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/-1/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/0/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset//size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/-1/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/0/size//ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/0/size/-1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/0/size/0/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/0/size/1/ipc/")); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/1/allocation/8192/offset/0/size/ 1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri(valid + "/tail")); + + auto badHex = valid; + badHex.back() = 'x'; + EXPECT_FALSE(parseGdrUri(badHex)); + badHex = valid; + badHex[badHex.find("/ipc/") + 5] = 'g'; + EXPECT_FALSE(parseGdrUri(badHex)); + EXPECT_FALSE(parseGdrUri(valid.substr(0, valid.size() - 2))); +} + +TEST(TestGdrUri, ParseRejectsNumericOverflowAndOutOfBoundsViews) { + auto handle = makeHandle(); + auto valid = formatGdrUri(0, 1, 0, 1, handle.data(), handle.size()); + auto hex = ipcHex(valid); + + auto deviceOverflow = std::to_string(static_cast(std::numeric_limits::max()) + 1); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/" + deviceOverflow + "/allocation/1/offset/0/size/1/ipc/" + hex)); + + constexpr auto numericOverflow = "184467440737095516160"; + EXPECT_FALSE( + parseGdrUri(std::string("gdr://v2/device/0/allocation/") + numericOverflow + "/offset/0/size/1/ipc/" + hex)); + EXPECT_FALSE( + parseGdrUri(std::string("gdr://v2/device/0/allocation/1/offset/") + numericOverflow + "/size/1/ipc/" + hex)); + EXPECT_FALSE( + parseGdrUri(std::string("gdr://v2/device/0/allocation/1/offset/0/size/") + numericOverflow + "/ipc/" + hex)); + + EXPECT_FALSE(parseGdrUri("gdr://v2/device/0/allocation/4096/offset/4097/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/0/allocation/4096/offset/4096/size/1/ipc/" + hex)); + EXPECT_FALSE(parseGdrUri("gdr://v2/device/0/allocation/4096/offset/2048/size/2049/ipc/" + hex)); +} + +} // namespace hf3fs::lib diff --git a/tests/fuse/TestIovTableGdr.cc b/tests/fuse/TestIovTableGdr.cc new file mode 100644 index 00000000..9fb5e6a3 --- /dev/null +++ b/tests/fuse/TestIovTableGdr.cc @@ -0,0 +1,544 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "client/storage/StorageClient.h" +#include "client/storage/StorageClientInMem.h" +#include "common/net/ib/SetupIB.h" +#include "fuse/FuseClients.h" +#include "fuse/IoRing.h" +#include "fuse/IovTable.h" +#include "fuse/PioV.h" +#include "tests/FakeMgmtdClient.h" +#include "tests/GtestHelpers.h" + +namespace hf3fs::fuse { + +class IovTableTestHelper { + public: + static Result insert(IovTable &table, std::shared_ptr entry) { + auto iovd = table.entries_->alloc(); + if (!iovd) { + return makeError(ClientAgentCode::kTooManyOpenFiles, "too many test iovs"); + } + + auto published = table.publishEntry(*iovd, std::move(entry)); + if (!published) { + table.entries_->dealloc(*iovd); + RETURN_ERROR(published); + } + return *iovd; + } + + static bool removeExpected(IovTable &table, int iovd, const std::shared_ptr &expected) { + std::unique_lock lock(table.mutex_); + return table.removeEntryLocked(iovd, expected); + } +}; + +namespace { + +meta::UserInfo user(uint32_t uid = 0, uint32_t gid = 0) { return meta::UserInfo{meta::Uid(uid), meta::Gid(gid)}; } + +class TestShm { + public: + explicit TestShm(size_t size) + : name_("/hf3fs-iov-table-" + Uuid::random().toHexString()), + path_(Path("/dev/shm") / name_.substr(1)) { + fd_ = shm_open(name_.c_str(), O_CREAT | O_EXCL | O_RDWR, 0600); + if (fd_ >= 0 && ftruncate(fd_, size) == 0) { + valid_ = true; + } + } + + ~TestShm() { + if (fd_ >= 0) { + close(fd_); + } + shm_unlink(name_.c_str()); + } + + bool valid() const { return valid_; } + const Path &path() const { return path_; } + + private: + std::string name_; + Path path_; + int fd_ = -1; + bool valid_ = false; +}; + +class TestIovTableGdr : public ::testing::Test { + protected: + TestIovTableGdr() + : mgmtd_(tests::FakeMgmtdClient::create()), + storageClient_(ClientId::random(), storageConfig_, *mgmtd_) {} + + auto addIovForParser(IovTable &table, const char *key, const Path &target) { + return table.addIov(key, target, 1234, user(), folly::Executor::KeepAlive<>{}, storageClient_); + } + + void expectParserError(IovTable &table, const char *key, std::string_view message) { + auto result = addIovForParser(table, key, Path("/dev/shm/unused")); + ASSERT_TRUE(result.hasError()); + EXPECT_THAT(std::string(result.error().message()), testing::HasSubstr(std::string(message))); + } + + storage::client::StorageClient::Config storageConfig_; + std::shared_ptr mgmtd_; + storage::client::StorageClientInMem storageClient_; +}; + +class TestIovTableGdrWithIB : public TestIovTableGdr { + public: + static void SetUpTestSuite() { net::test::SetupIB::SetUpTestSuite(); } +}; + +} // namespace + +TEST_F(TestIovTableGdr, ParsesAndValidatesGdrV2Keys) { + IovTable table; + + auto result = addIovForParser(table, "abcdef1234567890abcdef1234567890.gdr.d0", Path("gdr://invalid")); + ASSERT_TRUE(result.hasError()); + EXPECT_THAT(std::string(result.error().message()), testing::HasSubstr("failed to parse GDR target URI")); + + expectParserError(table, "", "invalid shm key"); + expectParserError(table, "abcdef1234567890abcdef1234567890..gdr.d0", "empty attr"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d", "invalid gpu device id"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d-1", "invalid gpu device id"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d1x", "invalid gpu device id"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gpu.d0", "invalid gdr attr"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d0.unknown", "unknown attr"); + expectParserError(table, "abcdef1234567890abcdef1234567890.d0", "gpu device id set for non-gdr key"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr", "gdr key missing gpu device id"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d0.b4096", "gdr key does not support"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d0.r4", "gdr key does not support"); + expectParserError(table, "abcdef1234567890abcdef1234567890.gdr.d0.t10", "gdr key does not support"); + expectParserError(table, "abcdef1234567890abcdef1234567890.r4.phx", "invalid priority"); + expectParserError(table, "abcdef1234567890abcdef1234567890.b0", "invalid block size"); + expectParserError(table, "abcdef1234567890abcdef1234567890.r", "invalid io batch size"); +} + +TEST_F(TestIovTableGdrWithIB, HostEntrySupportsMetadataLookupAndDeadPidRemoval) { + constexpr pid_t kPid = 4242; + constexpr int kIoRingIndex = 7; + const auto size = IoRing::bytesRequired(4); + TestShm testShm(size); + if (!testShm.valid()) { + GTEST_SKIP() << "POSIX shared memory is unavailable"; + } + + auto id = Uuid::random(); + auto key = id.toHexString() + ".r1"; + auto owner = user(17, 23); + IovTable table; + table.init(Path("/mnt/3fs"), 8); + + auto added = table.addIov(key.c_str(), testShm.path(), kPid, owner, folly::Executor::KeepAlive<>{}, storageClient_); + ASSERT_OK(added); + ASSERT_TRUE(added->second); + added->second->iorIndex = kIoRingIndex; + + auto iovd = table.iovDesc(added->first.id); + ASSERT_TRUE(iovd); + auto entry = table.entryAt(*iovd); + ASSERT_TRUE(entry); + EXPECT_FALSE(entry->isGpu()); + EXPECT_EQ(entry->key, key); + EXPECT_EQ(entry->id, id); + EXPECT_EQ(entry->target, testShm.path()); + EXPECT_EQ(entry->user, owner.uid); + EXPECT_EQ(entry->gid, owner.gid); + EXPECT_EQ(entry->pid, kPid); + + auto stat = table.statIov(*iovd, owner); + ASSERT_OK(stat); + EXPECT_EQ(stat->asSymlink().target, testShm.path()); + EXPECT_EQ(stat->acl.uid, owner.uid); + EXPECT_EQ(stat->acl.gid, owner.gid); + EXPECT_EQ(stat->acl.perm, meta::Permission(0400)); + ASSERT_ERROR(table.statIov(*iovd, user(18, 23)), MetaCode::kNoPermission); + + auto [dirEntries, inodes] = table.listIovs(owner); + ASSERT_EQ(dirEntries->size(), 4); + ASSERT_EQ(inodes->size(), 4); + EXPECT_EQ(dirEntries->back().name, key); + ASSERT_TRUE(inodes->back()); + EXPECT_EQ(inodes->back()->asSymlink().target, testShm.path()); + + std::vector> bufs; + std::vector first{{id, 8, 16}, {id, 32, 8}}; + table.lookupBufs(bufs, first, owner.uid); + ASSERT_EQ(bufs.size(), 2); + ASSERT_OK(bufs[0]); + ASSERT_OK(bufs[1]); + auto *firstPtr = ioBufPtr(bufs[0].value()); + EXPECT_EQ(firstPtr, added->second->bufStart + 8); + + std::vector> denied; + table.lookupBufs(denied, first, meta::Uid(18)); + ASSERT_EQ(denied.size(), first.size()); + ASSERT_ERROR(denied[0], MetaCode::kNoPermission); + ASSERT_ERROR(denied[1], MetaCode::kNoPermission); + ASSERT_ERROR(table.lookupBuf(IovLookupRequest{id, 0, 1}, meta::Uid(18)), MetaCode::kNoPermission); + + Uuid missing = Uuid::random(); + std::vector wrapped{{id, size - 4, 8}, {missing, 0, 1}}; + table.lookupBufs(bufs, wrapped, owner.uid); + ASSERT_EQ(bufs.size(), 4); + ASSERT_OK(bufs[0]); + EXPECT_EQ(ioBufPtr(bufs[0].value()), firstPtr); + EXPECT_TRUE(bufs[2].hasError()); + EXPECT_TRUE(bufs[3].hasError()); + + auto ioRingIndexes = table.removeIovsByPid(kPid); + ASSERT_EQ(ioRingIndexes, std::vector({kIoRingIndex})); + EXPECT_FALSE(table.entryAt(*iovd)); + EXPECT_TRUE(table.lookupIov(key.c_str(), owner).hasError()); + EXPECT_TRUE(table.lookupBuf(IovLookupRequest{id, 0, 1}, owner.uid).hasError()); +} + +TEST_F(TestIovTableGdrWithIB, HostRmIovReturnsIoRingBuffer) { + TestShm testShm(IoRing::bytesRequired(2)); + if (!testShm.valid()) { + GTEST_SKIP() << "POSIX shared memory is unavailable"; + } + + auto id = Uuid::random(); + auto key = id.toHexString() + ".r1"; + IovTable table; + table.init(Path("/mnt/3fs"), 2); + auto added = table.addIov(key.c_str(), testShm.path(), 123, user(), folly::Executor::KeepAlive<>{}, storageClient_); + ASSERT_OK(added); + + auto removed = table.rmIov(key.c_str(), user()); + ASSERT_OK(removed); + EXPECT_EQ(*removed, added->second); +} + +TEST_F(TestIovTableGdrWithIB, HostAddRejectsDuplicateKeyAndUuidWithoutCorruptingReverseMaps) { + TestShm testShm(IoRing::bytesRequired(2)); + if (!testShm.valid()) { + GTEST_SKIP() << "POSIX shared memory is unavailable"; + } + + auto id = Uuid::random(); + auto readKey = id.toHexString() + ".r1"; + auto writeKey = id.toHexString() + ".w1"; + auto owner = user(31, 32); + IovTable table; + table.init(Path("/mnt/3fs"), 1); + + auto first = + table.addIov(readKey.c_str(), testShm.path(), 1001, owner, folly::Executor::KeepAlive<>{}, storageClient_); + ASSERT_OK(first); + ASSERT_TRUE(first->second); + ASSERT_ERROR( + table.addIov(readKey.c_str(), testShm.path(), 1002, owner, folly::Executor::KeepAlive<>{}, storageClient_), + MetaCode::kExists); + ASSERT_ERROR( + table.addIov(writeKey.c_str(), testShm.path(), 1003, owner, folly::Executor::KeepAlive<>{}, storageClient_), + MetaCode::kExists); + + ASSERT_OK(table.lookupBuf(IovLookupRequest{id, 0, 1}, owner.uid)); + ASSERT_OK(table.rmIov(readKey.c_str(), owner)); + + auto replacement = + table.addIov(writeKey.c_str(), testShm.path(), 1004, owner, folly::Executor::KeepAlive<>{}, storageClient_); + ASSERT_OK(replacement); + ASSERT_ERROR(table.rmIov(writeKey.c_str(), owner, first->second), MetaCode::kNotFound); + ASSERT_OK(table.lookupIov(writeKey.c_str(), owner)); + ASSERT_OK(table.lookupBuf(IovLookupRequest{id, 0, 1}, owner.uid)); + + ASSERT_OK(table.rmIov(writeKey.c_str(), owner)); + EXPECT_TRUE(table.lookupBuf(IovLookupRequest{id, 0, 1}, owner.uid).hasError()); +} + +TEST_F(TestIovTableGdr, StalePidSnapshotCannotRemoveReusedDescriptor) { + IovTable table; + table.init(Path("/mnt/3fs"), 1); + auto owner = user(41, 42); + + auto oldId = Uuid::random(); + auto oldEntry = std::make_shared( + IovEntry{oldId.toHexString(), oldId, Path("/old"), owner.uid, owner.gid, 2001, IovBuffer{}}); + auto oldIovd = IovTableTestHelper::insert(table, oldEntry); + ASSERT_OK(oldIovd); + ASSERT_OK(table.rmIov(oldEntry->key.c_str(), owner)); + + auto newId = Uuid::random(); + auto newEntry = std::make_shared( + IovEntry{newId.toHexString(), newId, Path("/new"), owner.uid, owner.gid, 2002, IovBuffer{}}); + auto newIovd = IovTableTestHelper::insert(table, newEntry); + ASSERT_OK(newIovd); + ASSERT_EQ(*newIovd, *oldIovd); + + EXPECT_FALSE(IovTableTestHelper::removeExpected(table, *oldIovd, oldEntry)); + EXPECT_EQ(table.entryAt(*newIovd), newEntry); + EXPECT_TRUE(table.removeIovsByPid(2001).empty()); + EXPECT_EQ(table.entryAt(*newIovd), newEntry); +} + +TEST_F(TestIovTableGdr, ClosedReadHoleZerosOnlyMissingRanges) { + std::array data{}; + std::vector reads; + auto addRead = [&](uint32_t chunk, size_t dataOffset, Result result) { + auto chunkId = storage::ChunkId(meta::ChunkId(meta::InodeId{77}, 0, chunk).pack()); + auto read = + storageClient_ + .createReadIO({}, chunkId, 0, 4, data.data() + dataOffset, nullptr, reinterpret_cast(size_t{0})); + read.result.lengthInfo = std::move(result); + reads.push_back(std::move(read)); + }; + addRead(0, 0, 2); + addRead(1, 4, makeError(StorageClientCode::kChunkNotFound)); + addRead(2, 8, 4); + + std::vector result{0}; + std::vector> zeroed; + lib::agent::detail::finishReadResults( + result, + reads, + true, + [&](const storage::client::ReadIO &, size_t offset, size_t length) -> Result { + zeroed.emplace_back(offset, length); + return Void{}; + }); + + EXPECT_EQ(result, std::vector({12})); + EXPECT_EQ(zeroed, (std::vector>{{2, 2}, {0, 4}})); +} + +TEST_F(TestIovTableGdr, ReadHoleDoesNotLeakWhenNextIovStartsWithMissingChunk) { + std::array firstBuffer; + std::array secondBuffer; + firstBuffer.fill(0xAA); + secondBuffer.fill(0xBB); + std::vector reads; + auto addRead = [&](uint32_t chunk, size_t iovIdx, uint8_t *data, Result result) { + auto chunkId = storage::ChunkId(meta::ChunkId(meta::InodeId{80}, 0, chunk).pack()); + auto read = storageClient_.createReadIO({}, chunkId, 0, 4, data, nullptr, reinterpret_cast(iovIdx)); + read.result.lengthInfo = std::move(result); + reads.push_back(std::move(read)); + }; + addRead(0, 0, firstBuffer.data(), 2); + addRead(1, 1, secondBuffer.data(), makeError(StorageClientCode::kChunkNotFound)); + addRead(2, 1, secondBuffer.data() + 4, 4); + + std::vector result{0, 0}; + lib::agent::detail::finishReadResults( + result, + reads, + true, + [](const storage::client::ReadIO &io, size_t offset, size_t length) -> Result { + std::memset(io.data + offset, 0, length); + return Void{}; + }); + + EXPECT_EQ(result, (std::vector{2, 8})); + EXPECT_EQ(firstBuffer, (std::array{0xAA, 0xAA, 0xAA, 0xAA})); + EXPECT_EQ(secondBuffer, (std::array{0, 0, 0, 0, 0xBB, 0xBB, 0xBB, 0xBB})); +} + +TEST_F(TestIovTableGdr, ReadHoleDoesNotLeakWhenNextIovStartsWithShortRead) { + std::array firstBuffer; + std::array secondBuffer; + firstBuffer.fill(0xAA); + secondBuffer.fill(0xBB); + std::vector reads; + auto addRead = [&](uint32_t chunk, size_t iovIdx, uint8_t *data, uint32_t result) { + auto chunkId = storage::ChunkId(meta::ChunkId(meta::InodeId{81}, 0, chunk).pack()); + auto read = storageClient_.createReadIO({}, chunkId, 0, 4, data, nullptr, reinterpret_cast(iovIdx)); + read.result.lengthInfo = result; + reads.push_back(std::move(read)); + }; + addRead(0, 0, firstBuffer.data(), 2); + addRead(1, 1, secondBuffer.data(), 2); + addRead(2, 1, secondBuffer.data() + 4, 4); + + std::vector result{0, 0}; + lib::agent::detail::finishReadResults( + result, + reads, + true, + [](const storage::client::ReadIO &io, size_t offset, size_t length) -> Result { + std::memset(io.data + offset, 0, length); + return Void{}; + }); + + EXPECT_EQ(result, (std::vector{2, 8})); + EXPECT_EQ(firstBuffer, (std::array{0xAA, 0xAA, 0xAA, 0xAA})); + EXPECT_EQ(secondBuffer, (std::array{0xBB, 0xBB, 0, 0, 0xBB, 0xBB, 0xBB, 0xBB})); +} + +TEST_F(TestIovTableGdr, ReadHolePropagatesZeroFailureAndHonorsForbidMode) { + std::array data{}; + std::vector reads; + for (uint32_t chunk = 0; chunk < 2; ++chunk) { + auto chunkId = storage::ChunkId(meta::ChunkId(meta::InodeId{78}, 0, chunk).pack()); + auto read = + storageClient_ + .createReadIO({}, chunkId, 0, 4, data.data() + chunk * 4, nullptr, reinterpret_cast(size_t{0})); + read.result.lengthInfo = chunk == 0 ? 2 : 4; + reads.push_back(std::move(read)); + } + + std::vector result{0}; + size_t zeroCalls = 0; + lib::agent::detail::finishReadResults(result, + reads, + true, + [&](const storage::client::ReadIO &, size_t, size_t) -> Result { + ++zeroCalls; + return makeError(StorageClientCode::kRemoteIOError, + "injected GPU memset failure"); + }); + EXPECT_EQ(zeroCalls, 1u); + EXPECT_EQ(result[0], -static_cast(StorageClientCode::kRemoteIOError)); + + result = {0}; + zeroCalls = 0; + lib::agent::detail::finishReadResults(result, + reads, + false, + [&](const storage::client::ReadIO &, size_t, size_t) -> Result { + ++zeroCalls; + return Void{}; + }); + EXPECT_EQ(zeroCalls, 0u); + EXPECT_EQ(result[0], -static_cast(ClientAgentCode::kHoleInIoOutcome)); +} + +TEST_F(TestIovTableGdr, FullReadHoleRemainsEofAndIsNotZeroFilled) { + std::array data{}; + std::vector reads; + for (uint32_t chunk = 0; chunk < 2; ++chunk) { + auto chunkId = storage::ChunkId(meta::ChunkId(meta::InodeId{79}, 0, chunk).pack()); + auto read = + storageClient_ + .createReadIO({}, chunkId, 0, 4, data.data() + chunk * 4, nullptr, reinterpret_cast(size_t{0})); + if (chunk == 0) { + read.result.lengthInfo = 0; + } else { + read.result.lengthInfo = makeError(StorageClientCode::kChunkNotFound); + } + reads.push_back(std::move(read)); + } + + std::vector result{0}; + size_t zeroCalls = 0; + lib::agent::detail::finishReadResults(result, + reads, + true, + [&](const storage::client::ReadIO &, size_t, size_t) -> Result { + ++zeroCalls; + return Void{}; + }); + + EXPECT_EQ(result[0], 0); + EXPECT_EQ(zeroCalls, 0u); +} + +#ifdef HF3FS_ENABLE_GDR + +TEST_F(TestIovTableGdr, GpuPublicationRejectsDuplicateKeyAndUuidAndSupportsReverseRemoval) { + IovTable table; + table.init(Path("/mnt/3fs"), 3); + + std::shared_ptr noHardwareBuffer; + auto owner = user(51, 52); + auto id = Uuid::random(); + auto key = id.toHexString() + ".gdr.d0"; + Path target("gdr://v2/device/0/allocation/64/offset/0/size/64/ipc/" + std::string(128, '0')); + auto first = + std::make_shared(IovEntry{key, id, target, owner.uid, owner.gid, 3001, IovBuffer{noHardwareBuffer}}); + ASSERT_OK(IovTableTestHelper::insert(table, first)); + + auto differentId = Uuid::random(); + auto duplicateKey = std::make_shared( + IovEntry{key, differentId, target, owner.uid, owner.gid, 3002, IovBuffer{noHardwareBuffer}}); + ASSERT_ERROR(IovTableTestHelper::insert(table, duplicateKey), MetaCode::kExists); + + auto replacementKey = id.toHexString() + ".gdr.d1"; + auto duplicateId = std::make_shared( + IovEntry{replacementKey, id, target, owner.uid, owner.gid, 3003, IovBuffer{noHardwareBuffer}}); + ASSERT_ERROR(IovTableTestHelper::insert(table, duplicateId), MetaCode::kExists); + ASSERT_OK(table.lookupIov(key.c_str(), owner)); + + ASSERT_OK(table.rmIov(key.c_str(), owner)); + ASSERT_OK(IovTableTestHelper::insert(table, duplicateId)); + ASSERT_OK(table.lookupIov(replacementKey.c_str(), owner)); + ASSERT_OK(table.rmIov(replacementKey.c_str(), owner)); + EXPECT_TRUE(table.lookupIov(replacementKey.c_str(), owner).hasError()); +} + +TEST_F(TestIovTableGdr, TaggedGpuEntryUsesNonNullSlotAndUnifiedMetadataPaths) { + IovTable table; + table.init(Path("/mnt/3fs"), 4); + + auto id = Uuid::random(); + auto key = id.toHexString() + ".gdr.d0"; + Path target("gdr://v2/device/0/allocation/4096/offset/0/size/4096/ipc/" + std::string(128, '0')); + std::shared_ptr noHardwareBuffer; + auto entry = std::make_shared( + IovEntry{key, id, target, meta::Uid(7), meta::Gid(9), 5150, IovBuffer{noHardwareBuffer}}); + + auto iovd = IovTableTestHelper::insert(table, entry); + ASSERT_OK(iovd); + auto stored = table.entryAt(*iovd); + ASSERT_TRUE(stored); + EXPECT_TRUE(stored->isGpu()); + + auto stat = table.statIov(*iovd, user(7, 9)); + ASSERT_OK(stat); + EXPECT_EQ(stat->asSymlink().target, target); + EXPECT_EQ(stat->acl.uid, meta::Uid(7)); + EXPECT_EQ(stat->acl.gid, meta::Gid(9)); + + auto [dirEntries, inodes] = table.listIovs(user(7, 9)); + ASSERT_EQ(dirEntries->size(), 4); + EXPECT_EQ(dirEntries->back().name, key); + ASSERT_TRUE(inodes->back()); + EXPECT_EQ(inodes->back()->asSymlink().target, target); + ASSERT_ERROR(table.lookupBuf(IovLookupRequest{id, 0, 1}, meta::Uid(8)), MetaCode::kNoPermission); + + auto removed = table.rmIov(key.c_str(), user(7, 9)); + ASSERT_OK(removed); + EXPECT_FALSE(*removed); + EXPECT_FALSE(table.entryAt(*iovd)); + + auto deadId = Uuid::random(); + auto deadKey = deadId.toHexString() + ".gdr.d0"; + auto deadEntry = std::make_shared( + IovEntry{deadKey, deadId, target, meta::Uid(7), meta::Gid(9), 5150, IovBuffer{noHardwareBuffer}}); + auto deadIovd = IovTableTestHelper::insert(table, deadEntry); + ASSERT_OK(deadIovd); + EXPECT_TRUE(table.removeIovsByPid(5150).empty()); + EXPECT_FALSE(table.entryAt(*deadIovd)); + + auto clearId = Uuid::random(); + auto clearKey = clearId.toHexString() + ".gdr.d0"; + auto clearEntry = std::make_shared( + IovEntry{clearKey, clearId, target, meta::Uid(7), meta::Gid(9), 6160, IovBuffer{noHardwareBuffer}}); + auto clearIovd = IovTableTestHelper::insert(table, clearEntry); + ASSERT_OK(clearIovd); + table.clearGpuIovs(); + EXPECT_FALSE(table.entryAt(*clearIovd)); +} + +#endif + +} // namespace hf3fs::fuse diff --git a/tests/gdr/CMakeLists.txt b/tests/gdr/CMakeLists.txt new file mode 100644 index 00000000..8714352a --- /dev/null +++ b/tests/gdr/CMakeLists.txt @@ -0,0 +1,24 @@ +# GDR (GPU Direct RDMA) Tests +# +# All CUDA/GDR-dependent test sources are consolidated here for unified build control. +# Hardware-dependent tests use GTEST_SKIP when runtime support is unavailable. +# +# target_add_test() picks up local *.cc files. The FUSE parser/metadata test +# stays in tests/fuse and is added explicitly below. + +if(HF3FS_ENABLE_GDR) + # libfuse3 search paths (mirrored from src/fuse/CMakeLists.txt) + if(CMAKE_SYSTEM_PROCESSOR STREQUAL "x86_64") + link_directories(/usr/local/lib/x86_64-linux-gnu/ /usr/lib64 /usr/local/lib64) + elseif(CMAKE_SYSTEM_PROCESSOR STREQUAL "aarch64") + link_directories(/usr/local/lib/aarch64-linux-gnu/ /usr/lib64 /usr/local/lib64) + endif() + + target_add_test(test_gdr hf3fs_api common hf3fs_fuse fdb mgmtd-fbs) + target_sources(test_gdr PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../fuse/TestIovTableGdr.cc) + target_add_gdr_support(test_gdr SCOPE PRIVATE) + target_include_directories(test_gdr PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/.. # For tests/GtestHelpers.h + ${CMAKE_SOURCE_DIR}/src # For source headers + ) +endif() diff --git a/tests/gdr/TestCudaRDMABuf.cc b/tests/gdr/TestCudaRDMABuf.cc new file mode 100644 index 00000000..1ba8815e --- /dev/null +++ b/tests/gdr/TestCudaRDMABuf.cc @@ -0,0 +1,288 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "common/cuda/CudaMemory.h" +#include "common/net/ib/IBDevice.h" +#include "common/net/ib/RDMABuf.h" +#include "lib/common/CudaIpcMemory.h" +#include "lib/common/GpuShm.h" +#include "tests/GtestHelpers.h" + +namespace hf3fs::net { +namespace { + +std::optional firstCudaDevice() { + auto count = cuda::deviceCount(); + if (!count) { + return std::nullopt; + } + for (int deviceId = 0; deviceId < *count; ++deviceId) { + auto available = cuda::supportsIpc(deviceId); + if (available && *available) { + return deviceId; + } + } + return std::nullopt; +} + +uint8_t *deviceAddress(void *base, size_t offset) { + return reinterpret_cast(reinterpret_cast(base) + offset); +} + +bool ensureIbStarted() { + if (IBManager::initialized()) { + return !IBDevice::all().empty(); + } + static IBConfig config; + auto result = IBManager::start(config); + return result && !IBDevice::all().empty(); +} + +std::string currentExecutable() { + std::array path{}; + auto length = readlink("/proc/self/exe", path.data(), path.size() - 1); + if (length <= 0) { + return {}; + } + path[length] = '\0'; + return path.data(); +} + +std::string toHex(const lib::CudaIpcHandle &handle) { + constexpr char digits[] = "0123456789abcdef"; + std::string result; + result.resize(handle.size() * 2); + for (size_t i = 0; i < handle.size(); ++i) { + result[i * 2] = digits[handle[i] >> 4]; + result[i * 2 + 1] = digits[handle[i] & 0x0f]; + } + return result; +} + +bool fromHex(std::string_view text, lib::CudaIpcHandle &handle) { + if (text.size() != handle.size() * 2) { + return false; + } + auto nibble = [](char value) -> int { + if (value >= '0' && value <= '9') return value - '0'; + if (value >= 'a' && value <= 'f') return value - 'a' + 10; + if (value >= 'A' && value <= 'F') return value - 'A' + 10; + return -1; + }; + for (size_t i = 0; i < handle.size(); ++i) { + auto high = nibble(text[i * 2]); + auto low = nibble(text[i * 2 + 1]); + if (high < 0 || low < 0) { + return false; + } + handle[i] = static_cast((high << 4) | low); + } + return true; +} + +int runIpcMrImporter(const lib::CudaIpcExport &exported, size_t viewSize, int deviceId) { + auto executable = currentExecutable(); + if (executable.empty()) { + return -1; + } + + auto handle = toHex(exported.ipcHandle); + auto allocationSize = std::to_string(exported.allocationSize); + auto offset = std::to_string(exported.offset); + auto size = std::to_string(viewSize); + auto device = std::to_string(deviceId); + setenv("HF3FS_GDR_PROBE_HANDLE", handle.c_str(), 1); + setenv("HF3FS_GDR_PROBE_ALLOCATION_SIZE", allocationSize.c_str(), 1); + setenv("HF3FS_GDR_PROBE_OFFSET", offset.c_str(), 1); + setenv("HF3FS_GDR_PROBE_SIZE", size.c_str(), 1); + setenv("HF3FS_GDR_PROBE_DEVICE", device.c_str(), 1); + SCOPE_EXIT { + unsetenv("HF3FS_GDR_PROBE_HANDLE"); + unsetenv("HF3FS_GDR_PROBE_ALLOCATION_SIZE"); + unsetenv("HF3FS_GDR_PROBE_OFFSET"); + unsetenv("HF3FS_GDR_PROBE_SIZE"); + unsetenv("HF3FS_GDR_PROBE_DEVICE"); + }; + + auto child = fork(); + if (child == 0) { + execl(executable.c_str(), + executable.c_str(), + "--gtest_filter=TestCudaRDMABuf.IpcMrImporterProcess", + "--gtest_color=no", + static_cast(nullptr)); + _exit(127); + } + if (child < 0) { + return -1; + } + + int status = 0; + while (waitpid(child, &status, 0) < 0) { + if (errno != EINTR) { + return -1; + } + } + return WIFEXITED(status) ? WEXITSTATUS(status) : -1; +} + +} // namespace + +TEST(TestCudaMemory, RejectsInvalidDeviceViews) { + ASSERT_ERROR(cuda::inspectDeviceMemory(nullptr, 4096, 0), StatusCode::kInvalidArg); + ASSERT_ERROR(cuda::prepareGdrMemory(reinterpret_cast(0x1000), 0, 0), StatusCode::kInvalidArg); +} + +TEST(TestCudaMemory, PreparesWholeAllocationForGdr) { + auto deviceId = firstCudaDevice(); + if (!deviceId) { + GTEST_SKIP() << "No CUDA device available"; + } + + auto guard = cuda::ScopedDevice::create(*deviceId); + ASSERT_OK(guard); + void *allocation = nullptr; + ASSERT_EQ(cudaMalloc(&allocation, 4096), cudaSuccess); + SCOPE_EXIT { EXPECT_EQ(cudaFree(allocation), cudaSuccess); }; + + auto *viewPtr = deviceAddress(allocation, 512); + auto view = cuda::prepareGdrMemory(viewPtr, 1024, *deviceId); + ASSERT_OK(view); + EXPECT_EQ(view->allocationBase, allocation); + EXPECT_GE(view->allocationSize, 4096u); + EXPECT_EQ(view->offset, 512u); + EXPECT_EQ(view->deviceId, *deviceId); + + unsigned int syncMemops = 0; + ASSERT_EQ( + cuPointerGetAttribute(&syncMemops, CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, reinterpret_cast(viewPtr)), + CUDA_SUCCESS); + EXPECT_EQ(syncMemops, 1u); +} + +TEST(TestCudaRDMABuf, RegistersEveryActiveHcaAndSharesOwnerAcrossSubranges) { + auto deviceId = firstCudaDevice(); + if (!deviceId) { + GTEST_SKIP() << "No CUDA device available"; + } + if (!ensureIbStarted()) { + GTEST_SKIP() << "No active IB device available"; + } + + auto guard = cuda::ScopedDevice::create(*deviceId); + ASSERT_OK(guard); + void *allocation = nullptr; + ASSERT_EQ(cudaMalloc(&allocation, 4096), cudaSuccess); + SCOPE_EXIT { EXPECT_EQ(cudaFree(allocation), cudaSuccess); }; + + std::atomic ownerReleased = false; + auto owner = std::shared_ptr(reinterpret_cast(1), [&](void *) { ownerReleased.store(true); }); + auto rdmaBuf = RDMABuf::createFromCudaBuffer(static_cast(allocation), 4096, *deviceId, owner); + ASSERT_OK(rdmaBuf); + owner.reset(); + + EXPECT_TRUE(rdmaBuf->isDeviceMemory()); + EXPECT_EQ(rdmaBuf->cudaDeviceId(), *deviceId); + auto remote = rdmaBuf->toRemoteBuf(); + ASSERT_TRUE(remote); + for (const auto &device : IBDevice::all()) { + EXPECT_NE(rdmaBuf->getMR(device->id()), nullptr); + EXPECT_TRUE(remote.getRkey(device->id()).has_value()); + } + + auto subrange = rdmaBuf->subrange(128, 512); + ASSERT_TRUE(subrange); + EXPECT_TRUE(subrange.isDeviceMemory()); + EXPECT_EQ(subrange.cudaDeviceId(), *deviceId); + *rdmaBuf = RDMABuf(); + EXPECT_FALSE(ownerReleased.load()); + subrange = RDMABuf(); + EXPECT_TRUE(ownerReleased.load()); +} + +TEST(TestCudaRDMABuf, IpcMrImporterProcess) { + auto handleText = std::getenv("HF3FS_GDR_PROBE_HANDLE"); + if (!handleText) { + GTEST_SKIP() << "Subprocess-only CUDA IPC MR probe"; + } + ASSERT_TRUE(ensureIbStarted()); + + auto allocationSizeText = std::getenv("HF3FS_GDR_PROBE_ALLOCATION_SIZE"); + auto offsetText = std::getenv("HF3FS_GDR_PROBE_OFFSET"); + auto sizeText = std::getenv("HF3FS_GDR_PROBE_SIZE"); + auto deviceText = std::getenv("HF3FS_GDR_PROBE_DEVICE"); + ASSERT_NE(allocationSizeText, nullptr); + ASSERT_NE(offsetText, nullptr); + ASSERT_NE(sizeText, nullptr); + ASSERT_NE(deviceText, nullptr); + + lib::CudaIpcHandle handle{}; + ASSERT_TRUE(fromHex(handleText, handle)); + auto allocationSize = std::strtoull(allocationSizeText, nullptr, 10); + auto offset = std::strtoull(offsetText, nullptr, 10); + auto size = std::strtoull(sizeText, nullptr, 10); + auto deviceId = std::atoi(deviceText); + + auto guard = cuda::ScopedDevice::create(deviceId); + ASSERT_OK(guard); + auto gpuShm = lib::GpuShmBuf::create(handle, allocationSize, offset, size, deviceId); + ASSERT_OK(gpuShm); + auto *ioBuffer = (*gpuShm)->ioBuffer(); + ASSERT_NE(ioBuffer, nullptr); + EXPECT_TRUE(ioBuffer->isDeviceMemory()); + + unsigned int syncMemops = 0; + ASSERT_EQ(cuPointerGetAttribute(&syncMemops, + CU_POINTER_ATTRIBUTE_SYNC_MEMOPS, + reinterpret_cast(ioBuffer->data())), + CUDA_SUCCESS); + EXPECT_EQ(syncMemops, 1u); + + auto remote = ioBuffer->subrangeRemote(0, ioBuffer->size()); + ASSERT_TRUE(remote); + for (const auto &device : IBDevice::all()) { + EXPECT_TRUE(remote.getRkey(device->id()).has_value()); + } +} + +TEST(TestCudaRDMABuf, IpcImportedPointerRegistersWithEveryHca) { + auto deviceId = firstCudaDevice(); + if (!deviceId) { + GTEST_SKIP() << "No CUDA device available"; + } + if (!ensureIbStarted()) { + GTEST_SKIP() << "No active IB device available"; + } + if (currentExecutable().empty()) { + GTEST_SKIP() << "/proc/self/exe is unavailable"; + } + + auto guard = cuda::ScopedDevice::create(*deviceId); + ASSERT_OK(guard); + void *allocation = nullptr; + ASSERT_EQ(cudaMalloc(&allocation, 4096), cudaSuccess); + SCOPE_EXIT { EXPECT_EQ(cudaFree(allocation), cudaSuccess); }; + + auto *view = deviceAddress(allocation, 256); + auto exported = lib::exportCudaIpcMemory(view, 1024, *deviceId); + ASSERT_OK(exported); + EXPECT_EQ(runIpcMrImporter(*exported, 1024, *deviceId), 0); +} + +} // namespace hf3fs::net diff --git a/tests/gdr/TestGdrOffCompile.cc b/tests/gdr/TestGdrOffCompile.cc new file mode 100644 index 00000000..e6daef4a --- /dev/null +++ b/tests/gdr/TestGdrOffCompile.cc @@ -0,0 +1,60 @@ +#include +#include +#include +#include +#include +#include + +#include "common/cuda/CudaMemory.h" +#include "common/net/ib/RDMABuf.h" +#include "fuse/IovTable.h" +#include "fuse/IovTypes.h" +#include "lib/api/hf3fs_usrbio.h" +#include "lib/common/CudaIpcMemory.h" +#include "tests/GtestHelpers.h" + +#ifndef HF3FS_ENABLE_GDR + +namespace hf3fs { +namespace { + +using CreateDeviceFn = int (*)(hf3fs_iov *, const char *, size_t, size_t, int); + +static_assert(std::is_same_v); +static_assert(std::is_same_v); +static_assert(std::is_same_v>>); + +TEST(GdrOffCompileGuard, CudaIpcStubsAndHostIovTypesRemainUsable) { + EXPECT_FALSE(hf3fs_gdr_available()); + EXPECT_EQ(hf3fs_gdr_device_count(), 0); + + hf3fs_iov iov{}; + EXPECT_EQ(hf3fs_iovcreate_device(&iov, "/unused", 4096, 1, 0), -EINVAL); + + auto count = cuda::deviceCount(); + ASSERT_OK(count); + EXPECT_EQ(*count, 0); + + auto available = cuda::supportsIpc(0); + ASSERT_OK(available); + EXPECT_FALSE(*available); + + auto exported = lib::exportCudaIpcMemory(reinterpret_cast(uintptr_t{0x1000}), 4096, 0); + ASSERT_ERROR(exported, StatusCode::kNotImplemented); + + auto inspected = cuda::inspectDeviceMemory(reinterpret_cast(uintptr_t{0x1000}), 4096, 0); + ASSERT_ERROR(inspected, StatusCode::kNotImplemented); + + auto registered = net::RDMABuf::createFromCudaBuffer(reinterpret_cast(uintptr_t{0x1000}), 4096, 0); + ASSERT_ERROR(registered, StatusCode::kNotImplemented); + + fuse::IovEntry entry; + EXPECT_FALSE(entry.isGpu()); + + EXPECT_EQ(hf3fs_iovcreate_device(&iov, "/nonexistent", 4096, 0, 0), -ENOTSUP); +} + +} // namespace +} // namespace hf3fs + +#endif diff --git a/tests/gdr/TestUsrbIoGdr.cc b/tests/gdr/TestUsrbIoGdr.cc new file mode 100644 index 00000000..cb9137ca --- /dev/null +++ b/tests/gdr/TestUsrbIoGdr.cc @@ -0,0 +1,726 @@ +/** + * Scenario tests for Layer 3+4: UsrbIoGdr + UsrbIo (unified C API) + * + * Tests GPU IOV create/open/wrap/destroy, query functions, and sync. + * Covers core dispatch paths and unsupported-device behavior. + * + * Covers: REQ-L3-001, REQ-L3-003, REQ-L3-005 + * REQ-L4-001 through REQ-L4-006 + * INV-GDR-001, INV-GDR-002 + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#ifdef HF3FS_ENABLE_GDR +#include +#endif + +#include "common/utils/Uuid.h" +#include "lib/api/UsrbIoGdrInternal.h" +#include "lib/api/hf3fs_usrbio.h" +#include "lib/common/CudaIpcMemory.h" +#include "lib/common/GdrUri.h" +#include "tests/GtestHelpers.h" + +namespace { + +static bool hasGpu() { return hf3fs_gdr_available(); } + +uint8_t *deviceAddress(void *base, size_t offset) { + return reinterpret_cast(reinterpret_cast(base) + offset); +} + +// Temp directory for symlink testing +class TmpDir { + public: + TmpDir() { + const char *base = getenv("TMPDIR"); + path_ = std::string(base ? base : std::filesystem::temp_directory_path().c_str()) + "/gdr_test_XXXXXX"; + char *result = mkdtemp(path_.data()); + if (result) { + valid_ = true; + } + } + + ~TmpDir() { + if (valid_) { + std::filesystem::remove_all(path_); + } + } + + const char *path() const { return path_.c_str(); } + bool valid() const { return valid_; } + + private: + std::string path_; + bool valid_ = false; +}; + +// Helper to build a well-formed GDR URI +std::string buildGdrUri(int deviceId, size_t allocationSize, size_t offset, size_t size, const uint8_t ipcHandle[64]) { + char hex[129]; + for (int i = 0; i < 64; i++) { + snprintf(hex + i * 2, 3, "%02x", ipcHandle[i]); + } + hex[128] = '\0'; + return std::string("gdr://v2/device/") + std::to_string(deviceId) + "/allocation/" + std::to_string(allocationSize) + + "/offset/" + std::to_string(offset) + "/size/" + std::to_string(size) + "/ipc/" + hex; +} + +hf3fs::Uuid iovUuid(const struct hf3fs_iov &iov) { + hf3fs::Uuid id; + std::memcpy(id.data, iov.id, sizeof(id.data)); + return id; +} + +std::string publicationPath(const struct hf3fs_iov &iov, int deviceId) { + return std::string(iov.mount_point) + "/3fs-virt/iovs/" + iovUuid(iov).toHexString() + ".gdr.d" + + std::to_string(deviceId); +} + +bool isSymlink(const std::string &path) { + struct stat statbuf{}; + return lstat(path.c_str(), &statbuf) == 0 && S_ISLNK(statbuf.st_mode); +} + +std::string currentExecutable() { + std::array path{}; + auto length = readlink("/proc/self/exe", path.data(), path.size() - 1); + if (length <= 0) { + return {}; + } + path[length] = '\0'; + return path.data(); +} + +bool runImporterProcess(const struct hf3fs_iov &publisher, int deviceId, uint8_t expectedByte) { + auto executable = currentExecutable(); + if (executable.empty()) { + return false; + } + + auto size = std::to_string(publisher.size); + auto device = std::to_string(deviceId); + auto expected = std::to_string(expectedByte); + auto id = iovUuid(publisher).toHexString(); + setenv("HF3FS_GDR_IMPORT_MOUNT", publisher.mount_point, 1); + setenv("HF3FS_GDR_IMPORT_ID", id.c_str(), 1); + setenv("HF3FS_GDR_IMPORT_SIZE", size.c_str(), 1); + setenv("HF3FS_GDR_IMPORT_DEVICE", device.c_str(), 1); + setenv("HF3FS_GDR_IMPORT_EXPECTED_BYTE", expected.c_str(), 1); + SCOPE_EXIT { + unsetenv("HF3FS_GDR_IMPORT_MOUNT"); + unsetenv("HF3FS_GDR_IMPORT_ID"); + unsetenv("HF3FS_GDR_IMPORT_SIZE"); + unsetenv("HF3FS_GDR_IMPORT_DEVICE"); + unsetenv("HF3FS_GDR_IMPORT_EXPECTED_BYTE"); + }; + + auto child = fork(); + if (child == 0) { + execl(executable.c_str(), + executable.c_str(), + "--gtest_filter=TestUsrbIoGdrFixture.ImporterProcessKeepsPublisherPublication", + "--gtest_color=no", + static_cast(nullptr)); + _exit(127); + } + if (child < 0) { + return false; + } + + int status = 0; + while (waitpid(child, &status, 0) < 0) { + if (errno != EINTR) { + return false; + } + } + return WIFEXITED(status) && WEXITSTATUS(status) == 0; +} + +class TestUsrbIoGdrFixture : public ::testing::Test {}; + +} // namespace + +// ========================================================================== +// REQ-L4-001: iovcreate host path; iovcreate_device dispatches to GPU +// ========================================================================== + +// @tests SCN-L4-001-01b +TEST_F(TestUsrbIoGdrFixture, SCN_L4_001_01b_DeviceApiDispatch) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + + // WHEN: iovcreate_device with device 0 + int rc = hf3fs_iovcreate_device(&iov, "/nonexistent/mount", 4096, 0, 0); + + // THEN: Should fail (no mount) but exercise dispatch + EXPECT_NE(rc, 0); +} + +// @tests SCN-L4-001-02 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_001_02_DeviceCreateRejectsUnavailableGDR) { + if (hasGpu()) { + GTEST_SKIP() << "Test for machines without GPU"; + } + + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + + // WHEN: iovcreate_device, GDR unavailable + int rc = hf3fs_iovcreate_device(&iov, "/nonexistent/mount", 4096, 0, 0); + + // THEN: Device creation never silently substitutes host memory. + EXPECT_EQ(rc, -ENOTSUP); +} + +TEST_F(TestUsrbIoGdrFixture, DeviceApisRejectGpuBlockSize) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {}; + uint8_t buffer[4096] = {}; + + EXPECT_EQ(hf3fs_iovcreate_device(&iov, "/nonexistent/mount", 4096, 4096, 0), -EINVAL); + EXPECT_EQ(hf3fs_iovopen_device(&iov, id, "/nonexistent/mount", 4096, 4096, 0), -EINVAL); + EXPECT_EQ(hf3fs_iovwrap_device(&iov, buffer, id, "/nonexistent/mount", 4096, 4096, 0), -EINVAL); +} + +TEST_F(TestUsrbIoGdrFixture, PublicationRemovalRespectsOwnershipWithoutCuda) { + TmpDir tmpDir; + ASSERT_TRUE(tmpDir.valid()); + std::filesystem::create_directories(std::string(tmpDir.path()) + "/3fs-virt/iovs"); + + struct hf3fs_iov iov{}; + auto id = hf3fs::Uuid::random(); + std::memcpy(iov.id, id.data, sizeof(iov.id)); + std::strcpy(iov.mount_point, tmpDir.path()); + auto link = publicationPath(iov, 3); + ASSERT_EQ(symlink("gdr://test-publication", link.c_str()), 0); + ASSERT_TRUE(isSymlink(link)); + + EXPECT_EQ(hf3fs_iovunlink_gpu_publication_internal(&iov, 3, false), 0); + EXPECT_TRUE(isSymlink(link)); + + EXPECT_EQ(hf3fs_iovunlink_gpu_publication_internal(&iov, 3, true), 0); + EXPECT_FALSE(isSymlink(link)); + EXPECT_EQ(hf3fs_iovunlink_gpu_publication_internal(&iov, 3, true), -ENOENT); +} + +TEST_F(TestUsrbIoGdrFixture, CudaIpcCapabilityAndSuballocationExportDoNotRequireIbRegistration) { + if (!hasGpu()) { + GTEST_SKIP() << "No CUDA IPC-capable GPU available"; + } + + int deviceId = -1; + ASSERT_EQ(cudaGetDevice(&deviceId), cudaSuccess); + void *allocation = nullptr; + ASSERT_EQ(cudaMalloc(&allocation, 4096), cudaSuccess); + SCOPE_EXIT { EXPECT_EQ(cudaFree(allocation), cudaSuccess); }; + + auto *view = deviceAddress(allocation, 512); + auto exported = hf3fs::lib::exportCudaIpcMemory(view, 1024, deviceId); + ASSERT_OK(exported); + EXPECT_EQ(exported->allocationBase, allocation); + EXPECT_GE(exported->allocationSize, 1536u); + EXPECT_EQ(exported->offset, 512u); +} + +TEST_F(TestUsrbIoGdrFixture, CreateOwnsAndRemovesPublication) { + if (!hasGpu()) { + GTEST_SKIP() << "No CUDA IPC-capable GPU available"; + } + + TmpDir tmpDir; + ASSERT_TRUE(tmpDir.valid()); + std::filesystem::create_directories(std::string(tmpDir.path()) + "/3fs-virt/iovs"); + + int deviceId = -1; + ASSERT_EQ(cudaGetDevice(&deviceId), cudaSuccess); + struct hf3fs_iov iov{}; + ASSERT_EQ(hf3fs_iovcreate_device(&iov, tmpDir.path(), 4096, 0, deviceId), 0); + SCOPE_EXIT { + if (iov.iovh) { + hf3fs_iovdestroy(&iov); + } + }; + + auto link = publicationPath(iov, deviceId); + EXPECT_TRUE(isSymlink(link)); + EXPECT_EQ(hf3fs_iov_mem_type(&iov), HF3FS_MEM_DEVICE); + + hf3fs_iovdestroy(&iov); + EXPECT_FALSE(isSymlink(link)); +} + +TEST_F(TestUsrbIoGdrFixture, ExplicitUnlinkPreventsDestroyFromRemovingReplacementPublication) { + if (!hasGpu()) { + GTEST_SKIP() << "No CUDA IPC-capable GPU available"; + } + + TmpDir tmpDir; + ASSERT_TRUE(tmpDir.valid()); + std::filesystem::create_directories(std::string(tmpDir.path()) + "/3fs-virt/iovs"); + + int deviceId = -1; + ASSERT_EQ(cudaGetDevice(&deviceId), cudaSuccess); + struct hf3fs_iov iov{}; + ASSERT_EQ(hf3fs_iovcreate_device(&iov, tmpDir.path(), 4096, 0, deviceId), 0); + SCOPE_EXIT { + if (iov.iovh) { + hf3fs_iovdestroy(&iov); + } + }; + + auto link = publicationPath(iov, deviceId); + hf3fs_iovunlink(&iov); + ASSERT_FALSE(isSymlink(link)); + ASSERT_EQ(symlink("gdr://replacement-publication", link.c_str()), 0); + + hf3fs_iovdestroy(&iov); + EXPECT_EQ(iov.iovh, nullptr); + EXPECT_TRUE(isSymlink(link)); + EXPECT_EQ(unlink(link.c_str()), 0); +} + +TEST_F(TestUsrbIoGdrFixture, DestroyRetainsGpuHandleWhenPublicationUnlinkFails) { + if (!hasGpu()) { + GTEST_SKIP() << "No CUDA IPC-capable GPU available"; + } + + TmpDir tmpDir; + ASSERT_TRUE(tmpDir.valid()); + std::filesystem::create_directories(std::string(tmpDir.path()) + "/3fs-virt/iovs"); + + int deviceId = -1; + ASSERT_EQ(cudaGetDevice(&deviceId), cudaSuccess); + struct hf3fs_iov iov{}; + ASSERT_EQ(hf3fs_iovcreate_device(&iov, tmpDir.path(), 4096, 0, deviceId), 0); + SCOPE_EXIT { + if (iov.iovh) { + hf3fs_iovdestroy(&iov); + } + }; + + auto link = publicationPath(iov, deviceId); + ASSERT_EQ(unlink(link.c_str()), 0); + ASSERT_TRUE(std::filesystem::create_directory(link)); + + auto *handle = iov.iovh; + hf3fs_iovdestroy(&iov); + EXPECT_EQ(iov.iovh, handle); + EXPECT_EQ(hf3fs_iov_mem_type(&iov), HF3FS_MEM_DEVICE); + + ASSERT_TRUE(std::filesystem::remove(link)); + hf3fs_iovdestroy(&iov); + EXPECT_EQ(iov.iovh, nullptr); +} + +TEST_F(TestUsrbIoGdrFixture, ImporterProcessKeepsPublisherPublication) { + const char *mount = std::getenv("HF3FS_GDR_IMPORT_MOUNT"); + if (!mount) { + GTEST_SKIP() << "Subprocess-only importer check"; + } + + const char *idText = std::getenv("HF3FS_GDR_IMPORT_ID"); + const char *sizeText = std::getenv("HF3FS_GDR_IMPORT_SIZE"); + const char *deviceText = std::getenv("HF3FS_GDR_IMPORT_DEVICE"); + const char *expectedText = std::getenv("HF3FS_GDR_IMPORT_EXPECTED_BYTE"); + ASSERT_NE(idText, nullptr); + ASSERT_NE(sizeText, nullptr); + ASSERT_NE(deviceText, nullptr); + ASSERT_NE(expectedText, nullptr); + + auto parsedId = hf3fs::Uuid::fromHexString(idText); + ASSERT_OK(parsedId); + std::array id{}; + std::memcpy(id.data(), parsedId->data, id.size()); + auto size = static_cast(std::strtoull(sizeText, nullptr, 10)); + auto deviceId = std::atoi(deviceText); + auto expectedByte = static_cast(std::atoi(expectedText)); + + struct hf3fs_iov importer{}; + ASSERT_EQ(hf3fs_iovopen_device(&importer, id.data(), mount, size, 0, deviceId), 0); + SCOPE_EXIT { + if (importer.iovh) { + hf3fs_iovdestroy(&importer); + } + }; + + auto link = publicationPath(importer, deviceId); + ASSERT_TRUE(isSymlink(link)); + EXPECT_EQ(hf3fs_iovsync(&importer, 0), 0); + uint8_t actual = 0; + ASSERT_EQ(cudaMemcpy(&actual, importer.base, 1, cudaMemcpyDeviceToHost), cudaSuccess); + EXPECT_EQ(actual, expectedByte); + + hf3fs_iovdestroy(&importer); + EXPECT_TRUE(isSymlink(link)); +} + +TEST_F(TestUsrbIoGdrFixture, WrapPublishesBaseOffsetAndSupportsMultipleNonOwningImporters) { + if (!hasGpu()) { + GTEST_SKIP() << "No CUDA IPC-capable GPU available"; + } + if (currentExecutable().empty()) { + GTEST_SKIP() << "Subprocess execution through /proc/self/exe is unavailable"; + } + + TmpDir tmpDir; + ASSERT_TRUE(tmpDir.valid()); + std::filesystem::create_directories(std::string(tmpDir.path()) + "/3fs-virt/iovs"); + + int deviceId = -1; + ASSERT_EQ(cudaGetDevice(&deviceId), cudaSuccess); + void *allocation = nullptr; + ASSERT_EQ(cudaMalloc(&allocation, 4096), cudaSuccess); + SCOPE_EXIT { + if (allocation) { + EXPECT_EQ(cudaFree(allocation), cudaSuccess); + } + }; + ASSERT_EQ(cudaMemset(allocation, 0x11, 4096), cudaSuccess); + auto *view = deviceAddress(allocation, 512); + ASSERT_EQ(cudaMemset(view, 0x7A, 1024), cudaSuccess); + ASSERT_EQ(cudaDeviceSynchronize(), cudaSuccess); + + struct hf3fs_iov publisher{}; + auto id = hf3fs::Uuid::random(); + ASSERT_EQ(hf3fs_iovwrap_device(&publisher, + view, + reinterpret_cast(id.data), + tmpDir.path(), + 1024, + 0, + deviceId), + 0); + SCOPE_EXIT { + if (publisher.iovh) { + hf3fs_iovdestroy(&publisher); + } + }; + + auto link = publicationPath(publisher, deviceId); + ASSERT_TRUE(isSymlink(link)); + auto target = std::filesystem::read_symlink(link).string(); + auto parsed = hf3fs::lib::parseGdrUri(target); + ASSERT_TRUE(parsed); + EXPECT_EQ(parsed->deviceId, deviceId); + EXPECT_EQ(parsed->offset, 512u); + EXPECT_EQ(parsed->size, 1024u); + EXPECT_GE(parsed->allocationSize, parsed->offset + parsed->size); + + EXPECT_TRUE(runImporterProcess(publisher, deviceId, 0x7A)); + EXPECT_TRUE(isSymlink(link)); + EXPECT_TRUE(runImporterProcess(publisher, deviceId, 0x7A)); + EXPECT_TRUE(isSymlink(link)); + + hf3fs_iovdestroy(&publisher); + EXPECT_FALSE(isSymlink(link)); + + uint8_t actual = 0; + ASSERT_EQ(cudaMemcpy(&actual, view, 1, cudaMemcpyDeviceToHost), cudaSuccess); + EXPECT_EQ(actual, 0x7A); + ASSERT_EQ(cudaFree(allocation), cudaSuccess); + allocation = nullptr; +} + +// ========================================================================== +// REQ-L4-002: iovopen/iovwrap host path; _device variants for GPU +// ========================================================================== + +// @tests SCN-L4-002-00b +TEST_F(TestUsrbIoGdrFixture, SCN_L4_002_00b_IovWrapNegativeNumaHostPath) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {}; + // Use a valid host pointer so iovwrap succeeds (it doesn't validate mount existence) + uint8_t buf[64] = {}; + + // WHEN: iovwrap with negative numa (= host path, no NUMA binding) + int rc = hf3fs_iovwrap(&iov, buf, id, "/nonexistent", sizeof(buf), 0, -1); + + // THEN: Succeeds (iovwrap doesn't validate mount), host path, no GPU dispatch + EXPECT_EQ(rc, 0); + // Verify it's host memory, not GPU + EXPECT_EQ(hf3fs_iov_mem_type(&iov), HF3FS_MEM_HOST); +} + +TEST_F(TestUsrbIoGdrFixture, PrepIoUsesOverflowSafeOpaquePointerRanges) { + struct hf3fs_ior ior{}; + ior.iorh = reinterpret_cast(uintptr_t{1}); + ior.for_read = true; + struct hf3fs_iov iov{}; + + iov.base = reinterpret_cast(uintptr_t{0x1000}); + iov.size = 0x100; + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, nullptr, 0, 0, 1, nullptr), -EINVAL); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, iov.base, 0, 0, 0, nullptr), -EINVAL); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(uintptr_t{0x0fff}), 0, 0, 1, nullptr), -EINVAL); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(uintptr_t{0x1100}), 0, 0, 1, nullptr), -EINVAL); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(uintptr_t{0x10ff}), 0, 0, 2, nullptr), -EINVAL); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(uintptr_t{0x10ff}), 0, 0, 1, nullptr), -EBADF); + + auto maxAddress = std::numeric_limits::max(); + iov.base = reinterpret_cast(maxAddress - 0xff); + iov.size = 0xff; + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(maxAddress - 8), 0, 0, 8, nullptr), -EBADF); + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, reinterpret_cast(maxAddress - 8), 0, 0, 9, nullptr), -EINVAL); + + iov.base = reinterpret_cast(maxAddress - 7); + iov.size = 8; + EXPECT_EQ(hf3fs_prep_io(&ior, &iov, true, iov.base, 0, 0, 1, nullptr), -EINVAL); +} + +// @tests SCN-L4-002-01 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_002_01_IovOpenDeviceNoGdr) { + if (hasGpu()) { + GTEST_SKIP() << "Test for non-GPU environment"; + } + + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {}; + + // WHEN: iovopen_device, GDR unavailable + int rc = hf3fs_iovopen_device(&iov, id, "/nonexistent", 4096, 0, 0); + + // THEN: Returns -ENOTSUP + EXPECT_EQ(rc, -ENOTSUP); +} + +// @tests SCN-L4-002-02 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_002_02_IovWrapDeviceNoGdr) { + if (hasGpu()) { + GTEST_SKIP() << "Test for non-GPU environment"; + } + + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {}; + void *fakePtr = reinterpret_cast(0x1000); + + // WHEN: iovwrap_device, GDR unavailable + int rc = hf3fs_iovwrap_device(&iov, fakePtr, id, "/nonexistent", 4096, 0, 0); + + // THEN: Returns -ENOTSUP + EXPECT_EQ(rc, -ENOTSUP); +} + +// ========================================================================== +// REQ-L4-003: Unified iovdestroy Dispatch +// ========================================================================== + +// @tests SCN-L4-003-02 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_003_02_DestroyHostIov) { + // GIVEN: A zeroed iov (host-like, numa >= 0) + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = 0; + + // WHEN: iovdestroy is called + // THEN: No crash (best-effort cleanup on zeroed iov) + hf3fs_iovdestroy(&iov); + + // Verify iov was zeroed/cleaned + EXPECT_EQ(iov.base, nullptr); +} + +// ========================================================================== +// REQ-L4-004: Query Functions +// ========================================================================== + +// @tests SCN-L4-004-01 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_004_01_MemTypeHostIov) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = 0; + + // WHEN: mem_type query + enum hf3fs_mem_type type = hf3fs_iov_mem_type(&iov); + + // THEN: HF3FS_MEM_HOST + EXPECT_EQ(type, HF3FS_MEM_HOST); +} + +// @tests SCN-L4-004-01 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_004_01_DeviceIdHostIov) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = 0; + + // WHEN: device_id query on host iov + int devId = hf3fs_iov_device_id(&iov); + + // THEN: Returns -1 + EXPECT_EQ(devId, -1); +} + +// @tests SCN-L4-004-02 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_004_02_GdrAvailableLazyInit) { + // WHEN: hf3fs_gdr_available is called + bool avail = hf3fs_gdr_available(); + + // THEN: Returns a valid boolean, consistent across calls + bool avail2 = hf3fs_gdr_available(); + EXPECT_EQ(avail, avail2); +} + +// @tests REQ-L4-004 +TEST_F(TestUsrbIoGdrFixture, GdrDeviceCount) { + int count = hf3fs_gdr_device_count(); + if (hasGpu()) { + EXPECT_GT(count, 0); + } else { + EXPECT_EQ(count, 0); + } +} + +// ========================================================================== +// REQ-L4-006: iovsync Dispatch +// ========================================================================== + +// @tests SCN-L4-006-02 +TEST_F(TestUsrbIoGdrFixture, SCN_L4_006_02_SyncHostIov) { + // GIVEN: Host iov + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = 0; + + // WHEN: iovsync is called + int rc = hf3fs_iovsync(&iov, 1); + + // THEN: Returns 0 (no-op for host) + EXPECT_EQ(rc, 0); +} + +// @tests SCN-L4-006-02 +TEST_F(TestUsrbIoGdrFixture, SyncHostIovDirection0) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = 0; + + int rc = hf3fs_iovsync(&iov, 0); + EXPECT_EQ(rc, 0); +} + +// ========================================================================== +// REQ-L3-005: GDR URI Parsing (Pure Logic) +// ========================================================================== + +// @tests SCN-L3-005-01 +TEST_F(TestUsrbIoGdrFixture, SCN_L3_005_01_ValidUriFormatThroughIovopen) { + // GIVEN: A valid URI in symlink form + // Build a proper URI and verify format through the API path + uint8_t ipcHandle[64]; + for (int i = 0; i < 64; i++) ipcHandle[i] = static_cast(i * 3 + 7); + + std::string uri = buildGdrUri(0, 1073741824, 0, 1073741824, ipcHandle); + + // Verify URI has correct format + EXPECT_EQ(uri.substr(0, 14), "gdr://v2/devic"); + EXPECT_NE(uri.find("/device/0/"), std::string::npos); + EXPECT_NE(uri.find("/allocation/1073741824/"), std::string::npos); + EXPECT_NE(uri.find("/offset/0/"), std::string::npos); + EXPECT_NE(uri.find("/size/1073741824/"), std::string::npos); + EXPECT_NE(uri.find("/ipc/"), std::string::npos); + // IPC hex should be 128 chars + size_t ipcPos = uri.find("/ipc/"); + ASSERT_NE(ipcPos, std::string::npos); + std::string hexPart = uri.substr(ipcPos + 5); + EXPECT_EQ(hexPart.size(), 128u); + + // On GPU machines, test through iovopen with a crafted symlink + if (hasGpu()) { + TmpDir tmpDir; + if (tmpDir.valid()) { + // Create the symlink structure: {mount}/3fs-virt/iovs/{uuid}.gdr.d0 -> uri + std::string virtDir = std::string(tmpDir.path()) + "/3fs-virt/iovs"; + std::filesystem::create_directories(virtDir); + + // Create a symlink with the URI + std::string symlinkPath = virtDir + "/deadbeef12345678deadbeef12345678.gdr.d0"; + symlink(uri.c_str(), symlinkPath.c_str()); + + // Try to open — will fail because UUID doesn't match, but exercises parser + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {0xDE, 0xAD, 0xBE, 0xEF, 0x12, 0x34, 0x56, 0x78, 0xDE, 0xAD, 0xBE, 0xEF, 0x12, 0x34, 0x56, 0x78}; + int rc = hf3fs_iovopen_device(&iov, id, tmpDir.path(), 1073741824, 0, 0); + // May fail for various reasons, but should not crash + (void)rc; + } + } +} + +// ========================================================================== +// REQ-L3-003: GPU IOV Wrap (External Memory) +// ========================================================================== + +// @tests SCN-L3-003-01 +TEST_F(TestUsrbIoGdrFixture, SCN_L3_003_01_WrapExternalGpuPtr) { + if (!hasGpu()) { + // On CPU-only, iovwrap_device returns -ENOTSUP + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + uint8_t id[16] = {}; + void *fakePtr = reinterpret_cast(0x2000); + int rc = hf3fs_iovwrap_device(&iov, fakePtr, id, "/nonexistent", 4096, 0, 0); + EXPECT_EQ(rc, -ENOTSUP); + return; + } + + // On GPU: wrap would need real GPU pointer + mount + GTEST_SKIP() << "Requires real GPU memory and mount — integration test only"; +} + +// ========================================================================== +// INV-GDR-001: internal handle registry is the polymorphism discriminant +// ========================================================================== + +// @tests INV-GDR-001 +TEST_F(TestUsrbIoGdrFixture, INV_GDR_001_PolymorphismSafety) { + // GIVEN: An iov with the former magic NUMA value but no registered handle + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + iov.numa = -0x6472; + iov.iovh = nullptr; + + // WHEN: Query operations are called + // THEN: They must not crash and must recognize this as non-GPU + enum hf3fs_mem_type type = hf3fs_iov_mem_type(&iov); + EXPECT_EQ(type, HF3FS_MEM_HOST); + + int devId = hf3fs_iov_device_id(&iov); + EXPECT_EQ(devId, -1); + + // Sync remains a host no-op because NUMA does not classify the iov. + int rc = hf3fs_iovsync(&iov, 0); + EXPECT_EQ(rc, 0); +} + +// @tests INV-GDR-002 +TEST_F(TestUsrbIoGdrFixture, INV_GDR_002_NumaDoesNotClassifyGpuIov) { + struct hf3fs_iov iov; + memset(&iov, 0, sizeof(iov)); + + iov.numa = 0; + EXPECT_EQ(hf3fs_iov_mem_type(&iov), HF3FS_MEM_HOST); + + iov.numa = -0x6472; + EXPECT_EQ(hf3fs_iov_mem_type(&iov), HF3FS_MEM_HOST); + EXPECT_EQ(hf3fs_iov_device_id(&iov), -1); +} diff --git a/tests/storage/client/TestStorageClientSideError.cc b/tests/storage/client/TestStorageClientSideError.cc index 13981f73..b29e4829 100644 --- a/tests/storage/client/TestStorageClientSideError.cc +++ b/tests/storage/client/TestStorageClientSideError.cc @@ -1,8 +1,11 @@ +#include #include #include +#include #include "client/mgmtd/ICommonMgmtdClient.h" #include "client/storage/StorageClient.h" +#include "client/storage/StorageClientImpl.h" #include "common/net/Client.h" #include "tests/lib/UnitTestFabric.h" @@ -48,6 +51,79 @@ class TestStorageClientSideError : public UnitTestFabric, public ::testing::Test } }; +TEST(StorageClientChecksumRequest, DeviceChecksumUsesServerCompute) { + auto prepared = + prepareWriteChecksum(ChecksumType::CRC32C, true, true, reinterpret_cast(0x1000), 4096); + + ASSERT_OK(prepared); + EXPECT_EQ(prepared->checksum, (ChecksumInfo{ChecksumType::CRC32C, 0})); + EXPECT_TRUE(prepared->computeOnServer); +} + +TEST(StorageClientChecksumRequest, DeviceChecksumSetsRequestFeatureFlag) { + DebugOptions debug; + + auto flags = buildWriteFeatureFlagsFromOptions(debug, true); + EXPECT_TRUE(BITFLAGS_CONTAIN(flags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); + + flags = buildWriteFeatureFlagsFromOptions(debug, false); + EXPECT_FALSE(BITFLAGS_CONTAIN(flags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); +} + +TEST(StorageClientChecksumRequest, HostChecksumIsMaterializedLocally) { + const std::array data{1, 3, 5, 7, 9}; + auto prepared = prepareWriteChecksum(ChecksumType::CRC32C, true, false, data.data(), data.size()); + + ASSERT_OK(prepared); + EXPECT_EQ(prepared->checksum, ChecksumInfo::create(ChecksumType::CRC32C, data.data(), data.size())); + EXPECT_FALSE(prepared->computeOnServer); +} + +TEST(StorageClientChecksumRequest, DisabledVerificationUsesNone) { + auto prepared = + prepareWriteChecksum(ChecksumType::NONE, false, true, reinterpret_cast(0x1000), 4096); + + ASSERT_OK(prepared); + EXPECT_EQ(prepared->checksum, (ChecksumInfo{ChecksumType::NONE, 0})); + EXPECT_FALSE(prepared->computeOnServer); +} + +TEST(StorageClientChecksumRequest, ConfiguredNoneDisablesGpuChecksumVerification) { + auto prepared = prepareWriteChecksum(ChecksumType::NONE, true, true, reinterpret_cast(0x1000), 4096); + + ASSERT_OK(prepared); + EXPECT_EQ(prepared->checksum, (ChecksumInfo{ChecksumType::NONE, 0})); + EXPECT_FALSE(prepared->computeOnServer); + + DebugOptions debug; + auto flags = buildWriteFeatureFlagsFromOptions(debug, prepared->computeOnServer); + EXPECT_FALSE(BITFLAGS_CONTAIN(flags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); +} + +TEST_F(TestStorageClientSideError, ZeroHostIOBufferRange) { + std::vector data(64, 0xAB); + auto registered = storageClient_->registerIOBuffer(data.data(), data.size()); + ASSERT_OK(registered); + auto ioBuffer = std::move(*registered); + + ASSERT_OK(ioBuffer.zeroRange(8, 24)); + for (size_t i = 0; i < data.size(); ++i) { + EXPECT_EQ(data[i], i >= 8 && i < 32 ? 0 : 0xAB); + } + ASSERT_OK(ioBuffer.zeroRange(data.size(), 0)); + ASSERT_ERROR(ioBuffer.zeroRange(63, 2), StorageClientCode::kInvalidArg); + ASSERT_ERROR(ioBuffer.zeroRange(std::numeric_limits::max(), 0), StorageClientCode::kInvalidArg); +} + +TEST_F(TestStorageClientSideError, InvalidGpuRegistrationNeverReturnsHostBuffer) { + auto registered = storageClient_->registerGpuIOBuffer(reinterpret_cast(0x1000), 4096); + ASSERT_FALSE(registered); + + std::array host{}; + registered = storageClient_->registerGpuIOBuffer(host.data(), host.size()); + ASSERT_FALSE(registered); +} + TEST_F(TestStorageClientSideError, GetReplicationChainError) { updateRoutingInfo([&](auto &routingInfo) { auto &chainTable = *routingInfo.getChainTable(kTableId()); diff --git a/tests/storage/store/TestCommonStruct.cc b/tests/storage/store/TestCommonStruct.cc index 9043d86a..4732a50b 100644 --- a/tests/storage/store/TestCommonStruct.cc +++ b/tests/storage/store/TestCommonStruct.cc @@ -80,5 +80,25 @@ TEST(TestCommonStruct, Normal) { }); } +TEST(TestCommonStruct, ServerComputeChecksumFlagRoundTripsAsUint32) { + WriteReq writeReq; + BITFLAGS_SET(writeReq.featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM); + ASSERT_EQ(writeReq.featureFlags, uint32_t{16}); + + auto serialized = serde::serialize(writeReq); + WriteReq deserializedWrite; + ASSERT_OK(serde::deserialize(deserializedWrite, serialized)); + EXPECT_EQ(deserializedWrite.featureFlags, writeReq.featureFlags); + EXPECT_TRUE(BITFLAGS_CONTAIN(deserializedWrite.featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); + + UpdateReq updateReq; + updateReq.featureFlags = deserializedWrite.featureFlags; + serialized = serde::serialize(updateReq); + UpdateReq deserializedUpdate; + ASSERT_OK(serde::deserialize(deserializedUpdate, serialized)); + EXPECT_EQ(deserializedUpdate.featureFlags, updateReq.featureFlags); + EXPECT_TRUE(BITFLAGS_CONTAIN(deserializedUpdate.featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); +} + } // namespace } // namespace hf3fs::storage::test diff --git a/tests/storage/store/TestStorageTarget.cc b/tests/storage/store/TestStorageTarget.cc index a9259195..f4c01e25 100644 --- a/tests/storage/store/TestStorageTarget.cc +++ b/tests/storage/store/TestStorageTarget.cc @@ -1,10 +1,14 @@ +#include +#include #include #include #include "common/utils/Size.h" #include "fbs/storage/Common.h" #include "storage/aio/AioReadWorker.h" +#include "storage/service/StorageOperator.h" #include "storage/service/TargetMap.h" +#include "storage/store/ChunkReplica.h" #include "storage/store/ChunkStore.h" #include "storage/store/StorageTarget.h" #include "storage/update/UpdateWorker.h" @@ -13,6 +17,80 @@ namespace hf3fs::storage { namespace { +TEST(ServerComputedChecksum, ValidatesProtocolBeforeStorageUpdate) { + uint32_t featureFlags = 0; + BITFLAGS_SET(featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM); + + UpdateIO update; + update.updateType = UpdateType::REMOVE; + update.checksum.type = ChecksumType::CRC32C; + ASSERT_ERROR(materializeServerComputedChecksum(update, featureFlags, nullptr), StatusCode::kInvalidArg); + + update.updateType = UpdateType::WRITE; + update.length = 1; + update.checksum.type = ChecksumType::NONE; + ASSERT_ERROR(materializeServerComputedChecksum(update, featureFlags, nullptr), StatusCode::kInvalidArg); + + update.checksum.type = ChecksumType::CRC32C; + ASSERT_ERROR(materializeServerComputedChecksum(update, featureFlags, nullptr), StatusCode::kInvalidArg); +} + +TEST(ServerComputedChecksum, MaterializesAndPreservesFlagForNewServerForwarding) { + const std::array data{2, 3, 5, 7, 11, 13, 17}; + UpdateReq headReq; + headReq.payload.updateType = UpdateType::WRITE; + headReq.payload.length = data.size(); + headReq.payload.checksum = ChecksumInfo{ChecksumType::CRC32C, 0}; + headReq.options.fromClient = true; + BITFLAGS_SET(headReq.featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM); + + auto noFlag = headReq.payload; + ASSERT_OK(materializeServerComputedChecksum(noFlag, 0, nullptr)); + EXPECT_EQ(noFlag.checksum, (ChecksumInfo{ChecksumType::CRC32C, 0})); + + ASSERT_OK(materializeServerComputedChecksum(headReq.payload, headReq.featureFlags, data.data())); + auto expected = ChecksumInfo::create(ChecksumType::CRC32C, data.data(), data.size()); + EXPECT_EQ(headReq.payload.checksum, expected); + + auto successorReq = headReq; + successorReq.options.fromClient = false; + EXPECT_TRUE(BITFLAGS_CONTAIN(successorReq.featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM)); + ASSERT_OK(materializeServerComputedChecksum(successorReq.payload, successorReq.featureFlags, data.data())); + EXPECT_EQ(successorReq.payload.checksum, expected); +} + +TEST(ServerComputedChecksum, ChunkReplicaAppendPreservesMetadataAndCombinesChecksum) { + const std::array prefix{1, 2, 3, 4}; + const std::array suffix{5, 6, 7}; + std::array all{}; + std::copy(prefix.begin(), prefix.end(), all.begin()); + std::copy(suffix.begin(), suffix.end(), all.begin() + prefix.size()); + + ChunkInfo chunkInfo; + chunkInfo.meta.size = all.size(); + chunkInfo.meta.commitVer = ChunkVer{8}; + chunkInfo.meta.updateVer = ChunkVer{9}; + chunkInfo.meta.chainVer = ChainVer{10}; + chunkInfo.meta.chunkState = ChunkState::DIRTY; + auto prefixChecksum = ChecksumInfo::create(ChecksumType::CRC32C, prefix.data(), prefix.size()); + chunkInfo.meta.checksumType = prefixChecksum.type; + chunkInfo.meta.checksumValue = prefixChecksum.value; + + UpdateIO append; + append.updateType = UpdateType::WRITE; + append.offset = prefix.size(); + append.length = suffix.size(); + append.checksum = ChecksumInfo::create(ChecksumType::CRC32C, suffix.data(), suffix.size()); + + ASSERT_OK(ChunkReplica::updateChecksum(chunkInfo, append, prefix.size(), true)); + EXPECT_EQ(chunkInfo.meta.checksum(), ChecksumInfo::create(ChecksumType::CRC32C, all.data(), all.size())); + EXPECT_EQ(chunkInfo.meta.size, all.size()); + EXPECT_EQ(chunkInfo.meta.commitVer, ChunkVer{8}); + EXPECT_EQ(chunkInfo.meta.updateVer, ChunkVer{9}); + EXPECT_EQ(chunkInfo.meta.chainVer, ChainVer{10}); + EXPECT_EQ(chunkInfo.meta.chunkState, ChunkState::DIRTY); +} + TEST(TestStorageTarget, ReadWrite) { folly::test::TemporaryDirectory tmpPath; CPUExecutorGroup executor(8, ""); @@ -74,9 +152,17 @@ TEST(TestStorageTarget, ReadWrite) { writeIO.offset = 0; writeIO.updateVer = ChunkVer{1}; writeIO.updateType = UpdateType::WRITE; + writeIO.checksum = ChecksumInfo{ChecksumType::CRC32C, 0}; + + uint32_t featureFlags = 0; + BITFLAGS_SET(featureFlags, FeatureFlags::SERVER_COMPUTE_CHECKSUM); + auto data = reinterpret_cast(dataBytes.data()); + ASSERT_OK(materializeServerComputedChecksum(writeIO, featureFlags, data)); + auto expectedChecksum = ChecksumInfo::create(ChecksumType::CRC32C, data, chunkSize); + ASSERT_EQ(writeIO.checksum, expectedChecksum); UpdateJob updateJob(requestCtx, writeIO, {}, updateChunk, storageTarget); - updateJob.state().data = reinterpret_cast(dataBytes.data()); + updateJob.state().data = data; folly::coro::blockingWait(updateWorker.enqueue(&updateJob)); folly::coro::blockingWait(updateJob.complete()); @@ -84,6 +170,10 @@ TEST(TestStorageTarget, ReadWrite) { ASSERT_EQ(updateJob.result().lengthInfo.value(), chunkSize); ASSERT_EQ(updateJob.result().commitVer, ChunkVer{0}); ASSERT_EQ(updateJob.result().updateVer, ChunkVer{1}); + + auto metadata = storageTarget->queryChunk(chunkId); + ASSERT_OK(metadata); + EXPECT_EQ(metadata->checksum(), expectedChecksum); } ASSERT_EQ(storageTarget->usedSize(), chunkSize);