From 2cff93eaa2e79b4e78dadb59b4cca74bde14a22d Mon Sep 17 00:00:00 2001 From: Jiajun Huang Date: Fri, 6 Nov 2026 13:22:11 -0500 Subject: [PATCH] Rename to SZ-GPU SZ-GPU (FSZ) is the production GPU library for the SZ compressor family. This is version 1.0.1. The library, headers, C and C++ API, command-line tool, Python package, Fortran module and CMake package move from fsz to szgpu (SZGPU_ for macros, CMake options and environment variables), and the README, messages and comments name SZ-GPU. The container magic stays FSZ1 and the bitstream format is unchanged, so files written by version 1.0.0 decompress unchanged. The GPU kernels are unchanged. --- .gitignore | 2 +- CMakeLists.txt | 146 +++--- CONTRIBUTING.md | 18 +- Makefile | 28 +- README.md | 170 +++---- VERSION | 2 +- cmake/fszConfig.cmake.in | 8 - cmake/szgpuConfig.cmake.in | 8 + examples/example_hostptr.cu | 18 +- examples/example_roundtrip.cu | 12 +- .../{example_fsz.f90 => example_szgpu.f90} | 22 +- fortran/{fsz.f90 => szgpu.f90} | 256 +++++----- fortran/{test_fsz.f90 => test_szgpu.f90} | 80 ++-- include/fsz/fsz.h | 113 ----- include/fsz/fsz.hpp | 176 ------- include/szgpu/szgpu.h | 113 +++++ include/szgpu/szgpu.hpp | 176 +++++++ python/MANIFEST.in | 2 +- python/example.py | 22 +- python/pyproject.toml | 8 +- python/setup.py | 26 +- python/{fsz => szgpu}/__init__.py | 132 +++--- python/tests/{test_fsz.py => test_szgpu.py} | 116 ++--- src/fsz.cu | 442 ------------------ src/szgpu.cu | 442 ++++++++++++++++++ src/{fsz_internal.cuh => szgpu_internal.cuh} | 74 +-- src/{fsz_kernels.cuh => szgpu_kernels.cuh} | 144 +++--- tests/test_hostptr.cu | 124 ++--- tests/test_roundtrip.cu | 18 +- tools/{fsz.cu => szgpu.cu} | 108 ++--- ..._file_format.hpp => szgpu_file_format.hpp} | 10 +- 31 files changed, 1509 insertions(+), 1507 deletions(-) delete mode 100644 cmake/fszConfig.cmake.in create mode 100644 cmake/szgpuConfig.cmake.in rename fortran/{example_fsz.f90 => example_szgpu.f90} (67%) rename fortran/{fsz.f90 => szgpu.f90} (69%) rename fortran/{test_fsz.f90 => test_szgpu.f90} (71%) delete mode 100644 include/fsz/fsz.h delete mode 100644 include/fsz/fsz.hpp create mode 100644 include/szgpu/szgpu.h create mode 100644 include/szgpu/szgpu.hpp rename python/{fsz => szgpu}/__init__.py (84%) rename python/tests/{test_fsz.py => test_szgpu.py} (80%) delete mode 100644 src/fsz.cu create mode 100644 src/szgpu.cu rename src/{fsz_internal.cuh => szgpu_internal.cuh} (67%) rename src/{fsz_kernels.cuh => szgpu_kernels.cuh} (80%) rename tools/{fsz.cu => szgpu.cu} (91%) rename tools/{fsz_file_format.hpp => szgpu_file_format.hpp} (91%) diff --git a/.gitignore b/.gitignore index d71dcaa..8c69602 100644 --- a/.gitignore +++ b/.gitignore @@ -4,7 +4,7 @@ cmake-build/ *.o *.a *.so -*.fsz +*.szgpu *.dec.f32 *.log __pycache__/ diff --git a/CMakeLists.txt b/CMakeLists.txt index 69c2581..6b66598 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -5,12 +5,12 @@ cmake_minimum_required(VERSION 3.18) # automatically when the CUDA toolkit is new enough to know it. if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) set(CMAKE_CUDA_ARCHITECTURES 80 90) - set(_fsz_default_archs ON) + set(_szgpu_default_archs ON) endif() -project(fsz LANGUAGES CXX CUDA VERSION 1.0.0) +project(szgpu LANGUAGES CXX CUDA VERSION 1.0.1) -if(_fsz_default_archs AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8) +if(_szgpu_default_archs AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8) list(APPEND CMAKE_CUDA_ARCHITECTURES 100) endif() @@ -29,126 +29,126 @@ endif() find_package(CUDAToolkit REQUIRED) -option(FSZ_BUILD_SHARED "Also build the shared library (libfsz.so)" ON) -option(FSZ_BUILD_TOOLS "Build the fsz command-line tool" ON) -option(FSZ_BUILD_EXAMPLES "Build FSZ examples" ON) -option(FSZ_BUILD_TESTS "Build FSZ tests" ON) +option(SZGPU_BUILD_SHARED "Also build the shared library (libszgpu.so)" ON) +option(SZGPU_BUILD_TOOLS "Build the szgpu command-line tool" ON) +option(SZGPU_BUILD_EXAMPLES "Build SZ-GPU examples" ON) +option(SZGPU_BUILD_TESTS "Build SZ-GPU tests" ON) -set(FSZ_SOURCES src/fsz.cu) +set(SZGPU_SOURCES src/szgpu.cu) -# Static library (the default link target, exported as fsz::fsz). -add_library(fsz STATIC ${FSZ_SOURCES}) -add_library(fsz::fsz ALIAS fsz) +# Static library (the default link target, exported as szgpu::szgpu). +add_library(szgpu STATIC ${SZGPU_SOURCES}) +add_library(szgpu::szgpu ALIAS szgpu) # Shared library for dynamic linking and the Python bindings. -if(FSZ_BUILD_SHARED) - add_library(fsz_shared SHARED ${FSZ_SOURCES}) - set_target_properties(fsz_shared PROPERTIES OUTPUT_NAME fsz) +if(SZGPU_BUILD_SHARED) + add_library(szgpu_shared SHARED ${SZGPU_SOURCES}) + set_target_properties(szgpu_shared PROPERTIES OUTPUT_NAME szgpu) endif() -foreach(_fsz_lib IN ITEMS fsz fsz_shared) - if(NOT TARGET ${_fsz_lib}) +foreach(_szgpu_lib IN ITEMS szgpu szgpu_shared) + if(NOT TARGET ${_szgpu_lib}) continue() endif() - target_include_directories(${_fsz_lib} + target_include_directories(${_szgpu_lib} PUBLIC $ $ PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/src) - target_link_libraries(${_fsz_lib} PUBLIC CUDA::cudart) - set_target_properties(${_fsz_lib} PROPERTIES + target_link_libraries(${_szgpu_lib} PUBLIC CUDA::cudart) + set_target_properties(${_szgpu_lib} PROPERTIES CUDA_SEPARABLE_COMPILATION OFF POSITION_INDEPENDENT_CODE ON) - target_compile_options(${_fsz_lib} PRIVATE + target_compile_options(${_szgpu_lib} PRIVATE $<$:-O3>) endforeach() -# Command-line tool: the binary is named `fsz`. -if(FSZ_BUILD_TOOLS) - add_executable(fsz_cli tools/fsz.cu) - set_target_properties(fsz_cli PROPERTIES OUTPUT_NAME fsz) - target_include_directories(fsz_cli PRIVATE tools) - target_link_libraries(fsz_cli PRIVATE fsz::fsz) +# Command-line tool: the binary is named `szgpu`. +if(SZGPU_BUILD_TOOLS) + add_executable(szgpu_cli tools/szgpu.cu) + set_target_properties(szgpu_cli PROPERTIES OUTPUT_NAME szgpu) + target_include_directories(szgpu_cli PRIVATE tools) + target_link_libraries(szgpu_cli PRIVATE szgpu::szgpu) endif() -if(FSZ_BUILD_EXAMPLES) - add_executable(fsz_example_roundtrip examples/example_roundtrip.cu) - target_link_libraries(fsz_example_roundtrip PRIVATE fsz::fsz) +if(SZGPU_BUILD_EXAMPLES) + add_executable(szgpu_example_roundtrip examples/example_roundtrip.cu) + target_link_libraries(szgpu_example_roundtrip PRIVATE szgpu::szgpu) - add_executable(fsz_example_hostptr examples/example_hostptr.cu) - target_link_libraries(fsz_example_hostptr PRIVATE fsz::fsz) + add_executable(szgpu_example_hostptr examples/example_hostptr.cu) + target_link_libraries(szgpu_example_hostptr PRIVATE szgpu::szgpu) endif() -if(FSZ_BUILD_TESTS) +if(SZGPU_BUILD_TESTS) enable_testing() - add_executable(fsz_test_roundtrip tests/test_roundtrip.cu) - target_link_libraries(fsz_test_roundtrip PRIVATE fsz::fsz) - add_test(NAME roundtrip COMMAND fsz_test_roundtrip) - - add_executable(fsz_test_hostptr tests/test_hostptr.cu) - target_include_directories(fsz_test_hostptr PRIVATE src) - target_link_libraries(fsz_test_hostptr PRIVATE fsz::fsz) - add_test(NAME hostptr COMMAND fsz_test_hostptr) + add_executable(szgpu_test_roundtrip tests/test_roundtrip.cu) + target_link_libraries(szgpu_test_roundtrip PRIVATE szgpu::szgpu) + add_test(NAME roundtrip COMMAND szgpu_test_roundtrip) + + add_executable(szgpu_test_hostptr tests/test_hostptr.cu) + target_include_directories(szgpu_test_hostptr PRIVATE src) + target_link_libraries(szgpu_test_hostptr PRIVATE szgpu::szgpu) + add_test(NAME hostptr COMMAND szgpu_test_hostptr) endif() include(GNUInstallDirs) include(CMakePackageConfigHelpers) # Fortran interface: built whenever a Fortran compiler is available. -option(FSZ_BUILD_FORTRAN "Build the Fortran interface" ON) -if(FSZ_BUILD_FORTRAN) +option(SZGPU_BUILD_FORTRAN "Build the Fortran interface" ON) +if(SZGPU_BUILD_FORTRAN) include(CheckLanguage) check_language(Fortran) if(CMAKE_Fortran_COMPILER) enable_language(Fortran) - add_library(fsz_fortran STATIC fortran/fsz.f90) - target_link_libraries(fsz_fortran PUBLIC fsz::fsz stdc++) - set_target_properties(fsz_fortran PROPERTIES + add_library(szgpu_fortran STATIC fortran/szgpu.f90) + target_link_libraries(szgpu_fortran PUBLIC szgpu::szgpu stdc++) + set_target_properties(szgpu_fortran PROPERTIES Fortran_MODULE_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/fortran_modules) - target_include_directories(fsz_fortran INTERFACE + target_include_directories(szgpu_fortran INTERFACE $) - if(FSZ_BUILD_EXAMPLES) - add_executable(fsz_example_fortran fortran/example_fsz.f90) - target_link_libraries(fsz_example_fortran PRIVATE fsz_fortran) + if(SZGPU_BUILD_EXAMPLES) + add_executable(szgpu_example_fortran fortran/example_szgpu.f90) + target_link_libraries(szgpu_example_fortran PRIVATE szgpu_fortran) endif() - if(FSZ_BUILD_TESTS) - add_executable(fsz_test_fortran fortran/test_fsz.f90) - target_link_libraries(fsz_test_fortran PRIVATE fsz_fortran) - add_test(NAME fortran COMMAND fsz_test_fortran) + if(SZGPU_BUILD_TESTS) + add_executable(szgpu_test_fortran fortran/test_szgpu.f90) + target_link_libraries(szgpu_test_fortran PRIVATE szgpu_fortran) + add_test(NAME fortran COMMAND szgpu_test_fortran) endif() - install(TARGETS fsz_fortran + install(TARGETS szgpu_fortran ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR}) - install(FILES ${CMAKE_CURRENT_BINARY_DIR}/fortran_modules/fsz.mod - DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/fsz) + install(FILES ${CMAKE_CURRENT_BINARY_DIR}/fortran_modules/szgpu.mod + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/szgpu) else() - message(STATUS "FSZ: no Fortran compiler found; skipping the Fortran interface") + message(STATUS "SZ-GPU: no Fortran compiler found; skipping the Fortran interface") endif() endif() -install(TARGETS fsz EXPORT fszTargets +install(TARGETS szgpu EXPORT szgpuTargets ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}) -if(FSZ_BUILD_SHARED) - install(TARGETS fsz_shared EXPORT fszTargets +if(SZGPU_BUILD_SHARED) + install(TARGETS szgpu_shared EXPORT szgpuTargets LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}) endif() -if(FSZ_BUILD_TOOLS) - install(TARGETS fsz_cli RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}) +if(SZGPU_BUILD_TOOLS) + install(TARGETS szgpu_cli RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR}) endif() install(DIRECTORY include/ DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}) -install(EXPORT fszTargets - FILE fszTargets.cmake - NAMESPACE fsz:: - DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/fsz) +install(EXPORT szgpuTargets + FILE szgpuTargets.cmake + NAMESPACE szgpu:: + DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/szgpu) -configure_package_config_file(cmake/fszConfig.cmake.in - ${CMAKE_CURRENT_BINARY_DIR}/fszConfig.cmake - INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/fsz) +configure_package_config_file(cmake/szgpuConfig.cmake.in + ${CMAKE_CURRENT_BINARY_DIR}/szgpuConfig.cmake + INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/szgpu) write_basic_package_version_file( - ${CMAKE_CURRENT_BINARY_DIR}/fszConfigVersion.cmake + ${CMAKE_CURRENT_BINARY_DIR}/szgpuConfigVersion.cmake VERSION ${PROJECT_VERSION} COMPATIBILITY SameMajorVersion) install(FILES - ${CMAKE_CURRENT_BINARY_DIR}/fszConfig.cmake - ${CMAKE_CURRENT_BINARY_DIR}/fszConfigVersion.cmake - DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/fsz) + ${CMAKE_CURRENT_BINARY_DIR}/szgpuConfig.cmake + ${CMAKE_CURRENT_BINARY_DIR}/szgpuConfigVersion.cmake + DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/szgpu) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index 2e4385e..20dbd2a 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -1,4 +1,4 @@ -# Contributing to FSZ +# Contributing to SZ-GPU Bug reports, questions, and patches are welcome. @@ -8,12 +8,12 @@ Please include: - the GPU and its compute capability (`nvidia-smi --query-gpu=name,compute_cap --format=csv`), - the CUDA toolkit version (`nvcc --version`), -- how FSZ was built (CMake or Makefile, and any options), +- how SZ-GPU was built (CMake or Makefile, and any options), - the element count, dimensions, element type, and error bound, - what you expected and what happened. -A reproducer that runs from the `fsz` tool is the most useful form, for example -`fsz -i field.f32 -eb rel 1e-3 -d 512 512 512`. If the data cannot be shared, a +A reproducer that runs from the `szgpu` tool is the most useful form, for example +`szgpu -i field.f32 -eb rel 1e-3 -d 512 512 512`. If the data cannot be shared, a synthetic array that shows the same behavior works just as well. For anything that looks like wrong output, please say whether the error bound @@ -32,12 +32,12 @@ The suites cover the C and C++ APIs (`roundtrip`, `hostptr`) and the Fortran interface (`fortran`). The Python tests run separately: ```bash -PYTHONPATH=python python3 python/tests/test_fsz.py +PYTHONPATH=python python3 python/tests/test_szgpu.py ``` ## What a change needs -FSZ is an error-bounded compressor, so correctness has a specific meaning and +SZ-GPU is an error-bounded compressor, so correctness has a specific meaning and patches are held to it: 1. **The bound holds.** Every reconstructed value must be within @@ -46,9 +46,9 @@ patches are held to it: 2. **The stream does not move unintentionally.** A change meant to affect only performance, memory, or the build must leave the compressed bytes identical. The quickest check is to compress a field before and after and compare - checksums; `fsz -z -i field.f32 -o out.fsz -eb rel 1e-3` then `md5sum`. + checksums; `szgpu -z -i field.f32 -o out.szgpu -eb rel 1e-3` then `md5sum`. 3. **Deliberate format changes are versioned.** The container carries a version - field, and an existing `.fsz` file must keep decompressing. If it cannot, + field, and an existing `.szgpu` file must keep decompressing. If it cannot, that is a major version. 4. **New behavior comes with a test.** Add the case to the suite that covers the affected interface, and prefer a case that fails before the change. @@ -61,7 +61,7 @@ patches are held to it: - Follow the surrounding code: four-space indent, and the naming already in each file. - Comments explain why something is done, not what the line does. -- Keep the public headers free of implementation detail; `src/fsz_internal.cuh` +- Keep the public headers free of implementation detail; `src/szgpu_internal.cuh` is where format constants live. ## License diff --git a/Makefile b/Makefile index 48df924..75557d0 100644 --- a/Makefile +++ b/Makefile @@ -28,13 +28,13 @@ GENCODE += -maxrregcount=96 endif INCLUDE := -Iinclude -Isrc -SRC := src/fsz.cu -HDRS := include/fsz/fsz.h include/fsz/fsz.hpp src/fsz_internal.cuh src/fsz_kernels.cuh +SRC := src/szgpu.cu +HDRS := include/szgpu/szgpu.h include/szgpu/szgpu.hpp src/szgpu_internal.cuh src/szgpu_kernels.cuh BUILD ?= build -LIB := $(BUILD)/libfsz.a -SHLIB := $(BUILD)/libfsz.so -TOOL := $(BUILD)/fsz +LIB := $(BUILD)/libszgpu.a +SHLIB := $(BUILD)/libszgpu.so +TOOL := $(BUILD)/szgpu EXAMPLES := $(BUILD)/example_roundtrip $(BUILD)/example_hostptr TESTS := $(BUILD)/test_roundtrip $(BUILD)/test_hostptr @@ -45,10 +45,10 @@ all: lib tool examples tests $(BUILD): mkdir -p $(BUILD) -$(BUILD)/fsz.o: $(SRC) $(HDRS) | $(BUILD) +$(BUILD)/szgpu.o: $(SRC) $(HDRS) | $(BUILD) $(NVCC) $(CXXFLAGS) $(GENCODE) -c $(SRC) -o $@ $(INCLUDE) -Xcompiler -fPIC -$(LIB): $(BUILD)/fsz.o +$(LIB): $(BUILD)/szgpu.o ar rcs $@ $< $(SHLIB): $(SRC) $(HDRS) | $(BUILD) @@ -57,7 +57,7 @@ $(SHLIB): $(SRC) $(HDRS) | $(BUILD) lib: $(LIB) shared: $(SHLIB) -$(TOOL): tools/fsz.cu tools/fsz_file_format.hpp $(LIB) +$(TOOL): tools/szgpu.cu tools/szgpu_file_format.hpp $(LIB) $(NVCC) $(CXXFLAGS) $(GENCODE) $< -o $@ $(INCLUDE) -Itools $(LIB) tool: $(TOOL) @@ -71,7 +71,7 @@ $(BUILD)/test_%: tests/test_%.cu $(LIB) examples: $(EXAMPLES) tests: $(TESTS) -# Fortran interface (optional): builds the fsz module plus its example and +# Fortran interface (optional): builds the szgpu module plus its example and # test against the static library. Requires gfortran or a compatible FC. ifeq ($(origin FC),default) FC := gfortran @@ -79,18 +79,18 @@ endif CUDA_HOME ?= $(shell dirname $$(dirname $$(command -v nvcc))) fortran: $(LIB) | $(BUILD) - $(FC) -O2 -J$(BUILD) -c fortran/fsz.f90 -o $(BUILD)/fsz_mod.o - $(FC) -O2 -I$(BUILD) fortran/example_fsz.f90 $(BUILD)/fsz_mod.o $(LIB) \ + $(FC) -O2 -J$(BUILD) -c fortran/szgpu.f90 -o $(BUILD)/szgpu_mod.o + $(FC) -O2 -I$(BUILD) fortran/example_szgpu.f90 $(BUILD)/szgpu_mod.o $(LIB) \ -o $(BUILD)/example_fortran -L$(CUDA_HOME)/lib64 -lcudart -lstdc++ - $(FC) -O2 -I$(BUILD) fortran/test_fsz.f90 $(BUILD)/fsz_mod.o $(LIB) \ + $(FC) -O2 -I$(BUILD) fortran/test_szgpu.f90 $(BUILD)/szgpu_mod.o $(LIB) \ -o $(BUILD)/test_fortran -L$(CUDA_HOME)/lib64 -lcudart -lstdc++ PREFIX ?= /usr/local install: $(LIB) $(TOOL) - install -d $(PREFIX)/bin $(PREFIX)/lib $(PREFIX)/include/fsz + install -d $(PREFIX)/bin $(PREFIX)/lib $(PREFIX)/include/szgpu install -m 755 $(TOOL) $(PREFIX)/bin/ install -m 644 $(LIB) $(PREFIX)/lib/ - install -m 644 include/fsz/fsz.h include/fsz/fsz.hpp $(PREFIX)/include/fsz/ + install -m 644 include/szgpu/szgpu.h include/szgpu/szgpu.hpp $(PREFIX)/include/szgpu/ clean: rm -rf $(BUILD) diff --git a/README.md b/README.md index 76c1985..7592fa1 100644 --- a/README.md +++ b/README.md @@ -1,38 +1,39 @@ -# FSZ +# SZ-GPU (FSZ) -FSZ is an extremely fast GPU error-bounded lossy compressor in the SZ family. -It stands at the family's Pareto frontier: achieving both extreme throughput -and a high compression ratio simultaneously. +SZ-GPU (FSZ) is the production GPU library for the SZ compressor family. It +is an extremely fast error-bounded lossy compressor. It stands at the +family's Pareto frontier: achieving both extreme throughput and a high +compression ratio simultaneously. It is the compressor introduced in the +SC'26 paper *FSZ: Breaking the Prediction-Throughput Trade-off in GPU Lossy +Compression* ([arXiv:2607.15413](https://arxiv.org/abs/2607.15413)), +nominated for the Best Paper Award at SC'26. -You set an error bound, absolute or relative. FSZ compresses at up to +You set an error bound, absolute or relative. SZ-GPU compresses at up to 1.4 TB/s and decompresses at up to 1.7 TB/s (measured on a B200 GPU with the tool in this repository), and every reconstructed value is guaranteed to stay within the bound. The data can be anything floating point: simulation fields, sensor and instrument captures, checkpoints, machine learning tensors. -FSZ is the compressor introduced in the SC'26 paper -*FSZ: Breaking the Prediction-Throughput Trade-off in GPU Lossy Compression* -([arXiv:2607.15413](https://arxiv.org/abs/2607.15413)), nominated for the -Best Paper Award at SC'26. - What you get: - **A guaranteed bound.** Every reconstructed value stays within the absolute or relative error bound you set, on float32 and float64 alike, and the tool checks it on every run. -- **Nothing to tune.** FSZ adapts to the data on its own: no modes to +- **Nothing to tune.** SZ-GPU adapts to the data on its own: no modes to select, no flags to sweep, and compression ratios hold up on smooth, sparse, and unstructured data alike. -- **Every interface.** The `fsz` binary plus C, C++, Python, and Fortran +- **Every interface.** The `szgpu` binary plus C, C++, Python, and Fortran APIs. Python takes numpy arrays of any shape and hands them back with the shape restored; Fortran accepts `real(4)` and `real(8)` arrays up to - rank four. `.fsz` files are self-describing and decompress with no + rank four. `.szgpu` files are self-describing and decompress with no further arguments; a bare zero-overhead bitstream mode is also available. - **BSD 3-Clause licensed.** ## News +- **2026-11**: The project is renamed SZ-GPU, the production GPU library for + the SZ compressor family. SZ-GPU 1.0.1 is released. - **2026-08**: FSZ 1.0.0 is released. - **2026-07**: The FSZ paper is nominated for the Best Paper Award at SC'26. - **2026-07**: Paper accepted at SC'26. @@ -51,12 +52,12 @@ What you get: ### CMake ```bash -git clone https://github.com/JiajunHuang1999/FSZ.git -cd FSZ +git clone https://github.com/szcompressor/SZ-GPU.git +cd SZ-GPU cmake -S . -B build cmake --build build -j ctest --test-dir build --output-on-failure # optional, needs a GPU -sudo cmake --install build # optional: libfsz, headers, fsz tool +sudo cmake --install build # optional: libszgpu, headers, szgpu tool ``` The default build covers GPU generations from Ampere on: it embeds native @@ -66,21 +67,21 @@ or newer) together with PTX, so compute-capability 8.6/8.9 devices run the PTX. Override with `-DCMAKE_CUDA_ARCHITECTURES=90` (or `80;90;100`, etc.) to build for specific generations only. -Toggles: `-DFSZ_BUILD_SHARED=OFF` (skip `libfsz.so`), `-DFSZ_BUILD_TOOLS=OFF`, -`-DFSZ_BUILD_EXAMPLES=OFF`, `-DFSZ_BUILD_TESTS=OFF`. +Toggles: `-DSZGPU_BUILD_SHARED=OFF` (skip `libszgpu.so`), `-DSZGPU_BUILD_TOOLS=OFF`, +`-DSZGPU_BUILD_EXAMPLES=OFF`, `-DSZGPU_BUILD_TESTS=OFF`. Installed projects can consume the library with CMake's `find_package`: ```cmake -find_package(fsz REQUIRED) -target_link_libraries(myapp PRIVATE fsz::fsz) +find_package(szgpu REQUIRED) +target_link_libraries(myapp PRIVATE szgpu::szgpu) ``` ### Makefile (nvcc only) ```bash make ARCH=90 # ARCH auto-detected from nvidia-smi if omitted -./build/fsz -h +./build/szgpu -h ``` ## 3. Quick start @@ -88,7 +89,7 @@ make ARCH=90 # ARCH auto-detected from nvidia-smi if omitted Compress any raw float32 array with a relative error bound of `1e-3`: ```bash -./build/fsz -i velocity.f32 -eb rel 1e-3 +./build/szgpu -i velocity.f32 -eb rel 1e-3 ``` This runs an in-memory compress + decompress round trip and reports the @@ -96,29 +97,29 @@ compression ratio, throughput, and the measured maximum error against the bound. To actually write files: ```bash -./build/fsz -z -i velocity.f32 -o velocity.fsz -eb rel 1e-3 -d 512 512 512 -./build/fsz -x -i velocity.fsz -o velocity.out.f32 +./build/szgpu -z -i velocity.f32 -o velocity.szgpu -eb rel 1e-3 -d 512 512 512 +./build/szgpu -x -i velocity.szgpu -o velocity.out.f32 ``` -The `.fsz` file is self-describing: decompression needs no dimensions, sizes, +The `.szgpu` file is self-describing: decompression needs no dimensions, sizes, or error bounds on the command line. ## 4. Command-line tool ``` -fsz -i data.f32 -eb rel 1e-3 [-d D1 [D2 [D3]]] report mode: round trip + error check, writes nothing -fsz -z -i data.f32 -o data.fsz -eb rel 1e-3 [-d ...] compress to a self-describing .fsz file -fsz -x -i data.fsz -o data.out.f32 decompress (no other arguments needed) -fsz -t f64 -i data.d64 -eb rel 1e-3 any mode on float64 data +szgpu -i data.f32 -eb rel 1e-3 [-d D1 [D2 [D3]]] report mode: round trip + error check, writes nothing +szgpu -z -i data.f32 -o data.szgpu -eb rel 1e-3 [-d ...] compress to a self-describing .szgpu file +szgpu -x -i data.szgpu -o data.out.f32 decompress (no other arguments needed) +szgpu -t f64 -i data.d64 -eb rel 1e-3 any mode on float64 data ``` | Option | Meaning | |---|---| -| `-z` / `-x` | force compress / decompress. Without either, files starting with the FSZ magic are decompressed and raw inputs get the report mode. | +| `-z` / `-x` | force compress / decompress. Without either, files starting with the SZ-GPU magic are decompressed and raw inputs get the report mode. | | `-i`, `-o` | input / output path. Without `-o`, `-z` and `-x` run fully but write nothing. | | `-eb abs V` | absolute error bound `V`. | | `-eb rel V` | relative error bound: `V * (max - min)` of the input data. | -| `-t f32\|f64` | element type of a raw input file (default `f32`). Decompression of `.fsz` files reads the type from the header instead. | +| `-t f32\|f64` | element type of a raw input file (default `f32`). Decompression of `.szgpu` files reads the type from the header instead. | | `-d D1 [D2 [D3]]` | data dimensions in C order (`D1` varies slowest, like a NumPy shape). Also accepts `-d D1xD2xD3`. Optional; the product must match the file's element count. | | `--bare` | with `-z`: write the raw bitstream without the 56-byte header. Decompressing a bare stream needs `-n` (or `-d`) and `-eb abs`. | | `-n N` | element count, for decompressing bare streams. | @@ -133,18 +134,18 @@ file errors. Device-pointer API (data already on the GPU; zero extra copies): ```c -#include +#include -fsz_workspace_t* ws = NULL; -fsz_workspace_create(&ws, n_max); /* reusable across calls */ +szgpu_workspace_t* ws = NULL; +szgpu_workspace_create(&ws, n_max); /* reusable across calls */ -fsz_compress_result_t r; -fsz_compress(d_in, d_cmp, n, eb_abs, ws, /*stream=*/0, &r); +szgpu_compress_result_t r; +szgpu_compress(d_in, d_cmp, n, eb_abs, ws, /*stream=*/0, &r); /* r.cmp_size bytes of d_cmp now hold the bitstream */ -fsz_decompress(d_out, d_cmp, n, eb_abs, &r, ws, /*stream=*/0); +szgpu_decompress(d_out, d_cmp, n, eb_abs, &r, ws, /*stream=*/0); -fsz_workspace_destroy(ws); +szgpu_workspace_destroy(ws); ``` Host-pointer API (one call, no CUDA boilerplate; the library manages device @@ -152,60 +153,60 @@ memory and transfers internally): ```c size_t cmp_size = 0; -unsigned char* cmp = malloc(fsz_max_compressed_bytes(n)); -fsz_compress_hostptr(h_in, cmp, n, eb_abs, &cmp_size); +unsigned char* cmp = malloc(szgpu_max_compressed_bytes(n)); +szgpu_compress_hostptr(h_in, cmp, n, eb_abs, &cmp_size); -fsz_decompress_hostptr(h_out, cmp, cmp_size, n, eb_abs); +szgpu_decompress_hostptr(h_out, cmp, cmp_size, n, eb_abs); ``` Supporting calls: ```c -size_t fsz_max_compressed_bytes(size_t n); /* safe d_cmp / cmp capacity */ -size_t fsz_num_tiles(size_t n); -fsz_compress_result_t fsz_make_result(size_t n, size_t cmp_size); +size_t szgpu_max_compressed_bytes(size_t n); /* safe d_cmp / cmp capacity */ +size_t szgpu_num_tiles(size_t n); +szgpu_compress_result_t szgpu_make_result(size_t n, size_t cmp_size); /* rebuild the result struct for a stored bare stream */ -const char* fsz_version(void); -const char* fsz_status_string(fsz_status_t); +const char* szgpu_version(void); +const char* szgpu_status_string(szgpu_status_t); ``` -All `fsz_*` calls return `fsz_status_t`; `FSZ_STATUS_OK == 0`. +All `szgpu_*` calls return `szgpu_status_t`; `SZGPU_STATUS_OK == 0`. Float64 data uses the `_f64` variants of the same four entry points -(`fsz_compress_f64`, `fsz_decompress_f64`, `fsz_compress_hostptr_f64`, -`fsz_decompress_hostptr_f64`), which take `double` buffers and a `double` +(`szgpu_compress_f64`, `szgpu_decompress_f64`, `szgpu_compress_hostptr_f64`, +`szgpu_decompress_hostptr_f64`), which take `double` buffers and a `double` bound. Workspaces and the sizing helpers are shared between element types. Buffer rules: - `d_in` / `d_out`: `n * sizeof(float)` device bytes. -- `d_cmp`: `fsz_max_compressed_bytes(n)` device bytes, 4-byte aligned +- `d_cmp`: `szgpu_max_compressed_bytes(n)` device bytes, 4-byte aligned (any `cudaMalloc` allocation qualifies). - Workspace: created once with a maximum capacity and reused for any - `n <= fsz_workspace_capacity(ws)`. Allocate one workspace per concurrent + `n <= szgpu_workspace_capacity(ws)`. Allocate one workspace per concurrent stream; a workspace must not be shared by two in-flight calls. Passing `NULL` instead is also allowed: the call then creates and destroys a temporary workspace internally (one-shot mode; decompression is then synchronous). Reuse a workspace when compressing repeatedly. -- `fsz_compress` synchronizes the stream before returning (it must deliver - `cmp_size`); `fsz_decompress` is fully asynchronous on the given stream. +- `szgpu_compress` synchronizes the stream before returning (it must deliver + `cmp_size`); `szgpu_decompress` is fully asynchronous on the given stream. ## 6. C++ API ```cpp -#include +#include -fsz::Workspace ws(n_max); // RAII -auto r = fsz::compress(d_in, d_cmp, n, eb_abs, ws); -fsz::decompress(d_out, d_cmp, n, eb_abs, r, ws); +szgpu::Workspace ws(n_max); // RAII +auto r = szgpu::compress(d_in, d_cmp, n, eb_abs, ws); +szgpu::decompress(d_out, d_cmp, n, eb_abs, r, ws); // host-pointer one-shots -size_t cmp_size = fsz::compress_hostptr(h_in, cmp, n, eb_abs); -fsz::decompress_hostptr(h_out, cmp, cmp_size, n, eb_abs); +size_t cmp_size = szgpu::compress_hostptr(h_in, cmp, n, eb_abs); +szgpu::decompress_hostptr(h_out, cmp, cmp_size, n, eb_abs); ``` -Errors are reported as `fsz::Error` (derives from `std::runtime_error`). +Errors are reported as `szgpu::Error` (derives from `std::runtime_error`). All of these functions also accept `double*` buffers through overloads. ## 7. Python API @@ -216,19 +217,19 @@ pip install ./python The install compiles and bundles the shared library, so nothing needs to be built first; it only needs a CUDA toolkit on the machine. Set -`FSZ_CUDA_ARCHITECTURES` (for example `80;90;100`) to target specific GPU +`SZGPU_CUDA_ARCHITECTURES` (for example `80;90;100`) to target specific GPU generations. If no toolkit is present the install still succeeds and the -module falls back to `FSZ_LIB`, a repository build directory, or the system +module falls back to `SZGPU_LIB`, a repository build directory, or the system search path, naming all of them if none works. ```python import numpy as np -import fsz +import szgpu arr = np.fromfile("velocity.f32", dtype=np.float32).reshape(512, 512, 512) -blob = fsz.compress(arr, eb=1e-3, mode="rel") # bytes (self-describing) -out = fsz.decompress(blob) # float32 ndarray, shape restored +blob = szgpu.compress(arr, eb=1e-3, mode="rel") # bytes (self-describing) +out = szgpu.decompress(blob) # float32 ndarray, shape restored print(arr.shape == out.shape, np.abs(arr - out).max()) ``` @@ -236,35 +237,35 @@ print(arr.shape == out.shape, np.abs(arr - out).max()) Arrays of dtype float64 compress through the float64 path and decompress back as float64; the container records which type was written. -File helpers mirror the CLI: `fsz.compress_file(src, dst, eb, mode="rel", -dims=None, dtype=np.float32)` and `fsz.decompress_file(src, dst=None)`. The -package is a thin `ctypes` wrapper over `libfsz.so`: point it at a custom -build with the `FSZ_LIB` environment variable if the library is not +File helpers mirror the CLI: `szgpu.compress_file(src, dst, eb, mode="rel", +dims=None, dtype=np.float32)` and `szgpu.decompress_file(src, dst=None)`. The +package is a thin `ctypes` wrapper over `libszgpu.so`: point it at a custom +build with the `SZGPU_LIB` environment variable if the library is not installed system-wide. ## 8. Fortran API -The `fsz` module wraps the host-pointer calls with generic interfaces for +The `szgpu` module wraps the host-pointer calls with generic interfaces for `real(4)` and `real(8)` data; arrays of rank one to four are passed directly: ```fortran -use fsz +use szgpu real(4) :: field(512, 512, 512) integer(c_int8_t), allocatable :: cmp(:) integer(c_size_t) :: n, cmp_size integer(c_int) :: ierr n = int(size(field), c_size_t) -allocate(cmp(fsz_max_compressed_bytes(n))) +allocate(cmp(szgpu_max_compressed_bytes(n))) -call fsz_compress(field, n, 1.0e-3, cmp, cmp_size, ierr) ! absolute bound -call fsz_decompress(field, cmp, cmp_size, n, 1.0e-3, ierr) +call szgpu_compress(field, n, 1.0e-3, cmp, cmp_size, ierr) ! absolute bound +call szgpu_decompress(field, cmp, cmp_size, n, 1.0e-3, ierr) ``` -`ierr` is `FSZ_OK` (0) on success and `fsz_status_message(ierr)` describes +`ierr` is `SZGPU_OK` (0) on success and `szgpu_status_message(ierr)` describes a failure. CMake builds the module automatically when a Fortran compiler -is present (the `fsz_fortran` library and `fsz.mod`); with the Makefile, +is present (the `szgpu_fortran` library and `szgpu.mod`); with the Makefile, run `make fortran`. For data that is already on the GPU, pass device addresses as `type(c_ptr)` @@ -274,13 +275,13 @@ and `target data use_device_ptr`: ```fortran type(c_ptr) :: ws -type(fsz_result) :: res +type(szgpu_result) :: res -call fsz_workspace_create(ws, n, ierr) +call szgpu_workspace_create(ws, n, ierr) !$acc host_data use_device(field, cmp_buffer) -call fsz_compress_device(c_loc(field), c_loc(cmp_buffer), n, 1.0e-3, ws, res, ierr) +call szgpu_compress_device(c_loc(field), c_loc(cmp_buffer), n, 1.0e-3, ws, res, ierr) !$acc end host_data -call fsz_workspace_destroy(ws) +call szgpu_workspace_destroy(ws) ``` `res%cmp_size` gives the stream length. The element type follows the kind of @@ -319,7 +320,7 @@ single-field calls: | NYX/temperature | 537 MB | 1036 | 1303 | | SCALE-LETKF/T | 564 MB | 972 | 1226 | -Float64 fields, measured with the `fsz` tool on the same B200 at rel `1e-3` +Float64 fields, measured with the `szgpu` tool on the same B200 at rel `1e-3` (average of 5 timed runs after 3 warmups): | Field | Size | CR | Compress GB/s | Decompress GB/s | @@ -329,14 +330,15 @@ Float64 fields, measured with the `fsz` tool on the same B200 at rel `1e-3` ## 10. Citation +If you use SZ-GPU, please cite the FSZ paper: + ```bibtex @inproceedings{huang2026fsz, title = {{FSZ}: Breaking the Prediction-Throughput Trade-off in {GPU} Lossy Compression}, author = {Huang, Jiajun}, booktitle = {Proceedings of the International Conference for High Performance - Computing, Networking, Storage and Analysis (SC '26)}, - year = {2026}, - note = {To appear. arXiv:2607.15413} + Computing, Networking, Storage and Analysis}, + year = {2026} } ``` diff --git a/VERSION b/VERSION index 3eefcb9..7dea76e 100644 --- a/VERSION +++ b/VERSION @@ -1 +1 @@ -1.0.0 +1.0.1 diff --git a/cmake/fszConfig.cmake.in b/cmake/fszConfig.cmake.in deleted file mode 100644 index 1659257..0000000 --- a/cmake/fszConfig.cmake.in +++ /dev/null @@ -1,8 +0,0 @@ -@PACKAGE_INIT@ - -include(CMakeFindDependencyMacro) -find_dependency(CUDAToolkit) - -include("${CMAKE_CURRENT_LIST_DIR}/fszTargets.cmake") - -check_required_components(fsz) diff --git a/cmake/szgpuConfig.cmake.in b/cmake/szgpuConfig.cmake.in new file mode 100644 index 0000000..63daef7 --- /dev/null +++ b/cmake/szgpuConfig.cmake.in @@ -0,0 +1,8 @@ +@PACKAGE_INIT@ + +include(CMakeFindDependencyMacro) +find_dependency(CUDAToolkit) + +include("${CMAKE_CURRENT_LIST_DIR}/szgpuTargets.cmake") + +check_required_components(szgpu) diff --git a/examples/example_hostptr.cu b/examples/example_hostptr.cu index bdfd6ce..9d98cf8 100644 --- a/examples/example_hostptr.cu +++ b/examples/example_hostptr.cu @@ -1,5 +1,5 @@ -#include "fsz/fsz.h" +#include "szgpu/szgpu.h" #include #include @@ -18,18 +18,18 @@ int main() { * std::cos(0.04f * (float)y) + 0.1f * std::sin(0.03f * (float)z); - std::vector cmp(fsz_max_compressed_bytes(n)); + std::vector cmp(szgpu_max_compressed_bytes(n)); std::size_t cmp_size = 0; - fsz_status_t s = fsz_compress_hostptr(in.data(), cmp.data(), n, eb, &cmp_size); - if (s != FSZ_STATUS_OK) { - std::printf("compress failed: %s\n", fsz_status_string(s)); + szgpu_status_t s = szgpu_compress_hostptr(in.data(), cmp.data(), n, eb, &cmp_size); + if (s != SZGPU_STATUS_OK) { + std::printf("compress failed: %s\n", szgpu_status_string(s)); return 1; } std::vector out(n); - s = fsz_decompress_hostptr(out.data(), cmp.data(), cmp_size, n, eb); - if (s != FSZ_STATUS_OK) { - std::printf("decompress failed: %s\n", fsz_status_string(s)); + s = szgpu_decompress_hostptr(out.data(), cmp.data(), cmp_size, n, eb); + if (s != SZGPU_STATUS_OK) { + std::printf("decompress failed: %s\n", szgpu_status_string(s)); return 1; } @@ -39,7 +39,7 @@ int main() { if (e > max_err) max_err = e; } - std::printf("FSZ %s N=%zu abs_eb=%g\n", fsz_version(), n, (double)eb); + std::printf("SZ-GPU %s N=%zu abs_eb=%g\n", szgpu_version(), n, (double)eb); std::printf(" cmp_size = %zu B (from %zu B)\n", cmp_size, n * sizeof(float)); std::printf(" CR = %.3f\n", (double)(n * sizeof(float)) / (double)cmp_size); std::printf(" max_err = %.3e (bound %.3e)\n", max_err, (double)eb); diff --git a/examples/example_roundtrip.cu b/examples/example_roundtrip.cu index 2fdd3aa..9518aaa 100644 --- a/examples/example_roundtrip.cu +++ b/examples/example_roundtrip.cu @@ -1,4 +1,4 @@ -#include "fsz/fsz.hpp" +#include "szgpu/szgpu.hpp" #include #include @@ -17,7 +17,7 @@ int main(int argc, char** argv) { std::size_t n = (argc > 1) ? std::strtoull(argv[1], nullptr, 10) : (1u << 22); float eb = (argc > 2) ? std::strtof(argv[2], nullptr) : 1e-3f; - std::printf("FSZ %s N=%zu abs_eb=%g\n", fsz::version(), n, eb); + std::printf("SZ-GPU %s N=%zu abs_eb=%g\n", szgpu::version(), n, eb); std::vector h_in(n); for (std::size_t i = 0; i < n; ++i) @@ -28,14 +28,14 @@ int main(int argc, char** argv) { unsigned char* d_cmp = nullptr; check_cuda(cudaMalloc(&d_in, n * sizeof(float)), "alloc d_in"); check_cuda(cudaMalloc(&d_out, n * sizeof(float)), "alloc d_out"); - check_cuda(cudaMalloc(&d_cmp, fsz::max_compressed_bytes(n)), "alloc d_cmp"); + check_cuda(cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(n)), "alloc d_cmp"); check_cuda(cudaMemcpy(d_in, h_in.data(), n * sizeof(float), cudaMemcpyHostToDevice), "memcpy h2d"); - fsz::Workspace ws(n); + szgpu::Workspace ws(n); - auto cresult = fsz::compress(d_in, d_cmp, n, eb, ws); - fsz::decompress(d_out, d_cmp, n, eb, cresult, ws); + auto cresult = szgpu::compress(d_in, d_cmp, n, eb, ws); + szgpu::decompress(d_out, d_cmp, n, eb, cresult, ws); check_cuda(cudaDeviceSynchronize(), "synchronize"); std::vector h_out(n); diff --git a/fortran/example_fsz.f90 b/fortran/example_szgpu.f90 similarity index 67% rename from fortran/example_fsz.f90 rename to fortran/example_szgpu.f90 index 5aafdc1..a8ffca2 100644 --- a/fortran/example_fsz.f90 +++ b/fortran/example_szgpu.f90 @@ -1,6 +1,6 @@ -program example_fsz +program example_szgpu use iso_c_binding, only: c_size_t, c_int, c_int8_t - use fsz + use szgpu implicit none integer, parameter :: nx = 64, ny = 64, nz = 64 @@ -22,22 +22,22 @@ program example_fsz n = int(nx, c_size_t) * ny * nz eb = 1.0e-3 - allocate(cmp(fsz_max_compressed_bytes(n))) + allocate(cmp(szgpu_max_compressed_bytes(n))) - call fsz_compress(field, n, eb, cmp, cmp_size, ierr) - if (ierr /= FSZ_OK) then - print '(a)', 'compression failed: ' // fsz_status_message(ierr) + call szgpu_compress(field, n, eb, cmp, cmp_size, ierr) + if (ierr /= SZGPU_OK) then + print '(a)', 'compression failed: ' // szgpu_status_message(ierr) error stop 1 end if - call fsz_decompress(recon, cmp, cmp_size, n, eb, ierr) - if (ierr /= FSZ_OK) then - print '(a)', 'decompression failed: ' // fsz_status_message(ierr) + call szgpu_decompress(recon, cmp, cmp_size, n, eb, ierr) + if (ierr /= SZGPU_OK) then + print '(a)', 'decompression failed: ' // szgpu_status_message(ierr) error stop 1 end if max_err = maxval(abs(field - recon)) - print '(a,a)', 'FSZ : ', fsz_version() + print '(a,a)', 'SZ-GPU : ', szgpu_version() print '(a,f0.3)', 'ratio : ', real(n) * 4.0 / real(cmp_size) print '(a,es10.3,a,es10.3)', 'max error : ', max_err, ' bound ', eb if (max_err <= 1.01 * eb) then @@ -46,4 +46,4 @@ program example_fsz print '(a)', 'FAIL' error stop 1 end if -end program example_fsz +end program example_szgpu diff --git a/fortran/fsz.f90 b/fortran/szgpu.f90 similarity index 69% rename from fortran/fsz.f90 rename to fortran/szgpu.f90 index 9d3ad76..8bafa6d 100644 --- a/fortran/fsz.f90 +++ b/fortran/szgpu.f90 @@ -1,12 +1,12 @@ -! Fortran interface for the FSZ GPU error-bounded lossy compressor. +! Fortran interface for the SZ-GPU error-bounded lossy compressor. ! ! Host arrays. Generic interfaces for real(4) and real(8) arrays of rank one ! to four; pass the total element count. All device work happens inside the ! library, so callers never touch CUDA. ! -! use fsz -! call fsz_compress(field, n, 1.0e-3, cmp, cmp_size, ierr) -! call fsz_decompress(field, cmp, cmp_size, n, 1.0e-3, ierr) +! use szgpu +! call szgpu_compress(field, n, 1.0e-3, cmp, cmp_size, ierr) +! call szgpu_decompress(field, cmp, cmp_size, n, 1.0e-3, ierr) ! ! Data already on the GPU. Pass device addresses as type(c_ptr), which is what ! CUDA Fortran gives for a device array through c_devloc, and what OpenACC and @@ -14,64 +14,64 @@ ! element type is selected by the kind of the error bound. ! ! type(c_ptr) :: ws -! type(fsz_result) :: res -! call fsz_workspace_create(ws, n, ierr) -! call fsz_compress_device(d_in, d_cmp, n, 1.0e-3, ws, res, ierr) -! call fsz_decompress_device(d_out, d_cmp, n, 1.0e-3, res, ws, ierr) -! call fsz_workspace_destroy(ws) +! type(szgpu_result) :: res +! call szgpu_workspace_create(ws, n, ierr) +! call szgpu_compress_device(d_in, d_cmp, n, 1.0e-3, ws, res, ierr) +! call szgpu_decompress_device(d_out, d_cmp, n, 1.0e-3, res, ws, ierr) +! call szgpu_workspace_destroy(ws) ! ! A workspace may be c_null_ptr, in which case the call allocates and frees a ! temporary one; reuse a workspace when compressing repeatedly. The optional ! stream argument is a CUDA stream handle, defaulting to the default stream. ! -! ierr is FSZ_OK (0) on success; fsz_status_message(ierr) describes a -! failure. cmp must provide at least fsz_max_compressed_bytes(n) bytes. +! ierr is SZGPU_OK (0) on success; szgpu_status_message(ierr) describes a +! failure. cmp must provide at least szgpu_max_compressed_bytes(n) bytes. -module fsz +module szgpu use iso_c_binding, only: c_float, c_double, c_size_t, c_int, c_int8_t, & c_char, c_ptr, c_loc, c_f_pointer, c_associated, & c_null_char, c_null_ptr implicit none private - public :: fsz_compress, fsz_decompress - public :: fsz_compress_device, fsz_decompress_device - public :: fsz_workspace_create, fsz_workspace_destroy, fsz_workspace_capacity - public :: fsz_make_result - public :: fsz_max_compressed_bytes, fsz_num_tiles - public :: fsz_version, fsz_status_message - public :: fsz_result, FSZ_OK + public :: szgpu_compress, szgpu_decompress + public :: szgpu_compress_device, szgpu_decompress_device + public :: szgpu_workspace_create, szgpu_workspace_destroy, szgpu_workspace_capacity + public :: szgpu_make_result + public :: szgpu_max_compressed_bytes, szgpu_num_tiles + public :: szgpu_version, szgpu_status_message + public :: szgpu_result, SZGPU_OK - integer(c_int), parameter :: FSZ_OK = 0_c_int + integer(c_int), parameter :: SZGPU_OK = 0_c_int - type, bind(c) :: fsz_result + type, bind(c) :: szgpu_result integer(c_size_t) :: cmp_size = 0_c_size_t integer(c_size_t) :: num_tiles = 0_c_size_t integer(c_size_t) :: data_offset = 0_c_size_t - end type fsz_result + end type szgpu_result interface - function c_max_bytes(n) bind(c, name='fsz_max_compressed_bytes') result(r) + function c_max_bytes(n) bind(c, name='szgpu_max_compressed_bytes') result(r) import :: c_size_t implicit none integer(c_size_t), value :: n integer(c_size_t) :: r end function c_max_bytes - function c_num_tiles(n) bind(c, name='fsz_num_tiles') result(r) + function c_num_tiles(n) bind(c, name='szgpu_num_tiles') result(r) import :: c_size_t implicit none integer(c_size_t), value :: n integer(c_size_t) :: r end function c_num_tiles - function c_version() bind(c, name='fsz_version') result(p) + function c_version() bind(c, name='szgpu_version') result(p) import :: c_ptr implicit none type(c_ptr) :: p end function c_version - function c_status_string(s) bind(c, name='fsz_status_string') result(p) + function c_status_string(s) bind(c, name='szgpu_status_string') result(p) import :: c_int, c_ptr implicit none integer(c_int), value :: s @@ -79,7 +79,7 @@ function c_status_string(s) bind(c, name='fsz_status_string') result(p) end function c_status_string function c_compress_r4(values, cmp, n, eb, out_size) & - bind(c, name='fsz_compress_hostptr') result(st) + bind(c, name='szgpu_compress_hostptr') result(st) import :: c_float, c_size_t, c_int, c_int8_t, c_ptr implicit none type(c_ptr), value :: values @@ -91,7 +91,7 @@ function c_compress_r4(values, cmp, n, eb, out_size) & end function c_compress_r4 function c_decompress_r4(values, cmp, cmp_size, n, eb) & - bind(c, name='fsz_decompress_hostptr') result(st) + bind(c, name='szgpu_decompress_hostptr') result(st) import :: c_float, c_size_t, c_int, c_int8_t, c_ptr implicit none type(c_ptr), value :: values @@ -103,7 +103,7 @@ function c_decompress_r4(values, cmp, cmp_size, n, eb) & end function c_decompress_r4 function c_compress_r8(values, cmp, n, eb, out_size) & - bind(c, name='fsz_compress_hostptr_f64') result(st) + bind(c, name='szgpu_compress_hostptr_f64') result(st) import :: c_double, c_size_t, c_int, c_int8_t, c_ptr implicit none type(c_ptr), value :: values @@ -115,7 +115,7 @@ function c_compress_r8(values, cmp, n, eb, out_size) & end function c_compress_r8 function c_decompress_r8(values, cmp, cmp_size, n, eb) & - bind(c, name='fsz_decompress_hostptr_f64') result(st) + bind(c, name='szgpu_decompress_hostptr_f64') result(st) import :: c_double, c_size_t, c_int, c_int8_t, c_ptr implicit none type(c_ptr), value :: values @@ -126,7 +126,7 @@ function c_decompress_r8(values, cmp, cmp_size, n, eb) & integer(c_int) :: st end function c_decompress_r8 - function c_ws_create(ws, max_n) bind(c, name='fsz_workspace_create') result(st) + function c_ws_create(ws, max_n) bind(c, name='szgpu_workspace_create') result(st) import :: c_ptr, c_size_t, c_int implicit none type(c_ptr), intent(out) :: ws @@ -134,32 +134,32 @@ function c_ws_create(ws, max_n) bind(c, name='fsz_workspace_create') result(st) integer(c_int) :: st end function c_ws_create - subroutine c_ws_destroy(ws) bind(c, name='fsz_workspace_destroy') + subroutine c_ws_destroy(ws) bind(c, name='szgpu_workspace_destroy') import :: c_ptr implicit none type(c_ptr), value :: ws end subroutine c_ws_destroy - function c_ws_capacity(ws) bind(c, name='fsz_workspace_capacity') result(r) + function c_ws_capacity(ws) bind(c, name='szgpu_workspace_capacity') result(r) import :: c_ptr, c_size_t implicit none type(c_ptr), value :: ws integer(c_size_t) :: r end function c_ws_capacity - function c_make_result(n, cmp_size) bind(c, name='fsz_make_result') result(r) + function c_make_result(n, cmp_size) bind(c, name='szgpu_make_result') result(r) import :: c_size_t - import :: fsz_result + import :: szgpu_result implicit none integer(c_size_t), value :: n integer(c_size_t), value :: cmp_size - type(fsz_result) :: r + type(szgpu_result) :: r end function c_make_result function c_dev_compress_r4(d_in, d_cmp, n, eb, ws, stream, res) & - bind(c, name='fsz_compress') result(st) + bind(c, name='szgpu_compress') result(st) import :: c_ptr, c_size_t, c_float, c_int - import :: fsz_result + import :: szgpu_result implicit none type(c_ptr), value :: d_in type(c_ptr), value :: d_cmp @@ -167,14 +167,14 @@ function c_dev_compress_r4(d_in, d_cmp, n, eb, ws, stream, res) & real(c_float), value :: eb type(c_ptr), value :: ws type(c_ptr), value :: stream - type(fsz_result), intent(out) :: res + type(szgpu_result), intent(out) :: res integer(c_int) :: st end function c_dev_compress_r4 function c_dev_compress_r8(d_in, d_cmp, n, eb, ws, stream, res) & - bind(c, name='fsz_compress_f64') result(st) + bind(c, name='szgpu_compress_f64') result(st) import :: c_ptr, c_size_t, c_double, c_int - import :: fsz_result + import :: szgpu_result implicit none type(c_ptr), value :: d_in type(c_ptr), value :: d_cmp @@ -182,66 +182,66 @@ function c_dev_compress_r8(d_in, d_cmp, n, eb, ws, stream, res) & real(c_double), value :: eb type(c_ptr), value :: ws type(c_ptr), value :: stream - type(fsz_result), intent(out) :: res + type(szgpu_result), intent(out) :: res integer(c_int) :: st end function c_dev_compress_r8 function c_dev_decompress_r4(d_out, d_cmp, n, eb, res, ws, stream) & - bind(c, name='fsz_decompress') result(st) + bind(c, name='szgpu_decompress') result(st) import :: c_ptr, c_size_t, c_float, c_int - import :: fsz_result + import :: szgpu_result implicit none type(c_ptr), value :: d_out type(c_ptr), value :: d_cmp integer(c_size_t), value :: n real(c_float), value :: eb - type(fsz_result), intent(in) :: res + type(szgpu_result), intent(in) :: res type(c_ptr), value :: ws type(c_ptr), value :: stream integer(c_int) :: st end function c_dev_decompress_r4 function c_dev_decompress_r8(d_out, d_cmp, n, eb, res, ws, stream) & - bind(c, name='fsz_decompress_f64') result(st) + bind(c, name='szgpu_decompress_f64') result(st) import :: c_ptr, c_size_t, c_double, c_int - import :: fsz_result + import :: szgpu_result implicit none type(c_ptr), value :: d_out type(c_ptr), value :: d_cmp integer(c_size_t), value :: n real(c_double), value :: eb - type(fsz_result), intent(in) :: res + type(szgpu_result), intent(in) :: res type(c_ptr), value :: ws type(c_ptr), value :: stream integer(c_int) :: st end function c_dev_decompress_r8 end interface - interface fsz_compress_device - module procedure fsz_compress_device_r4, fsz_compress_device_r8 - end interface fsz_compress_device + interface szgpu_compress_device + module procedure szgpu_compress_device_r4, szgpu_compress_device_r8 + end interface szgpu_compress_device - interface fsz_decompress_device - module procedure fsz_decompress_device_r4, fsz_decompress_device_r8 - end interface fsz_decompress_device + interface szgpu_decompress_device + module procedure szgpu_decompress_device_r4, szgpu_decompress_device_r8 + end interface szgpu_decompress_device - interface fsz_compress - module procedure fsz_compress_r4_1, fsz_compress_r4_2, & - fsz_compress_r4_3, fsz_compress_r4_4, & - fsz_compress_r8_1, fsz_compress_r8_2, & - fsz_compress_r8_3, fsz_compress_r8_4 - end interface fsz_compress + interface szgpu_compress + module procedure szgpu_compress_r4_1, szgpu_compress_r4_2, & + szgpu_compress_r4_3, szgpu_compress_r4_4, & + szgpu_compress_r8_1, szgpu_compress_r8_2, & + szgpu_compress_r8_3, szgpu_compress_r8_4 + end interface szgpu_compress - interface fsz_decompress - module procedure fsz_decompress_r4_1, fsz_decompress_r4_2, & - fsz_decompress_r4_3, fsz_decompress_r4_4, & - fsz_decompress_r8_1, fsz_decompress_r8_2, & - fsz_decompress_r8_3, fsz_decompress_r8_4 - end interface fsz_decompress + interface szgpu_decompress + module procedure szgpu_decompress_r4_1, szgpu_decompress_r4_2, & + szgpu_decompress_r4_3, szgpu_decompress_r4_4, & + szgpu_decompress_r8_1, szgpu_decompress_r8_2, & + szgpu_decompress_r8_3, szgpu_decompress_r8_4 + end interface szgpu_decompress contains - subroutine fsz_compress_r4_1(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r4_1(values, n, eb_abs, cmp, cmp_size, ierr) real(c_float), intent(in), contiguous, target :: values(:) integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs @@ -249,9 +249,9 @@ subroutine fsz_compress_r4_1(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r4(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r4_1 + end subroutine szgpu_compress_r4_1 - subroutine fsz_compress_r4_2(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r4_2(values, n, eb_abs, cmp, cmp_size, ierr) real(c_float), intent(in), contiguous, target :: values(:, :) integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs @@ -259,9 +259,9 @@ subroutine fsz_compress_r4_2(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r4(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r4_2 + end subroutine szgpu_compress_r4_2 - subroutine fsz_compress_r4_3(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r4_3(values, n, eb_abs, cmp, cmp_size, ierr) real(c_float), intent(in), contiguous, target :: values(:, :, :) integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs @@ -269,9 +269,9 @@ subroutine fsz_compress_r4_3(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r4(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r4_3 + end subroutine szgpu_compress_r4_3 - subroutine fsz_compress_r4_4(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r4_4(values, n, eb_abs, cmp, cmp_size, ierr) real(c_float), intent(in), contiguous, target :: values(:, :, :, :) integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs @@ -279,9 +279,9 @@ subroutine fsz_compress_r4_4(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r4(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r4_4 + end subroutine szgpu_compress_r4_4 - subroutine fsz_decompress_r4_1(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r4_1(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(out), contiguous, target :: values(:) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -289,9 +289,9 @@ subroutine fsz_decompress_r4_1(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r4(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r4_1 + end subroutine szgpu_decompress_r4_1 - subroutine fsz_decompress_r4_2(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r4_2(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(out), contiguous, target :: values(:, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -299,9 +299,9 @@ subroutine fsz_decompress_r4_2(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r4(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r4_2 + end subroutine szgpu_decompress_r4_2 - subroutine fsz_decompress_r4_3(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r4_3(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(out), contiguous, target :: values(:, :, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -309,9 +309,9 @@ subroutine fsz_decompress_r4_3(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r4(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r4_3 + end subroutine szgpu_decompress_r4_3 - subroutine fsz_decompress_r4_4(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r4_4(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(out), contiguous, target :: values(:, :, :, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -319,9 +319,9 @@ subroutine fsz_decompress_r4_4(values, cmp, cmp_size, n, eb_abs, ierr) real(c_float), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r4(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r4_4 + end subroutine szgpu_decompress_r4_4 - subroutine fsz_compress_r8_1(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r8_1(values, n, eb_abs, cmp, cmp_size, ierr) real(c_double), intent(in), contiguous, target :: values(:) integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs @@ -329,9 +329,9 @@ subroutine fsz_compress_r8_1(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r8(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r8_1 + end subroutine szgpu_compress_r8_1 - subroutine fsz_compress_r8_2(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r8_2(values, n, eb_abs, cmp, cmp_size, ierr) real(c_double), intent(in), contiguous, target :: values(:, :) integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs @@ -339,9 +339,9 @@ subroutine fsz_compress_r8_2(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r8(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r8_2 + end subroutine szgpu_compress_r8_2 - subroutine fsz_compress_r8_3(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r8_3(values, n, eb_abs, cmp, cmp_size, ierr) real(c_double), intent(in), contiguous, target :: values(:, :, :) integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs @@ -349,9 +349,9 @@ subroutine fsz_compress_r8_3(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r8(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r8_3 + end subroutine szgpu_compress_r8_3 - subroutine fsz_compress_r8_4(values, n, eb_abs, cmp, cmp_size, ierr) + subroutine szgpu_compress_r8_4(values, n, eb_abs, cmp, cmp_size, ierr) real(c_double), intent(in), contiguous, target :: values(:, :, :, :) integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs @@ -359,9 +359,9 @@ subroutine fsz_compress_r8_4(values, n, eb_abs, cmp, cmp_size, ierr) integer(c_size_t), intent(out) :: cmp_size integer(c_int), intent(out) :: ierr ierr = c_compress_r8(c_loc(values), cmp, n, eb_abs, cmp_size) - end subroutine fsz_compress_r8_4 + end subroutine szgpu_compress_r8_4 - subroutine fsz_decompress_r8_1(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r8_1(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(out), contiguous, target :: values(:) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -369,9 +369,9 @@ subroutine fsz_decompress_r8_1(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r8(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r8_1 + end subroutine szgpu_decompress_r8_1 - subroutine fsz_decompress_r8_2(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r8_2(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(out), contiguous, target :: values(:, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -379,9 +379,9 @@ subroutine fsz_decompress_r8_2(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r8(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r8_2 + end subroutine szgpu_decompress_r8_2 - subroutine fsz_decompress_r8_3(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r8_3(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(out), contiguous, target :: values(:, :, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -389,9 +389,9 @@ subroutine fsz_decompress_r8_3(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r8(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r8_3 + end subroutine szgpu_decompress_r8_3 - subroutine fsz_decompress_r8_4(values, cmp, cmp_size, n, eb_abs, ierr) + subroutine szgpu_decompress_r8_4(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(out), contiguous, target :: values(:, :, :, :) integer(c_int8_t), intent(in) :: cmp(*) integer(c_size_t), intent(in) :: cmp_size @@ -399,72 +399,72 @@ subroutine fsz_decompress_r8_4(values, cmp, cmp_size, n, eb_abs, ierr) real(c_double), intent(in) :: eb_abs integer(c_int), intent(out) :: ierr ierr = c_decompress_r8(c_loc(values), cmp, cmp_size, n, eb_abs) - end subroutine fsz_decompress_r8_4 + end subroutine szgpu_decompress_r8_4 - subroutine fsz_workspace_create(ws, max_n, ierr) + subroutine szgpu_workspace_create(ws, max_n, ierr) type(c_ptr), intent(out) :: ws integer(c_size_t), intent(in) :: max_n integer(c_int), intent(out) :: ierr ierr = c_ws_create(ws, max_n) - end subroutine fsz_workspace_create + end subroutine szgpu_workspace_create - subroutine fsz_workspace_destroy(ws) + subroutine szgpu_workspace_destroy(ws) type(c_ptr), intent(inout) :: ws call c_ws_destroy(ws) ws = c_null_ptr - end subroutine fsz_workspace_destroy + end subroutine szgpu_workspace_destroy - function fsz_workspace_capacity(ws) result(r) + function szgpu_workspace_capacity(ws) result(r) type(c_ptr), intent(in) :: ws integer(c_size_t) :: r r = c_ws_capacity(ws) - end function fsz_workspace_capacity + end function szgpu_workspace_capacity ! Rebuilds the descriptor of a stream known only by its element count and ! byte size, for decompressing data read back from storage. - function fsz_make_result(n, cmp_size) result(r) + function szgpu_make_result(n, cmp_size) result(r) integer(c_size_t), intent(in) :: n integer(c_size_t), intent(in) :: cmp_size - type(fsz_result) :: r + type(szgpu_result) :: r r = c_make_result(n, cmp_size) - end function fsz_make_result + end function szgpu_make_result - subroutine fsz_compress_device_r4(d_in, d_cmp, n, eb_abs, ws, res, ierr, stream) + subroutine szgpu_compress_device_r4(d_in, d_cmp, n, eb_abs, ws, res, ierr, stream) type(c_ptr), intent(in) :: d_in type(c_ptr), intent(in) :: d_cmp integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs type(c_ptr), intent(in) :: ws - type(fsz_result), intent(out) :: res + type(szgpu_result), intent(out) :: res integer(c_int), intent(out) :: ierr type(c_ptr), intent(in), optional :: stream type(c_ptr) :: s s = c_null_ptr if (present(stream)) s = stream ierr = c_dev_compress_r4(d_in, d_cmp, n, eb_abs, ws, s, res) - end subroutine fsz_compress_device_r4 + end subroutine szgpu_compress_device_r4 - subroutine fsz_compress_device_r8(d_in, d_cmp, n, eb_abs, ws, res, ierr, stream) + subroutine szgpu_compress_device_r8(d_in, d_cmp, n, eb_abs, ws, res, ierr, stream) type(c_ptr), intent(in) :: d_in type(c_ptr), intent(in) :: d_cmp integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs type(c_ptr), intent(in) :: ws - type(fsz_result), intent(out) :: res + type(szgpu_result), intent(out) :: res integer(c_int), intent(out) :: ierr type(c_ptr), intent(in), optional :: stream type(c_ptr) :: s s = c_null_ptr if (present(stream)) s = stream ierr = c_dev_compress_r8(d_in, d_cmp, n, eb_abs, ws, s, res) - end subroutine fsz_compress_device_r8 + end subroutine szgpu_compress_device_r8 - subroutine fsz_decompress_device_r4(d_out, d_cmp, n, eb_abs, res, ws, ierr, stream) + subroutine szgpu_decompress_device_r4(d_out, d_cmp, n, eb_abs, res, ws, ierr, stream) type(c_ptr), intent(in) :: d_out type(c_ptr), intent(in) :: d_cmp integer(c_size_t), intent(in) :: n real(c_float), intent(in) :: eb_abs - type(fsz_result), intent(in) :: res + type(szgpu_result), intent(in) :: res type(c_ptr), intent(in) :: ws integer(c_int), intent(out) :: ierr type(c_ptr), intent(in), optional :: stream @@ -472,14 +472,14 @@ subroutine fsz_decompress_device_r4(d_out, d_cmp, n, eb_abs, res, ws, ierr, stre s = c_null_ptr if (present(stream)) s = stream ierr = c_dev_decompress_r4(d_out, d_cmp, n, eb_abs, res, ws, s) - end subroutine fsz_decompress_device_r4 + end subroutine szgpu_decompress_device_r4 - subroutine fsz_decompress_device_r8(d_out, d_cmp, n, eb_abs, res, ws, ierr, stream) + subroutine szgpu_decompress_device_r8(d_out, d_cmp, n, eb_abs, res, ws, ierr, stream) type(c_ptr), intent(in) :: d_out type(c_ptr), intent(in) :: d_cmp integer(c_size_t), intent(in) :: n real(c_double), intent(in) :: eb_abs - type(fsz_result), intent(in) :: res + type(szgpu_result), intent(in) :: res type(c_ptr), intent(in) :: ws integer(c_int), intent(out) :: ierr type(c_ptr), intent(in), optional :: stream @@ -487,30 +487,30 @@ subroutine fsz_decompress_device_r8(d_out, d_cmp, n, eb_abs, res, ws, ierr, stre s = c_null_ptr if (present(stream)) s = stream ierr = c_dev_decompress_r8(d_out, d_cmp, n, eb_abs, res, ws, s) - end subroutine fsz_decompress_device_r8 + end subroutine szgpu_decompress_device_r8 - function fsz_max_compressed_bytes(n) result(r) + function szgpu_max_compressed_bytes(n) result(r) integer(c_size_t), intent(in) :: n integer(c_size_t) :: r r = c_max_bytes(n) - end function fsz_max_compressed_bytes + end function szgpu_max_compressed_bytes - function fsz_num_tiles(n) result(r) + function szgpu_num_tiles(n) result(r) integer(c_size_t), intent(in) :: n integer(c_size_t) :: r r = c_num_tiles(n) - end function fsz_num_tiles + end function szgpu_num_tiles - function fsz_version() result(v) + function szgpu_version() result(v) character(len=:), allocatable :: v v = to_f_string(c_version()) - end function fsz_version + end function szgpu_version - function fsz_status_message(s) result(v) + function szgpu_status_message(s) result(v) integer(c_int), intent(in) :: s character(len=:), allocatable :: v v = to_f_string(c_status_string(s)) - end function fsz_status_message + end function szgpu_status_message function to_f_string(p) result(v) type(c_ptr), intent(in) :: p @@ -530,4 +530,4 @@ function to_f_string(p) result(v) end do end function to_f_string -end module fsz +end module szgpu diff --git a/fortran/test_fsz.f90 b/fortran/test_szgpu.f90 similarity index 71% rename from fortran/test_fsz.f90 rename to fortran/test_szgpu.f90 index bf11c3f..e19bea4 100644 --- a/fortran/test_fsz.f90 +++ b/fortran/test_szgpu.f90 @@ -1,7 +1,7 @@ -program test_fsz +program test_szgpu use iso_c_binding, only: c_size_t, c_int, c_int8_t, c_ptr, c_loc, c_null_ptr, & c_float, c_associated - use fsz + use szgpu implicit none interface @@ -61,12 +61,12 @@ subroutine gate_r4() do i = 1, n a(i) = sin(0.0017 * real(i)) + 0.05 * cos(0.13 * real(i)) end do - allocate(cmp(fsz_max_compressed_bytes(n))) - call fsz_compress(a, n, eb, cmp, cmp_size, ierr) - call check('compress-r4', ierr == FSZ_OK .and. cmp_size > 0) + allocate(cmp(szgpu_max_compressed_bytes(n))) + call szgpu_compress(a, n, eb, cmp, cmp_size, ierr) + call check('compress-r4', ierr == SZGPU_OK .and. cmp_size > 0) b = 0.0 - call fsz_decompress(b, cmp, cmp_size, n, eb, ierr) - call check('roundtrip-r4', ierr == FSZ_OK .and. & + call szgpu_decompress(b, cmp, cmp_size, n, eb, ierr) + call check('roundtrip-r4', ierr == SZGPU_OK .and. & maxval(abs(a - b)) <= 1.01 * eb) print '(a,f0.3)', ' ratio r4: ', real(n) * 4.0 / real(cmp_size) end subroutine gate_r4 @@ -82,12 +82,12 @@ subroutine gate_r8() do i = 1, n a(i) = sin(0.0017d0 * real(i, 8)) + 0.05d0 * cos(0.13d0 * real(i, 8)) end do - allocate(cmp(fsz_max_compressed_bytes(n))) - call fsz_compress(a, n, eb, cmp, cmp_size, ierr) - call check('compress-r8', ierr == FSZ_OK .and. cmp_size > 0) + allocate(cmp(szgpu_max_compressed_bytes(n))) + call szgpu_compress(a, n, eb, cmp, cmp_size, ierr) + call check('compress-r8', ierr == SZGPU_OK .and. cmp_size > 0) b = 0.0d0 - call fsz_decompress(b, cmp, cmp_size, n, eb, ierr) - call check('roundtrip-r8', ierr == FSZ_OK .and. & + call szgpu_decompress(b, cmp, cmp_size, n, eb, ierr) + call check('roundtrip-r8', ierr == SZGPU_OK .and. & maxval(abs(a - b)) <= 1.01d0 * eb) print '(a,f0.3)', ' ratio r8: ', real(n) * 8.0 / real(cmp_size) end subroutine gate_r8 @@ -102,12 +102,12 @@ subroutine gate_small() do i = 1, n a(i) = sin(0.05 * real(i)) + 0.2 * cos(0.31 * real(i)) end do - allocate(cmp(fsz_max_compressed_bytes(n))) - call fsz_compress(a, n, eb, cmp, cmp_size, ierr) + allocate(cmp(szgpu_max_compressed_bytes(n))) + call szgpu_compress(a, n, eb, cmp, cmp_size, ierr) b = 0.0 - call fsz_decompress(b, cmp, cmp_size, n, eb, ierr) - call check('small-array', ierr == FSZ_OK .and. & - cmp_size > fsz_num_tiles(n) * 8 .and. & + call szgpu_decompress(b, cmp, cmp_size, n, eb, ierr) + call check('small-array', ierr == SZGPU_OK .and. & + cmp_size > szgpu_num_tiles(n) * 8 .and. & maxval(abs(a - b)) <= 1.01 * eb) end subroutine gate_small @@ -127,11 +127,11 @@ subroutine gate_rank3() end do end do n = int(nx, c_size_t) * ny * nz - allocate(cmp(fsz_max_compressed_bytes(n))) - call fsz_compress(a, n, eb, cmp, cmp_size, ierr) + allocate(cmp(szgpu_max_compressed_bytes(n))) + call szgpu_compress(a, n, eb, cmp, cmp_size, ierr) b = 0.0 - call fsz_decompress(b, cmp, cmp_size, n, eb, ierr) - call check('rank-3', ierr == FSZ_OK .and. & + call szgpu_decompress(b, cmp, cmp_size, n, eb, ierr) + call check('rank-3', ierr == SZGPU_OK .and. & maxval(abs(a - b)) <= 1.01 * eb) end subroutine gate_rank3 @@ -141,9 +141,9 @@ subroutine gate_error() integer(c_size_t) :: cmp_size integer(c_int) :: ierr a = 1.0 - call fsz_compress(a, 16_c_size_t, 0.0, cmp, cmp_size, ierr) - call check('error-status', ierr /= FSZ_OK .and. & - len(fsz_status_message(ierr)) > 0) + call szgpu_compress(a, 16_c_size_t, 0.0, cmp, cmp_size, ierr) + call check('error-status', ierr /= SZGPU_OK .and. & + len(szgpu_status_message(ierr)) > 0) end subroutine gate_error subroutine gate_device() @@ -153,7 +153,7 @@ subroutine gate_device() integer(c_size_t) :: host_size, cap, i integer(c_int) :: ierr, cst type(c_ptr) :: d_in, d_out, d_cmp, ws - type(fsz_result) :: res + type(szgpu_result) :: res real(4), parameter :: eb = 1.0e-3 allocate(a(n), b(n)) @@ -161,10 +161,10 @@ subroutine gate_device() a(i) = sin(0.0017 * real(i)) + 0.05 * cos(0.13 * real(i)) end do - cap = fsz_max_compressed_bytes(n) + cap = szgpu_max_compressed_bytes(n) allocate(cmp_host(cap)) - call fsz_compress(a, n, eb, cmp_host, host_size, ierr) - call check('device-setup-host', ierr == FSZ_OK) + call szgpu_compress(a, n, eb, cmp_host, host_size, ierr) + call check('device-setup-host', ierr == SZGPU_OK) cst = cudaMalloc(d_in, n * 4_c_size_t) if (cst /= 0) then @@ -176,30 +176,30 @@ subroutine gate_device() cst = cudaMemcpy(d_in, c_loc(a), n * 4_c_size_t, H2D) call check('device-h2d', cst == 0) - call fsz_workspace_create(ws, n, ierr) - call check('workspace-create', ierr == FSZ_OK .and. c_associated(ws) .and. & - fsz_workspace_capacity(ws) == n) + call szgpu_workspace_create(ws, n, ierr) + call check('workspace-create', ierr == SZGPU_OK .and. c_associated(ws) .and. & + szgpu_workspace_capacity(ws) == n) - call fsz_compress_device(d_in, d_cmp, n, eb, ws, res, ierr) - call check('compress-device', ierr == FSZ_OK .and. res%cmp_size == host_size) + call szgpu_compress_device(d_in, d_cmp, n, eb, ws, res, ierr) + call check('compress-device', ierr == SZGPU_OK .and. res%cmp_size == host_size) - call fsz_decompress_device(d_out, d_cmp, n, eb, res, ws, ierr) + call szgpu_decompress_device(d_out, d_cmp, n, eb, res, ws, ierr) cst = cudaDeviceSynchronize() b = 0.0 cst = cudaMemcpy(c_loc(b), d_out, n * 4_c_size_t, D2H) - call check('decompress-device', ierr == FSZ_OK .and. cst == 0 .and. & + call check('decompress-device', ierr == SZGPU_OK .and. cst == 0 .and. & maxval(abs(real(a, 8) - real(b, 8))) <= 1.01d0 * real(eb, 8)) block - type(fsz_result) :: res2 - call fsz_compress_device(d_in, d_cmp, n, eb, c_null_ptr, res2, ierr) - call check('compress-device-nullws', ierr == FSZ_OK .and. & + type(szgpu_result) :: res2 + call szgpu_compress_device(d_in, d_cmp, n, eb, c_null_ptr, res2, ierr) + call check('compress-device-nullws', ierr == SZGPU_OK .and. & res2%cmp_size == host_size) end block print '(a,i0,a)', ' device stream ', res%cmp_size, ' bytes, matches the host path' - call fsz_workspace_destroy(ws) + call szgpu_workspace_destroy(ws) cst = cudaFree(d_in); cst = cudaFree(d_out); cst = cudaFree(d_cmp) end subroutine gate_device @@ -214,4 +214,4 @@ subroutine check(name, ok) end if end subroutine check -end program test_fsz +end program test_szgpu diff --git a/include/fsz/fsz.h b/include/fsz/fsz.h deleted file mode 100644 index 46992cd..0000000 --- a/include/fsz/fsz.h +++ /dev/null @@ -1,113 +0,0 @@ -#ifndef FSZ_H -#define FSZ_H - -#include -#include - -#ifdef __cplusplus -extern "C" { -#endif - -#define FSZ_VERSION_MAJOR 1 -#define FSZ_VERSION_MINOR 0 -#define FSZ_VERSION_PATCH 0 -#define FSZ_VERSION_STRING "1.0.0" - -typedef enum { - FSZ_STATUS_OK = 0, - FSZ_STATUS_INVALID_ARGUMENT = 1, - FSZ_STATUS_OUT_OF_MEMORY = 2, - FSZ_STATUS_WORKSPACE_TOO_SMALL = 3, - FSZ_STATUS_CUDA_ERROR = 4, - FSZ_STATUS_TOO_LARGE = 5 -} fsz_status_t; - -typedef struct fsz_workspace fsz_workspace_t; - -typedef struct { - size_t cmp_size; - size_t num_tiles; - size_t data_offset; -} fsz_compress_result_t; - -const char* fsz_version(void); -const char* fsz_status_string(fsz_status_t status); - -size_t fsz_num_tiles(size_t n_elements); -size_t fsz_max_compressed_bytes(size_t n_elements); - -fsz_status_t fsz_workspace_create(fsz_workspace_t** out_ws, size_t max_n_elements); -void fsz_workspace_destroy(fsz_workspace_t* ws); -size_t fsz_workspace_capacity(const fsz_workspace_t* ws); - -fsz_status_t fsz_compress( - const float* d_in, - unsigned char* d_cmp, - size_t n_elements, - float abs_error_bound, - fsz_workspace_t* ws, - cudaStream_t stream, - fsz_compress_result_t* out_result); - -fsz_status_t fsz_decompress( - float* d_out, - const unsigned char* d_cmp, - size_t n_elements, - float abs_error_bound, - const fsz_compress_result_t* result, - fsz_workspace_t* ws, - cudaStream_t stream); - -fsz_compress_result_t fsz_make_result(size_t n_elements, size_t cmp_size); - -fsz_status_t fsz_compress_hostptr( - const float* in, - unsigned char* cmp, - size_t n_elements, - float abs_error_bound, - size_t* out_cmp_size); - -fsz_status_t fsz_decompress_hostptr( - float* out, - const unsigned char* cmp, - size_t cmp_size, - size_t n_elements, - float abs_error_bound); - -fsz_status_t fsz_compress_f64( - const double* d_in, - unsigned char* d_cmp, - size_t n_elements, - double abs_error_bound, - fsz_workspace_t* ws, - cudaStream_t stream, - fsz_compress_result_t* out_result); - -fsz_status_t fsz_decompress_f64( - double* d_out, - const unsigned char* d_cmp, - size_t n_elements, - double abs_error_bound, - const fsz_compress_result_t* result, - fsz_workspace_t* ws, - cudaStream_t stream); - -fsz_status_t fsz_compress_hostptr_f64( - const double* in, - unsigned char* cmp, - size_t n_elements, - double abs_error_bound, - size_t* out_cmp_size); - -fsz_status_t fsz_decompress_hostptr_f64( - double* out, - const unsigned char* cmp, - size_t cmp_size, - size_t n_elements, - double abs_error_bound); - -#ifdef __cplusplus -} -#endif - -#endif diff --git a/include/fsz/fsz.hpp b/include/fsz/fsz.hpp deleted file mode 100644 index 10f2e2c..0000000 --- a/include/fsz/fsz.hpp +++ /dev/null @@ -1,176 +0,0 @@ -#ifndef FSZ_HPP -#define FSZ_HPP - -#include "fsz/fsz.h" - -#include -#include -#include - -namespace fsz { - -class Error : public std::runtime_error { -public: - Error(fsz_status_t status, const std::string& where) - : std::runtime_error(where + ": " + fsz_status_string(status)), - status_(status) {} - fsz_status_t status() const noexcept { return status_; } -private: - fsz_status_t status_; -}; - -inline void check(fsz_status_t s, const char* where) { - if (s != FSZ_STATUS_OK) throw Error(s, where); -} - -class Workspace { -public: - explicit Workspace(std::size_t max_n_elements) { - check(fsz_workspace_create(&ws_, max_n_elements), "fsz_workspace_create"); - } - ~Workspace() { if (ws_) fsz_workspace_destroy(ws_); } - - Workspace(const Workspace&) = delete; - Workspace& operator=(const Workspace&) = delete; - Workspace(Workspace&& other) noexcept : ws_(other.ws_) { other.ws_ = nullptr; } - Workspace& operator=(Workspace&& other) noexcept { - if (this != &other) { - if (ws_) fsz_workspace_destroy(ws_); - ws_ = other.ws_; other.ws_ = nullptr; - } - return *this; - } - - fsz_workspace_t* handle() noexcept { return ws_; } - std::size_t capacity() const noexcept { return fsz_workspace_capacity(ws_); } - -private: - fsz_workspace_t* ws_ = nullptr; -}; - -inline fsz_compress_result_t compress( - const float* d_in, unsigned char* d_cmp, - std::size_t n, float abs_error_bound, - Workspace& ws, cudaStream_t stream = 0) -{ - fsz_compress_result_t r{}; - check(fsz_compress(d_in, d_cmp, n, abs_error_bound, ws.handle(), stream, &r), - "fsz_compress"); - return r; -} - -inline void decompress( - float* d_out, const unsigned char* d_cmp, - std::size_t n, float abs_error_bound, - const fsz_compress_result_t& result, - Workspace& ws, cudaStream_t stream = 0) -{ - check(fsz_decompress(d_out, d_cmp, n, abs_error_bound, &result, ws.handle(), stream), - "fsz_decompress"); -} - -inline fsz_compress_result_t compress( - const float* d_in, unsigned char* d_cmp, - std::size_t n, float abs_error_bound, cudaStream_t stream = 0) -{ - fsz_compress_result_t r{}; - check(fsz_compress(d_in, d_cmp, n, abs_error_bound, nullptr, stream, &r), - "fsz_compress"); - return r; -} - -inline fsz_compress_result_t compress( - const double* d_in, unsigned char* d_cmp, - std::size_t n, double abs_error_bound, cudaStream_t stream = 0) -{ - fsz_compress_result_t r{}; - check(fsz_compress_f64(d_in, d_cmp, n, abs_error_bound, nullptr, stream, &r), - "fsz_compress_f64"); - return r; -} - -inline void decompress( - float* d_out, const unsigned char* d_cmp, - std::size_t n, float abs_error_bound, - const fsz_compress_result_t& result, cudaStream_t stream = 0) -{ - check(fsz_decompress(d_out, d_cmp, n, abs_error_bound, &result, nullptr, stream), - "fsz_decompress"); -} - -inline void decompress( - double* d_out, const unsigned char* d_cmp, - std::size_t n, double abs_error_bound, - const fsz_compress_result_t& result, cudaStream_t stream = 0) -{ - check(fsz_decompress_f64(d_out, d_cmp, n, abs_error_bound, &result, nullptr, stream), - "fsz_decompress_f64"); -} - -inline std::size_t compress_hostptr( - const float* in, unsigned char* cmp, - std::size_t n, float abs_error_bound) -{ - std::size_t cmp_size = 0; - check(fsz_compress_hostptr(in, cmp, n, abs_error_bound, &cmp_size), - "fsz_compress_hostptr"); - return cmp_size; -} - -inline void decompress_hostptr( - float* out, const unsigned char* cmp, std::size_t cmp_size, - std::size_t n, float abs_error_bound) -{ - check(fsz_decompress_hostptr(out, cmp, cmp_size, n, abs_error_bound), - "fsz_decompress_hostptr"); -} - -inline fsz_compress_result_t make_result(std::size_t n, std::size_t cmp_size) { - return fsz_make_result(n, cmp_size); -} - -inline fsz_compress_result_t compress( - const double* d_in, unsigned char* d_cmp, - std::size_t n, double abs_error_bound, - Workspace& ws, cudaStream_t stream = 0) -{ - fsz_compress_result_t r{}; - check(fsz_compress_f64(d_in, d_cmp, n, abs_error_bound, ws.handle(), stream, &r), - "fsz_compress_f64"); - return r; -} - -inline void decompress( - double* d_out, const unsigned char* d_cmp, - std::size_t n, double abs_error_bound, - const fsz_compress_result_t& result, - Workspace& ws, cudaStream_t stream = 0) -{ - check(fsz_decompress_f64(d_out, d_cmp, n, abs_error_bound, &result, ws.handle(), stream), - "fsz_decompress_f64"); -} - -inline std::size_t compress_hostptr( - const double* in, unsigned char* cmp, std::size_t n, double abs_error_bound) -{ - std::size_t cmp_size = 0; - check(fsz_compress_hostptr_f64(in, cmp, n, abs_error_bound, &cmp_size), - "fsz_compress_hostptr_f64"); - return cmp_size; -} - -inline void decompress_hostptr( - double* out, const unsigned char* cmp, std::size_t cmp_size, - std::size_t n, double abs_error_bound) -{ - check(fsz_decompress_hostptr_f64(out, cmp, cmp_size, n, abs_error_bound), - "fsz_decompress_hostptr_f64"); -} - -inline std::size_t max_compressed_bytes(std::size_t n) { return fsz_max_compressed_bytes(n); } -inline std::size_t num_tiles(std::size_t n) { return fsz_num_tiles(n); } -inline const char* version() { return fsz_version(); } - -} - -#endif diff --git a/include/szgpu/szgpu.h b/include/szgpu/szgpu.h new file mode 100644 index 0000000..c4c7f45 --- /dev/null +++ b/include/szgpu/szgpu.h @@ -0,0 +1,113 @@ +#ifndef SZGPU_H +#define SZGPU_H + +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +#define SZGPU_VERSION_MAJOR 1 +#define SZGPU_VERSION_MINOR 0 +#define SZGPU_VERSION_PATCH 1 +#define SZGPU_VERSION_STRING "1.0.1" + +typedef enum { + SZGPU_STATUS_OK = 0, + SZGPU_STATUS_INVALID_ARGUMENT = 1, + SZGPU_STATUS_OUT_OF_MEMORY = 2, + SZGPU_STATUS_WORKSPACE_TOO_SMALL = 3, + SZGPU_STATUS_CUDA_ERROR = 4, + SZGPU_STATUS_TOO_LARGE = 5 +} szgpu_status_t; + +typedef struct szgpu_workspace szgpu_workspace_t; + +typedef struct { + size_t cmp_size; + size_t num_tiles; + size_t data_offset; +} szgpu_compress_result_t; + +const char* szgpu_version(void); +const char* szgpu_status_string(szgpu_status_t status); + +size_t szgpu_num_tiles(size_t n_elements); +size_t szgpu_max_compressed_bytes(size_t n_elements); + +szgpu_status_t szgpu_workspace_create(szgpu_workspace_t** out_ws, size_t max_n_elements); +void szgpu_workspace_destroy(szgpu_workspace_t* ws); +size_t szgpu_workspace_capacity(const szgpu_workspace_t* ws); + +szgpu_status_t szgpu_compress( + const float* d_in, + unsigned char* d_cmp, + size_t n_elements, + float abs_error_bound, + szgpu_workspace_t* ws, + cudaStream_t stream, + szgpu_compress_result_t* out_result); + +szgpu_status_t szgpu_decompress( + float* d_out, + const unsigned char* d_cmp, + size_t n_elements, + float abs_error_bound, + const szgpu_compress_result_t* result, + szgpu_workspace_t* ws, + cudaStream_t stream); + +szgpu_compress_result_t szgpu_make_result(size_t n_elements, size_t cmp_size); + +szgpu_status_t szgpu_compress_hostptr( + const float* in, + unsigned char* cmp, + size_t n_elements, + float abs_error_bound, + size_t* out_cmp_size); + +szgpu_status_t szgpu_decompress_hostptr( + float* out, + const unsigned char* cmp, + size_t cmp_size, + size_t n_elements, + float abs_error_bound); + +szgpu_status_t szgpu_compress_f64( + const double* d_in, + unsigned char* d_cmp, + size_t n_elements, + double abs_error_bound, + szgpu_workspace_t* ws, + cudaStream_t stream, + szgpu_compress_result_t* out_result); + +szgpu_status_t szgpu_decompress_f64( + double* d_out, + const unsigned char* d_cmp, + size_t n_elements, + double abs_error_bound, + const szgpu_compress_result_t* result, + szgpu_workspace_t* ws, + cudaStream_t stream); + +szgpu_status_t szgpu_compress_hostptr_f64( + const double* in, + unsigned char* cmp, + size_t n_elements, + double abs_error_bound, + size_t* out_cmp_size); + +szgpu_status_t szgpu_decompress_hostptr_f64( + double* out, + const unsigned char* cmp, + size_t cmp_size, + size_t n_elements, + double abs_error_bound); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/include/szgpu/szgpu.hpp b/include/szgpu/szgpu.hpp new file mode 100644 index 0000000..2f6b9f8 --- /dev/null +++ b/include/szgpu/szgpu.hpp @@ -0,0 +1,176 @@ +#ifndef SZGPU_HPP +#define SZGPU_HPP + +#include "szgpu/szgpu.h" + +#include +#include +#include + +namespace szgpu { + +class Error : public std::runtime_error { +public: + Error(szgpu_status_t status, const std::string& where) + : std::runtime_error(where + ": " + szgpu_status_string(status)), + status_(status) {} + szgpu_status_t status() const noexcept { return status_; } +private: + szgpu_status_t status_; +}; + +inline void check(szgpu_status_t s, const char* where) { + if (s != SZGPU_STATUS_OK) throw Error(s, where); +} + +class Workspace { +public: + explicit Workspace(std::size_t max_n_elements) { + check(szgpu_workspace_create(&ws_, max_n_elements), "szgpu_workspace_create"); + } + ~Workspace() { if (ws_) szgpu_workspace_destroy(ws_); } + + Workspace(const Workspace&) = delete; + Workspace& operator=(const Workspace&) = delete; + Workspace(Workspace&& other) noexcept : ws_(other.ws_) { other.ws_ = nullptr; } + Workspace& operator=(Workspace&& other) noexcept { + if (this != &other) { + if (ws_) szgpu_workspace_destroy(ws_); + ws_ = other.ws_; other.ws_ = nullptr; + } + return *this; + } + + szgpu_workspace_t* handle() noexcept { return ws_; } + std::size_t capacity() const noexcept { return szgpu_workspace_capacity(ws_); } + +private: + szgpu_workspace_t* ws_ = nullptr; +}; + +inline szgpu_compress_result_t compress( + const float* d_in, unsigned char* d_cmp, + std::size_t n, float abs_error_bound, + Workspace& ws, cudaStream_t stream = 0) +{ + szgpu_compress_result_t r{}; + check(szgpu_compress(d_in, d_cmp, n, abs_error_bound, ws.handle(), stream, &r), + "szgpu_compress"); + return r; +} + +inline void decompress( + float* d_out, const unsigned char* d_cmp, + std::size_t n, float abs_error_bound, + const szgpu_compress_result_t& result, + Workspace& ws, cudaStream_t stream = 0) +{ + check(szgpu_decompress(d_out, d_cmp, n, abs_error_bound, &result, ws.handle(), stream), + "szgpu_decompress"); +} + +inline szgpu_compress_result_t compress( + const float* d_in, unsigned char* d_cmp, + std::size_t n, float abs_error_bound, cudaStream_t stream = 0) +{ + szgpu_compress_result_t r{}; + check(szgpu_compress(d_in, d_cmp, n, abs_error_bound, nullptr, stream, &r), + "szgpu_compress"); + return r; +} + +inline szgpu_compress_result_t compress( + const double* d_in, unsigned char* d_cmp, + std::size_t n, double abs_error_bound, cudaStream_t stream = 0) +{ + szgpu_compress_result_t r{}; + check(szgpu_compress_f64(d_in, d_cmp, n, abs_error_bound, nullptr, stream, &r), + "szgpu_compress_f64"); + return r; +} + +inline void decompress( + float* d_out, const unsigned char* d_cmp, + std::size_t n, float abs_error_bound, + const szgpu_compress_result_t& result, cudaStream_t stream = 0) +{ + check(szgpu_decompress(d_out, d_cmp, n, abs_error_bound, &result, nullptr, stream), + "szgpu_decompress"); +} + +inline void decompress( + double* d_out, const unsigned char* d_cmp, + std::size_t n, double abs_error_bound, + const szgpu_compress_result_t& result, cudaStream_t stream = 0) +{ + check(szgpu_decompress_f64(d_out, d_cmp, n, abs_error_bound, &result, nullptr, stream), + "szgpu_decompress_f64"); +} + +inline std::size_t compress_hostptr( + const float* in, unsigned char* cmp, + std::size_t n, float abs_error_bound) +{ + std::size_t cmp_size = 0; + check(szgpu_compress_hostptr(in, cmp, n, abs_error_bound, &cmp_size), + "szgpu_compress_hostptr"); + return cmp_size; +} + +inline void decompress_hostptr( + float* out, const unsigned char* cmp, std::size_t cmp_size, + std::size_t n, float abs_error_bound) +{ + check(szgpu_decompress_hostptr(out, cmp, cmp_size, n, abs_error_bound), + "szgpu_decompress_hostptr"); +} + +inline szgpu_compress_result_t make_result(std::size_t n, std::size_t cmp_size) { + return szgpu_make_result(n, cmp_size); +} + +inline szgpu_compress_result_t compress( + const double* d_in, unsigned char* d_cmp, + std::size_t n, double abs_error_bound, + Workspace& ws, cudaStream_t stream = 0) +{ + szgpu_compress_result_t r{}; + check(szgpu_compress_f64(d_in, d_cmp, n, abs_error_bound, ws.handle(), stream, &r), + "szgpu_compress_f64"); + return r; +} + +inline void decompress( + double* d_out, const unsigned char* d_cmp, + std::size_t n, double abs_error_bound, + const szgpu_compress_result_t& result, + Workspace& ws, cudaStream_t stream = 0) +{ + check(szgpu_decompress_f64(d_out, d_cmp, n, abs_error_bound, &result, ws.handle(), stream), + "szgpu_decompress_f64"); +} + +inline std::size_t compress_hostptr( + const double* in, unsigned char* cmp, std::size_t n, double abs_error_bound) +{ + std::size_t cmp_size = 0; + check(szgpu_compress_hostptr_f64(in, cmp, n, abs_error_bound, &cmp_size), + "szgpu_compress_hostptr_f64"); + return cmp_size; +} + +inline void decompress_hostptr( + double* out, const unsigned char* cmp, std::size_t cmp_size, + std::size_t n, double abs_error_bound) +{ + check(szgpu_decompress_hostptr_f64(out, cmp, cmp_size, n, abs_error_bound), + "szgpu_decompress_hostptr_f64"); +} + +inline std::size_t max_compressed_bytes(std::size_t n) { return szgpu_max_compressed_bytes(n); } +inline std::size_t num_tiles(std::size_t n) { return szgpu_num_tiles(n); } +inline const char* version() { return szgpu_version(); } + +} + +#endif diff --git a/python/MANIFEST.in b/python/MANIFEST.in index 945ef64..ac59a85 100644 --- a/python/MANIFEST.in +++ b/python/MANIFEST.in @@ -1,2 +1,2 @@ include README.md -recursive-include fsz/_cuda *.cu *.cuh *.h *.hpp +recursive-include szgpu/_cuda *.cu *.cuh *.h *.hpp diff --git a/python/example.py b/python/example.py index 0ceefa4..e999245 100644 --- a/python/example.py +++ b/python/example.py @@ -1,5 +1,5 @@ #!/usr/bin/env python3 -"""A short tour of the FSZ Python bindings. +"""A short tour of the SZ-GPU Python bindings. Builds a smooth 128 x 128 x 128 float32 field, compresses it with a relative error bound of 1e-3, decompresses it, and reports what that cost and what it @@ -15,14 +15,14 @@ import numpy as np try: - import fsz + import szgpu except ImportError: sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))) - import fsz + import szgpu def smooth_field(n=128): - """A smooth, wavy scalar field of shape (n, n, n), the kind of data FSZ is for.""" + """A smooth, wavy scalar field of shape (n, n, n), the kind of data SZ-GPU is for.""" axis = np.linspace(0.0, 2.0 * np.pi, n, dtype=np.float32) z, y, x = np.meshgrid(axis, axis, axis, indexing="ij") field = (np.sin(x) * np.cos(y) @@ -32,8 +32,8 @@ def smooth_field(n=128): def main(): - print("FSZ Python bindings %s (library %s)" - % (fsz.__version__, fsz.library_version())) + print("SZ-GPU Python bindings %s (library %s)" + % (szgpu.__version__, szgpu.library_version())) print() field = smooth_field(128) @@ -45,11 +45,11 @@ def main(): print() t0 = time.time() - blob = fsz.compress(field, eb_rel, mode="rel") + blob = szgpu.compress(field, eb_rel, mode="rel") t_compress = time.time() - t0 t0 = time.time() - back = fsz.decompress(blob) + back = szgpu.decompress(blob) t_decompress = time.time() - t0 value_range = np.float32(field.max() - field.min()) @@ -62,7 +62,7 @@ def main(): within = max_err <= tolerance * float(abs_eb) print("compressed : %d bytes (%d header + %d bitstream) in %.3f s" - % (len(blob), fsz.HEADER_SIZE, len(blob) - fsz.HEADER_SIZE, t_compress)) + % (len(blob), szgpu.HEADER_SIZE, len(blob) - szgpu.HEADER_SIZE, t_compress)) print("ratio : %.2fx smaller than the raw float32 data" % ratio) print("decompressed: %.3f s" % t_decompress) print() @@ -80,8 +80,8 @@ def main(): print("A blob carries its own shape and bound, so a file round trip needs") print("no extra bookkeeping:") - print(" fsz.compress_file(src, dst, 1e-3, dims=(128, 128, 128))") - print(" field = fsz.decompress_file(dst)") + print(" szgpu.compress_file(src, dst, 1e-3, dims=(128, 128, 128))") + print(" field = szgpu.decompress_file(dst)") return 0 if (within and back.shape == field.shape) else 1 diff --git a/python/pyproject.toml b/python/pyproject.toml index 30827d8..def0c0d 100644 --- a/python/pyproject.toml +++ b/python/pyproject.toml @@ -3,9 +3,9 @@ requires = ["setuptools>=61"] build-backend = "setuptools.build_meta" [project] -name = "fsz" -version = "1.0.0" -description = "Python bindings for the FSZ GPU error-bounded lossy compressor" +name = "szgpu" +version = "1.0.1" +description = "Python bindings for the SZ-GPU error-bounded lossy compressor" requires-python = ">=3.8" license = "BSD-3-Clause" dependencies = ["numpy"] @@ -18,4 +18,4 @@ classifiers = [ ] [tool.setuptools] -packages = ["fsz"] +packages = ["szgpu"] diff --git a/python/setup.py b/python/setup.py index 2b17964..e26adf6 100644 --- a/python/setup.py +++ b/python/setup.py @@ -1,10 +1,10 @@ -"""Build libfsz.so into the package at install time. +"""Build libszgpu.so into the package at install time. The Python API is a ctypes wrapper, so the wheel needs the shared library. If a CUDA toolkit is available the library is compiled here and installed alongside the module, which makes ``pip install`` self-contained. If nvcc is missing the build is skipped with a warning rather than failing: the module can still find -a library through FSZ_LIB, a repository build directory, or the system search +a library through SZGPU_LIB, a repository build directory, or the system search path, and reporting that at import time is more useful than refusing to install. Sources come from the repository when installing from a checkout, or from the @@ -21,7 +21,7 @@ from setuptools.command.sdist import sdist as _sdist HERE = os.path.dirname(os.path.abspath(__file__)) -VENDOR = os.path.join(HERE, "fsz", "_cuda") +VENDOR = os.path.join(HERE, "szgpu", "_cuda") # Architectures to embed. PTX for the newest entry is added as well, so GPU # generations released after this build still run. @@ -55,7 +55,7 @@ def _find_nvcc(): def _archs(): - override = os.environ.get("FSZ_CUDA_ARCHITECTURES") + override = os.environ.get("SZGPU_CUDA_ARCHITECTURES") if override: return [a for a in override.replace(",", ";").split(";") if a] return list(DEFAULT_ARCHS) @@ -64,12 +64,12 @@ def _archs(): def _build_library(target_dir): src_dir, inc_dir = _source_dirs() if not src_dir: - print("FSZ: CUDA sources not found; skipping the library build.") + print("SZ-GPU: CUDA sources not found; skipping the library build.") return False nvcc = _find_nvcc() if not nvcc: - print("FSZ: nvcc not found; skipping the library build. Set FSZ_LIB to " - "an existing libfsz.so, or install a CUDA toolkit and reinstall.") + print("SZ-GPU: nvcc not found; skipping the library build. Set SZGPU_LIB to " + "an existing libszgpu.so, or install a CUDA toolkit and reinstall.") return False archs = _archs() @@ -79,16 +79,16 @@ def _build_library(target_dir): gencode += ["-gencode", "arch=compute_%s,code=compute_%s" % (archs[-1], archs[-1])] os.makedirs(target_dir, exist_ok=True) - out = os.path.join(target_dir, "libfsz.so") + out = os.path.join(target_dir, "libszgpu.so") cmd = [nvcc, "-O3", "--std=c++17", "-shared", "-Xcompiler", "-fPIC"] + gencode + [ - os.path.join(src_dir, "fsz.cu"), "-o", out, + os.path.join(src_dir, "szgpu.cu"), "-o", out, "-I", inc_dir, "-I", src_dir, ] - print("FSZ: building the shared library for sm_%s" % ", sm_".join(archs)) + print("SZ-GPU: building the shared library for sm_%s" % ", sm_".join(archs)) try: subprocess.check_call(cmd) except (subprocess.CalledProcessError, OSError) as exc: - print("FSZ: the library build failed (%s); continuing without a bundled " + print("SZ-GPU: the library build failed (%s); continuing without a bundled " "library." % exc, file=sys.stderr) return False return True @@ -97,7 +97,7 @@ def _build_library(target_dir): class build_py(_build_py): def run(self): super().run() - target = os.path.join(self.build_lib, "fsz") + target = os.path.join(self.build_lib, "szgpu") _build_library(target) @@ -118,5 +118,5 @@ def run(self): setup( cmdclass={"build_py": build_py, "sdist": sdist}, - package_data={"fsz": ["libfsz.so", "_cuda/src/*", "_cuda/include/fsz/*"]}, + package_data={"szgpu": ["libszgpu.so", "_cuda/src/*", "_cuda/include/szgpu/*"]}, ) diff --git a/python/fsz/__init__.py b/python/szgpu/__init__.py similarity index 84% rename from python/fsz/__init__.py rename to python/szgpu/__init__.py index 82d3872..bff9cee 100644 --- a/python/fsz/__init__.py +++ b/python/szgpu/__init__.py @@ -1,15 +1,15 @@ -"""Python bindings for the FSZ GPU error-bounded lossy compressor. +"""Python bindings for the SZ-GPU error-bounded lossy compressor. -FSZ compresses float32 and float64 arrays on the GPU under an absolute +SZ-GPU compresses float32 and float64 arrays on the GPU under an absolute error bound: every reconstructed value lands within that bound of the original, up to rounding on the last bit. The library checks itself against the bound with one percent of slack, and so should you. These bindings talk to the shared library through ctypes and numpy, with no other dependencies. - import numpy as np, fsz + import numpy as np, szgpu field = np.sin(np.linspace(0, 8, 64 * 64 * 64, dtype=np.float32)).reshape(64, 64, 64) - blob = fsz.compress(field, 1e-3, mode="rel") # bytes, self-describing - back = fsz.decompress(blob) # float32, original shape + blob = szgpu.compress(field, 1e-3, mode="rel") # bytes, self-describing + back = szgpu.decompress(blob) # float32, original shape print(field.nbytes / len(blob), np.abs(field - back).max()) The element type follows the input array: a float64 array compresses as @@ -20,7 +20,7 @@ :func:`compress` returns a self-describing container: a 56 byte header holding the shape, the element type, the absolute error bound and the bitstream length, followed by the bitstream exactly as the library produced it. The layout is the -one the ``fsz`` command line tool reads and writes, so blobs move between the +one the ``szgpu`` command line tool reads and writes, so blobs move between the two freely. Arrays of any dimensionality are accepted, up to the three dimensions that the @@ -28,9 +28,9 @@ The shared library is located at import time, in this order: -1. the path in the ``FSZ_LIB`` environment variable, -2. ``build/libfsz.so`` beside the package in a repository checkout, -3. whatever ``ctypes.util.find_library("fsz")`` turns up. +1. the path in the ``SZGPU_LIB`` environment variable, +2. ``build/libszgpu.so`` beside the package in a repository checkout, +3. whatever ``ctypes.util.find_library("szgpu")`` turns up. """ import ctypes @@ -40,7 +40,7 @@ import numpy as np -__version__ = "1.0.0" +__version__ = "1.0.1" __all__ = [ "compress", @@ -69,33 +69,33 @@ MAGIC = b"FSZ1" assert struct.calcsize(HEADER_FORMAT) == 56, ( - "the FSZ header layout must pack to exactly 56 bytes, got %d" + "the SZ-GPU header layout must pack to exactly 56 bytes, got %d" % struct.calcsize(HEADER_FORMAT) ) def _repo_build_candidate(): - """Path to build/libfsz.so as laid out in a repository checkout.""" + """Path to build/libszgpu.so as laid out in a repository checkout.""" here = os.path.dirname(os.path.abspath(__file__)) - return os.path.normpath(os.path.join(here, "..", "..", "build", "libfsz.so")) + return os.path.normpath(os.path.join(here, "..", "..", "build", "libszgpu.so")) def _load_library(): - """Find and open libfsz.so, or raise ImportError naming every place tried.""" + """Find and open libszgpu.so, or raise ImportError naming every place tried.""" attempts = [] - env_path = os.environ.get("FSZ_LIB") + env_path = os.environ.get("SZGPU_LIB") if env_path: - attempts.append(("FSZ_LIB environment variable", env_path)) + attempts.append(("SZGPU_LIB environment variable", env_path)) else: - attempts.append(("FSZ_LIB environment variable", None)) + attempts.append(("SZGPU_LIB environment variable", None)) - bundled = os.path.join(os.path.dirname(os.path.abspath(__file__)), "libfsz.so") + bundled = os.path.join(os.path.dirname(os.path.abspath(__file__)), "libszgpu.so") attempts.append(("library built into the installed package", bundled if os.path.exists(bundled) else None)) attempts.append(("build directory beside the package", _repo_build_candidate())) - attempts.append(("system search path", ctypes.util.find_library("fsz"))) + attempts.append(("system search path", ctypes.util.find_library("szgpu"))) notes = [] for source, path in attempts: @@ -111,12 +111,12 @@ def _load_library(): notes.append(" %s: %s (%s)" % (source, path, exc)) raise ImportError( - "the FSZ shared library could not be loaded. Tried, in order:\n" + "the SZ-GPU shared library could not be loaded. Tried, in order:\n" + "\n".join(notes) + "\n\nBuild it with\n" - " cmake -S -B /build -DFSZ_BUILD_SHARED=ON\n" - " cmake --build /build --target fsz_shared -j4\n" - "or point FSZ_LIB at an existing libfsz.so. If the library exists but " + " cmake -S -B /build -DSZGPU_BUILD_SHARED=ON\n" + " cmake --build /build --target szgpu_shared -j4\n" + "or point SZGPU_LIB at an existing libszgpu.so. If the library exists but " "fails to load, check that the CUDA runtime (libcudart) is on " "LD_LIBRARY_PATH." ) @@ -129,60 +129,60 @@ def _load_library(): _c_double_p = ctypes.POINTER(ctypes.c_double) _c_ubyte_p = ctypes.POINTER(ctypes.c_ubyte) -_lib.fsz_version.argtypes = [] -_lib.fsz_version.restype = ctypes.c_char_p +_lib.szgpu_version.argtypes = [] +_lib.szgpu_version.restype = ctypes.c_char_p -_lib.fsz_status_string.argtypes = [ctypes.c_int] -_lib.fsz_status_string.restype = ctypes.c_char_p +_lib.szgpu_status_string.argtypes = [ctypes.c_int] +_lib.szgpu_status_string.restype = ctypes.c_char_p -_lib.fsz_num_tiles.argtypes = [ctypes.c_size_t] -_lib.fsz_num_tiles.restype = ctypes.c_size_t +_lib.szgpu_num_tiles.argtypes = [ctypes.c_size_t] +_lib.szgpu_num_tiles.restype = ctypes.c_size_t -_lib.fsz_max_compressed_bytes.argtypes = [ctypes.c_size_t] -_lib.fsz_max_compressed_bytes.restype = ctypes.c_size_t +_lib.szgpu_max_compressed_bytes.argtypes = [ctypes.c_size_t] +_lib.szgpu_max_compressed_bytes.restype = ctypes.c_size_t -_lib.fsz_compress_hostptr.argtypes = [ +_lib.szgpu_compress_hostptr.argtypes = [ _c_float_p, # in _c_ubyte_p, # cmp ctypes.c_size_t, # n_elements ctypes.c_float, # abs_error_bound _c_size_p, # out_cmp_size ] -_lib.fsz_compress_hostptr.restype = ctypes.c_int +_lib.szgpu_compress_hostptr.restype = ctypes.c_int -_lib.fsz_decompress_hostptr.argtypes = [ +_lib.szgpu_decompress_hostptr.argtypes = [ _c_float_p, # out _c_ubyte_p, # cmp ctypes.c_size_t, # cmp_size ctypes.c_size_t, # n_elements ctypes.c_float, # abs_error_bound ] -_lib.fsz_decompress_hostptr.restype = ctypes.c_int +_lib.szgpu_decompress_hostptr.restype = ctypes.c_int -_lib.fsz_compress_hostptr_f64.argtypes = [ +_lib.szgpu_compress_hostptr_f64.argtypes = [ _c_double_p, # in _c_ubyte_p, # cmp ctypes.c_size_t, # n_elements ctypes.c_double, # abs_error_bound _c_size_p, # out_cmp_size ] -_lib.fsz_compress_hostptr_f64.restype = ctypes.c_int +_lib.szgpu_compress_hostptr_f64.restype = ctypes.c_int -_lib.fsz_decompress_hostptr_f64.argtypes = [ +_lib.szgpu_decompress_hostptr_f64.argtypes = [ _c_double_p, # out _c_ubyte_p, # cmp ctypes.c_size_t, # cmp_size ctypes.c_size_t, # n_elements ctypes.c_double, # abs_error_bound ] -_lib.fsz_decompress_hostptr_f64.restype = ctypes.c_int +_lib.szgpu_decompress_hostptr_f64.restype = ctypes.c_int -FSZ_STATUS_OK = 0 +SZGPU_STATUS_OK = 0 def _status_text(status): """Human readable text for a library status code.""" - text = _lib.fsz_status_string(ctypes.c_int(int(status))) + text = _lib.szgpu_status_string(ctypes.c_int(int(status))) if text is None: return "unknown status" return text.decode("utf-8", "replace") @@ -190,7 +190,7 @@ def _status_text(status): def _check(status, what): """Raise RuntimeError unless ``status`` is the library's OK code.""" - if int(status) != FSZ_STATUS_OK: + if int(status) != SZGPU_STATUS_OK: raise RuntimeError( "%s failed: %s (status %d)" % (what, _status_text(status), int(status)) ) @@ -209,7 +209,7 @@ def _as_input_array(arr): if arr.dtype != np.float32 and arr.dtype != np.float64: raise TypeError( - "FSZ compresses float32 and float64 data, but the input array has " + "SZ-GPU compresses float32 and float64 data, but the input array has " "dtype %s. Convert it explicitly, for example " "arr.astype(np.float32), so the precision loss is your choice and " "not a silent cast." % arr.dtype @@ -221,7 +221,7 @@ def _as_input_array(arr): ) if arr.ndim > MAX_HEADER_DIMS: raise ValueError( - "the input array has %d dimensions, and the FSZ container records at " + "the input array has %d dimensions, and the SZ-GPU container records at " "most %d. Reshape it first, for example " "arr.reshape(arr.shape[0], arr.shape[1], -1), and reshape the result " "of decompress() back afterwards." % (arr.ndim, MAX_HEADER_DIMS) @@ -335,7 +335,7 @@ def _parse_header(data): """ if data.size < HEADER_SIZE: raise ValueError( - "this is not an FSZ container: %d bytes is shorter than the %d byte " + "this is not an SZ-GPU container: %d bytes is shorter than the %d byte " "header." % (data.size, HEADER_SIZE) ) @@ -344,7 +344,7 @@ def _parse_header(data): if magic != MAGIC: raise ValueError( - "this is not an FSZ container: expected the magic %r, found %r. A " + "this is not an SZ-GPU container: expected the magic %r, found %r. A " "bare bitstream has no header, so decompress it with " "decompress_raw(payload, n, abs_eb)." % (MAGIC, magic) ) @@ -396,19 +396,19 @@ def library_version(): """Return the version string reported by the shared library. >>> library_version() - '1.0.0' + '1.0.1' """ - return _lib.fsz_version().decode("utf-8", "replace") + return _lib.szgpu_version().decode("utf-8", "replace") def num_tiles(n): """Return the number of tiles the library splits ``n`` elements into.""" - return int(_lib.fsz_num_tiles(ctypes.c_size_t(int(n)))) + return int(_lib.szgpu_num_tiles(ctypes.c_size_t(int(n)))) def max_compressed_bytes(n): """Return the worst case compressed size in bytes for ``n`` elements.""" - return int(_lib.fsz_max_compressed_bytes(ctypes.c_size_t(int(n)))) + return int(_lib.szgpu_max_compressed_bytes(ctypes.c_size_t(int(n)))) def compress_raw(arr, abs_eb): @@ -449,23 +449,23 @@ def compress_raw(arr, abs_eb): written = ctypes.c_size_t(0) if f64: - status = _lib.fsz_compress_hostptr_f64( + status = _lib.szgpu_compress_hostptr_f64( arr.ctypes.data_as(_c_double_p), buf.ctypes.data_as(_c_ubyte_p), ctypes.c_size_t(n), ctypes.c_double(float(abs_eb)), ctypes.byref(written), ) - _check(status, "fsz_compress_hostptr_f64") + _check(status, "szgpu_compress_hostptr_f64") else: - status = _lib.fsz_compress_hostptr( + status = _lib.szgpu_compress_hostptr( arr.ctypes.data_as(_c_float_p), buf.ctypes.data_as(_c_ubyte_p), ctypes.c_size_t(n), ctypes.c_float(float(abs_eb)), ctypes.byref(written), ) - _check(status, "fsz_compress_hostptr") + _check(status, "szgpu_compress_hostptr") return buf[: written.value].tobytes() @@ -505,7 +505,7 @@ def decompress_raw(payload, n, abs_eb, dtype=np.float32): dtype = np.dtype(dtype) if dtype != np.float32 and dtype != np.float64: raise TypeError( - "FSZ reconstructs float32 and float64 data, but dtype is %s." % dtype + "SZ-GPU reconstructs float32 and float64 data, but dtype is %s." % dtype ) f64 = dtype == np.float64 @@ -518,23 +518,23 @@ def decompress_raw(payload, n, abs_eb, dtype=np.float32): out = np.empty(n, dtype=dtype) if f64: - status = _lib.fsz_decompress_hostptr_f64( + status = _lib.szgpu_decompress_hostptr_f64( out.ctypes.data_as(_c_double_p), data.ctypes.data_as(_c_ubyte_p), ctypes.c_size_t(int(data.size)), ctypes.c_size_t(n), ctypes.c_double(float(abs_eb)), ) - _check(status, "fsz_decompress_hostptr_f64") + _check(status, "szgpu_decompress_hostptr_f64") else: - status = _lib.fsz_decompress_hostptr( + status = _lib.szgpu_decompress_hostptr( out.ctypes.data_as(_c_float_p), data.ctypes.data_as(_c_ubyte_p), ctypes.c_size_t(int(data.size)), ctypes.c_size_t(n), ctypes.c_float(float(abs_eb)), ) - _check(status, "fsz_decompress_hostptr") + _check(status, "szgpu_decompress_hostptr") return out @@ -544,7 +544,7 @@ def compress(arr, eb, mode="rel"): The blob is a 56 byte header followed by the compressed bitstream. The header records the shape, the element type, the absolute error bound and the bitstream length, which is everything :func:`decompress` needs, and it - is the same container the ``fsz`` command line tool reads. + is the same container the ``szgpu`` command line tool reads. Args: arr: float32 or float64 numpy array of 1 to 3 dimensions. A @@ -615,10 +615,10 @@ def decompress(blob): def compress_file(src, dst, eb, mode="rel", dims=None, dtype=np.float32): - """Compress a raw float32 or float64 file into an FSZ container on disk. + """Compress a raw float32 or float64 file into an SZ-GPU container on disk. ``src`` is read as a flat stream of values in ``dtype``, the layout that - raw binary dumps and the ``fsz`` command line tool use. Pass ``dims`` to + raw binary dumps and the ``szgpu`` command line tool use. Pass ``dims`` to record the shape in the header so :func:`decompress_file` hands the array back in its original form. @@ -648,7 +648,7 @@ def compress_file(src, dst, eb, mode="rel", dims=None, dtype=np.float32): dtype = np.dtype(dtype) if dtype != np.float32 and dtype != np.float64: raise TypeError( - "FSZ compresses float32 and float64 data, but dtype is %s." % dtype + "SZ-GPU compresses float32 and float64 data, but dtype is %s." % dtype ) arr = np.fromfile(src, dtype=dtype) @@ -683,11 +683,11 @@ def compress_file(src, dst, eb, mode="rel", dims=None, dtype=np.float32): def decompress_file(src, dst=None): - """Reconstruct the array in an FSZ container file. + """Reconstruct the array in an SZ-GPU container file. Args: src: path of the container written by :func:`compress_file` or by the - ``fsz`` command line tool. + ``szgpu`` command line tool. dst: optional path for the reconstructed data, written raw in C order, in the element type the container records. @@ -696,7 +696,7 @@ def decompress_file(src, dst=None): the container, float32 or float64. Raises: - ValueError: if the file is not a readable FSZ container. + ValueError: if the file is not a readable SZ-GPU container. RuntimeError: if the library reports a failure. """ with open(src, "rb") as f: diff --git a/python/tests/test_fsz.py b/python/tests/test_szgpu.py similarity index 80% rename from python/tests/test_fsz.py rename to python/tests/test_szgpu.py index c1b0940..787a074 100644 --- a/python/tests/test_fsz.py +++ b/python/tests/test_szgpu.py @@ -1,10 +1,10 @@ -"""Tests for the FSZ Python bindings. +"""Tests for the SZ-GPU Python bindings. A plain script, so it needs nothing beyond numpy and the bindings themselves. Every case prints [PASS] or [FAIL]; the exit status is 0 only when all of them pass. Run it from the repository root: - python python/tests/test_fsz.py + python python/tests/test_szgpu.py """ import os @@ -17,10 +17,10 @@ import numpy as np try: - import fsz + import szgpu except ImportError: sys.path.insert(0, os.path.dirname(os.path.dirname(os.path.abspath(__file__)))) - import fsz + import szgpu SHAPE = (97, 113, 131) @@ -64,12 +64,12 @@ def relative_bound_f64(arr, eb_rel): def header_flags(blob): """The flags word of a container header.""" - return struct.unpack(fsz.HEADER_FORMAT, blob[:fsz.HEADER_SIZE])[3] + return struct.unpack(szgpu.HEADER_FORMAT, blob[:szgpu.HEADER_SIZE])[3] def header_eb_f64(blob): """The float64 error bound a container records, read at full precision.""" - return struct.unpack(" 1.0, "compression ratio %.3f is not above 1" % info["cr"] - back = fsz.decompress_file(cmp_path, out_path) + back = szgpu.decompress_file(cmp_path, out_path) assert back.shape == SHAPE, \ "decompress_file returned shape %s, expected %s" % (back.shape, SHAPE) @@ -246,33 +246,33 @@ def test_file_cycle(): @case("float64 compress_file / decompress_file cycle on a .d64") def test_f64_file_cycle(): original = field_f64() - workdir = tempfile.mkdtemp(prefix="fsz_test_") + workdir = tempfile.mkdtemp(prefix="szgpu_test_") try: raw_path = os.path.join(workdir, "field.d64") - cmp_path = os.path.join(workdir, "field.fsz") + cmp_path = os.path.join(workdir, "field.szgpu") out_path = os.path.join(workdir, "field.out.d64") original.tofile(raw_path) assert os.path.getsize(raw_path) == original.size * 8, \ "the raw file is not 8 bytes per value" - info = fsz.compress_file(raw_path, cmp_path, ABS_EB, mode="abs", + info = szgpu.compress_file(raw_path, cmp_path, ABS_EB, mode="abs", dims=SHAPE, dtype=np.float64) assert info["n"] == original.size, \ "compress_file counted %d elements, expected %d" % (info["n"], original.size) - assert info["cmp_size"] == os.path.getsize(cmp_path) - fsz.HEADER_SIZE, \ + assert info["cmp_size"] == os.path.getsize(cmp_path) - szgpu.HEADER_SIZE, \ "cmp_size %d disagrees with the file size" % info["cmp_size"] expected_cr = float(original.nbytes) / float(info["cmp_size"]) assert abs(info["cr"] - expected_cr) < 1e-9, \ "cr %.6f is not the float64 ratio %.6f" % (info["cr"], expected_cr) with open(cmp_path, "rb") as f: - head = f.read(fsz.HEADER_SIZE) - assert header_flags(head) & fsz.FLAG_F64, \ + head = f.read(szgpu.HEADER_SIZE) + assert header_flags(head) & szgpu.FLAG_F64, \ "the container on disk does not flag float64 data" - back = fsz.decompress_file(cmp_path, out_path) + back = szgpu.decompress_file(cmp_path, out_path) assert back.dtype == np.float64, \ "decompress_file returned dtype %s, expected float64" % back.dtype assert back.shape == SHAPE, \ @@ -298,17 +298,17 @@ def test_mixed_dtypes(): f32 = field() f64 = field_f64() - blob32 = fsz.compress(f32, ABS_EB, mode="abs") - blob64 = fsz.compress(f64, ABS_EB, mode="abs") + blob32 = szgpu.compress(f32, ABS_EB, mode="abs") + blob64 = szgpu.compress(f64, ABS_EB, mode="abs") - assert header_flags(blob32) & fsz.FLAG_F64 == 0, \ + assert header_flags(blob32) & szgpu.FLAG_F64 == 0, \ "a float32 container has the float64 flag set" - assert header_flags(blob64) & fsz.FLAG_F64 == fsz.FLAG_F64, \ + assert header_flags(blob64) & szgpu.FLAG_F64 == szgpu.FLAG_F64, \ "a float64 container does not have the float64 flag set" assert blob32 != blob64, "the two element types produced the same blob" - back32 = fsz.decompress(blob32) - back64 = fsz.decompress(blob64) + back32 = szgpu.decompress(blob32) + back64 = szgpu.decompress(blob64) assert back32.dtype == np.float32, "float32 blob came back as %s" % back32.dtype assert back64.dtype == np.float64, "float64 blob came back as %s" % back64.dtype @@ -322,8 +322,8 @@ def test_mixed_dtypes(): assert max_err <= TOLERANCE * ABS_EB, \ "%s max error %.6e exceeds %.6e" % (name, max_err, TOLERANCE * ABS_EB) - payload64 = blob64[fsz.HEADER_SIZE:] - from_raw = fsz.decompress_raw(payload64, f64.size, ABS_EB, dtype=np.float64) + payload64 = blob64[szgpu.HEADER_SIZE:] + from_raw = szgpu.decompress_raw(payload64, f64.size, ABS_EB, dtype=np.float64) assert from_raw.dtype == np.float64, "decompress_raw returned %s" % from_raw.dtype assert from_raw.tobytes() == back64.ravel().tobytes(), \ "the raw and container paths disagree on the float64 stream" @@ -335,16 +335,16 @@ def test_mixed_dtypes(): @case("compress_raw / decompress_raw agree with the container path") def test_raw_consistency(): original = field() - blob = fsz.compress(original, ABS_EB, mode="abs") - payload = fsz.compress_raw(original, np.float32(ABS_EB)) + blob = szgpu.compress(original, ABS_EB, mode="abs") + payload = szgpu.compress_raw(original, np.float32(ABS_EB)) - assert payload == blob[fsz.HEADER_SIZE:], \ + assert payload == blob[szgpu.HEADER_SIZE:], \ ("the bare bitstream is %d bytes and the container payload is %d; " "they must be identical" - % (len(payload), len(blob) - fsz.HEADER_SIZE)) + % (len(payload), len(blob) - szgpu.HEADER_SIZE)) - from_raw = fsz.decompress_raw(payload, original.size, np.float32(ABS_EB)) - from_blob = fsz.decompress(blob).ravel() + from_raw = szgpu.decompress_raw(payload, original.size, np.float32(ABS_EB)) + from_blob = szgpu.decompress(blob).ravel() assert from_raw.shape == from_blob.shape, \ "shapes differ: %s and %s" % (from_raw.shape, from_blob.shape) @@ -355,49 +355,49 @@ def test_raw_consistency(): @case("error cases are reported, not crashed on") def test_error_cases(): - exc = expect_raises(TypeError, fsz.compress, + exc = expect_raises(TypeError, szgpu.compress, np.zeros((8, 8), dtype=np.int32), 1e-3) print(" int32 input -> TypeError: %s" % str(exc).split(".")[0]) - exc = expect_raises(ValueError, fsz.compress, + exc = expect_raises(ValueError, szgpu.compress, np.zeros((2, 3, 4, 5), dtype=np.float32), 1e-3) print(" 4 axes -> ValueError: %s" % str(exc).split(".")[0]) - exc = expect_raises(ValueError, fsz.compress, + exc = expect_raises(ValueError, szgpu.compress, np.full((32, 32), 2.5, dtype=np.float32), 1e-3, "rel") print(" constant + rel -> ValueError: %s" % str(exc).split(",")[0]) - exc = expect_raises(ValueError, fsz.compress, + exc = expect_raises(ValueError, szgpu.compress, np.zeros((8, 8), dtype=np.float32), 0.0, "abs") print(" eb 0 -> ValueError: %s" % str(exc).split(",")[0]) - exc = expect_raises(ValueError, fsz.compress, + exc = expect_raises(ValueError, szgpu.compress, np.zeros((8, 8), dtype=np.float32), 1e-3, "sideways") print(" unknown mode -> ValueError: %s" % str(exc).split(",")[0]) - blob = fsz.compress(np.asarray(np.arange(4096), dtype=np.float32), 1e-2, "abs") + blob = szgpu.compress(np.asarray(np.arange(4096), dtype=np.float32), 1e-2, "abs") - exc = expect_raises((ValueError, RuntimeError), fsz.decompress, - blob[:fsz.HEADER_SIZE - 20]) + exc = expect_raises((ValueError, RuntimeError), szgpu.decompress, + blob[:szgpu.HEADER_SIZE - 20]) print(" stub blob -> %s" % type(exc).__name__) - exc = expect_raises((ValueError, RuntimeError), fsz.decompress, - blob[:fsz.HEADER_SIZE + (len(blob) - fsz.HEADER_SIZE) // 2]) + exc = expect_raises((ValueError, RuntimeError), szgpu.decompress, + blob[:szgpu.HEADER_SIZE + (len(blob) - szgpu.HEADER_SIZE) // 2]) print(" cut payload -> %s: %s" % (type(exc).__name__, str(exc).split(";")[0])) - exc = expect_raises((ValueError, RuntimeError), fsz.decompress, + exc = expect_raises((ValueError, RuntimeError), szgpu.decompress, b"NOPE" + blob[4:]) print(" bad magic -> %s" % type(exc).__name__) - assert fsz.decompress(blob).shape == (4096,), "a good blob stopped working" + assert szgpu.decompress(blob).shape == (4096,), "a good blob stopped working" @case("compression ratio on smooth data clears 5x") def test_compression_ratio(): original = field() for mode, eb in (("abs", ABS_EB), ("rel", REL_EB)): - blob = fsz.compress(original, eb, mode=mode) - payload_size = len(blob) - fsz.HEADER_SIZE + blob = szgpu.compress(original, eb, mode=mode) + payload_size = len(blob) - szgpu.HEADER_SIZE cr = (original.size * 4.0) / payload_size print(" %s %g: %d bytes -> %d bytes, cr = %.3f" % (mode, eb, original.nbytes, payload_size, cr)) @@ -406,8 +406,8 @@ def test_compression_ratio(): def main(): - print("FSZ Python bindings %s, library %s" - % (fsz.__version__, fsz.library_version())) + print("SZ-GPU Python bindings %s, library %s" + % (szgpu.__version__, szgpu.library_version())) print("%d cases" % len(_CASES)) print() diff --git a/src/fsz.cu b/src/fsz.cu deleted file mode 100644 index 5e9fed2..0000000 --- a/src/fsz.cu +++ /dev/null @@ -1,442 +0,0 @@ -#include "fsz/fsz.h" - -#include "fsz_internal.cuh" -#include "fsz_kernels.cuh" - -#include -#include -#include -#include -#include - -using fsz_detail::fsz_compress_kernel_T; -using fsz_detail::fsz_decompress_kernel_T; - -struct fsz_workspace { - unsigned int* d_workspace; - unsigned int* cmpOffset; - unsigned int* locOffset; - int* flag; - unsigned int* h_cmpSize; - int gsize_max; - int gsize_sm_max; - int cmpOffSize; - size_t max_n_elements; -}; - -static int fsz_grid_size(size_t numTiles) { - return (int)((numTiles + FSZ_TBLOCK_SZ * FSZ_TILES_PER_THR - 1) - / (FSZ_TBLOCK_SZ * FSZ_TILES_PER_THR)); -} - -static int fsz_grid_size_sm(size_t numTiles) { - return (int)((numTiles + FSZ_TBLOCK_SZ * FSZ_TILES_PER_THR_SM - 1) - / (FSZ_TBLOCK_SZ * FSZ_TILES_PER_THR_SM)); -} - -static void fsz_hostptr_release(void* d_a, void* d_b, fsz_workspace_t* ws) { - if (d_a) cudaFree(d_a); - if (d_b) cudaFree(d_b); - if (ws) fsz_workspace_destroy(ws); -} - -extern "C" const char* fsz_version(void) { - return FSZ_VERSION_STRING; -} - -extern "C" const char* fsz_status_string(fsz_status_t s) { - switch (s) { - case FSZ_STATUS_OK: return "OK"; - case FSZ_STATUS_INVALID_ARGUMENT: return "invalid argument"; - case FSZ_STATUS_OUT_OF_MEMORY: return "out of memory"; - case FSZ_STATUS_WORKSPACE_TOO_SMALL: return "workspace too small"; - case FSZ_STATUS_CUDA_ERROR: return "CUDA error"; - case FSZ_STATUS_TOO_LARGE: return "compressed stream too large"; - } - return "unknown"; -} - -extern "C" size_t fsz_num_tiles(size_t n) { - return (n + FSZ_TILE_SIZE - 1) / FSZ_TILE_SIZE; -} - -extern "C" size_t fsz_max_compressed_bytes(size_t n) { - return fsz_num_tiles(n) * FSZ_MAX_TILE_BYTES; -} - -extern "C" fsz_compress_result_t fsz_make_result(size_t n, size_t cmp_size) { - size_t numTiles = fsz_num_tiles(n); - fsz_compress_result_t r; - r.cmp_size = cmp_size; - r.num_tiles = numTiles; - r.data_offset = numTiles * FSZ_BLKS_PER_TILE; - return r; -} - -extern "C" fsz_status_t fsz_workspace_create(fsz_workspace_t** out_ws, size_t max_n) { - if (!out_ws) return FSZ_STATUS_INVALID_ARGUMENT; - *out_ws = nullptr; - if (max_n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - - fsz_workspace* w = new(std::nothrow) fsz_workspace(); - if (!w) return FSZ_STATUS_OUT_OF_MEMORY; - - size_t numTiles = fsz_num_tiles(max_n); - w->max_n_elements = max_n; - w->gsize_max = fsz_grid_size(numTiles); - w->gsize_sm_max = fsz_grid_size_sm(numTiles); - - int maxgs = w->gsize_max > w->gsize_sm_max ? w->gsize_max : w->gsize_sm_max; - w->cmpOffSize = maxgs + 1; - - cudaError_t e = cudaMalloc(&w->d_workspace, sizeof(unsigned int) * w->cmpOffSize * 3); - if (e != cudaSuccess) { delete w; return FSZ_STATUS_CUDA_ERROR; } - - w->cmpOffset = w->d_workspace; - w->locOffset = w->d_workspace + w->cmpOffSize; - w->flag = (int*)(w->d_workspace + 2 * w->cmpOffSize); - - e = cudaHostAlloc(&w->h_cmpSize, sizeof(unsigned int), cudaHostAllocDefault); - if (e != cudaSuccess) { - cudaFree(w->d_workspace); - delete w; - return FSZ_STATUS_CUDA_ERROR; - } - - *out_ws = w; - return FSZ_STATUS_OK; -} - -extern "C" void fsz_workspace_destroy(fsz_workspace_t* ws) { - if (!ws) return; - if (ws->d_workspace) cudaFree(ws->d_workspace); - if (ws->h_cmpSize) cudaFreeHost(ws->h_cmpSize); - delete ws; -} - -extern "C" size_t fsz_workspace_capacity(const fsz_workspace_t* ws) { - return ws ? ws->max_n_elements : 0; -} - -template -static fsz_status_t fsz_compress_with_ws( - const FpT* d_in, - unsigned char* d_cmp, - size_t n, - FpT eb, - fsz_workspace_t* ws, - cudaStream_t stream, - fsz_compress_result_t* out_result) -{ - if (!d_in || !d_cmp || !ws || !out_result) return FSZ_STATUS_INVALID_ARGUMENT; - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - if (!(eb > (FpT)0.0f)) return FSZ_STATUS_INVALID_ARGUMENT; - if (n > ws->max_n_elements) return FSZ_STATUS_WORKSPACE_TOO_SMALL; - - size_t numTiles = fsz_num_tiles(n); - size_t numBlocks = numTiles * FSZ_BLKS_PER_TILE; - size_t dataOfs = numBlocks; - - bool use_small = (n <= FSZ_SMALL_THRESHOLD); - int gs = use_small ? fsz_grid_size_sm(numTiles) : fsz_grid_size(numTiles); - - cudaError_t e = cudaMemsetAsync(ws->flag, 0, sizeof(int) * (gs + 1), stream); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - cudaGetLastError(); - - if (use_small) { - fsz_compress_kernel_T<<>>( - d_in, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, - eb, n, numTiles, dataOfs); - } else { - fsz_compress_kernel_T<<>>( - d_in, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, - eb, n, numTiles, dataOfs); - } - if (cudaGetLastError() != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - const unsigned int* d_total = (gs == 1) ? (ws->locOffset + 1) - : (ws->cmpOffset + gs); - e = cudaMemcpyAsync(ws->h_cmpSize, d_total, - sizeof(unsigned int), cudaMemcpyDeviceToHost, stream); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - e = cudaStreamSynchronize(stream); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - if ((unsigned long long)numTiles * (FSZ_MAX_TILE_BYTES - FSZ_BLKS_PER_TILE) - > 0xffffffffull) { - std::vector parts((size_t)gs); - e = cudaMemcpy(parts.data(), ws->locOffset + 1, - sizeof(unsigned int) * (size_t)gs, cudaMemcpyDeviceToHost); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - unsigned long long total = 0; - for (int i = 0; i < gs; ++i) total += parts[(size_t)i]; - if (total > 0xffffffffull) return FSZ_STATUS_TOO_LARGE; - } - - out_result->cmp_size = dataOfs + *ws->h_cmpSize; - out_result->num_tiles = numTiles; - out_result->data_offset = dataOfs; - return FSZ_STATUS_OK; -} - -template -static fsz_status_t fsz_compress_core( - const FpT* d_in, - unsigned char* d_cmp, - size_t n, - FpT eb, - fsz_workspace_t* ws, - cudaStream_t stream, - fsz_compress_result_t* out_result) -{ - if (ws) return fsz_compress_with_ws(d_in, d_cmp, n, eb, ws, stream, out_result); - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - - fsz_workspace_t* tmp = nullptr; - fsz_status_t s = fsz_workspace_create(&tmp, n); - if (s != FSZ_STATUS_OK) return s; - s = fsz_compress_with_ws(d_in, d_cmp, n, eb, tmp, stream, out_result); - fsz_workspace_destroy(tmp); - return s; -} - -extern "C" fsz_status_t fsz_compress( - const float* d_in, unsigned char* d_cmp, size_t n, float eb, - fsz_workspace_t* ws, cudaStream_t stream, fsz_compress_result_t* out_result) -{ - return fsz_compress_core(d_in, d_cmp, n, eb, ws, stream, out_result); -} - -extern "C" fsz_status_t fsz_compress_f64( - const double* d_in, unsigned char* d_cmp, size_t n, double eb, - fsz_workspace_t* ws, cudaStream_t stream, fsz_compress_result_t* out_result) -{ - return fsz_compress_core(d_in, d_cmp, n, eb, ws, stream, out_result); -} - -template -static fsz_status_t fsz_decompress_with_ws( - FpT* d_out, - const unsigned char* d_cmp, - size_t n, - FpT eb, - const fsz_compress_result_t* result, - fsz_workspace_t* ws, - cudaStream_t stream) -{ - if (!d_out || !d_cmp || !ws || !result) return FSZ_STATUS_INVALID_ARGUMENT; - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - if (!(eb > (FpT)0.0f)) return FSZ_STATUS_INVALID_ARGUMENT; - if (n > ws->max_n_elements) return FSZ_STATUS_WORKSPACE_TOO_SMALL; - - const size_t expect_tiles = fsz_num_tiles(n); - if (result->num_tiles != expect_tiles) return FSZ_STATUS_INVALID_ARGUMENT; - if (result->data_offset != expect_tiles * FSZ_BLKS_PER_TILE) - return FSZ_STATUS_INVALID_ARGUMENT; - - bool use_small = (n <= FSZ_SMALL_THRESHOLD); - int gs = use_small ? fsz_grid_size_sm(result->num_tiles) - : fsz_grid_size(result->num_tiles); - - cudaError_t e = cudaMemsetAsync(ws->flag, 0, sizeof(int) * (gs + 1), stream); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - size_t origSize = n * sizeof(FpT); - bool high_cr = (result->cmp_size > 0) && (origSize / result->cmp_size > 100); - if (high_cr) { - e = cudaMemsetAsync(d_out, 0, origSize, stream); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - } - - cudaGetLastError(); - - if (use_small) { - fsz_decompress_kernel_T<<>>( - d_out, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, - eb, n, result->num_tiles, result->data_offset, high_cr); - } else { - fsz_decompress_kernel_T<<>>( - d_out, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, - eb, n, result->num_tiles, result->data_offset, high_cr); - } - if (cudaGetLastError() != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - return FSZ_STATUS_OK; -} - -template -static fsz_status_t fsz_decompress_core( - FpT* d_out, - const unsigned char* d_cmp, - size_t n, - FpT eb, - const fsz_compress_result_t* result, - fsz_workspace_t* ws, - cudaStream_t stream) -{ - if (ws) return fsz_decompress_with_ws(d_out, d_cmp, n, eb, result, ws, stream); - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - - fsz_workspace_t* tmp = nullptr; - fsz_status_t s = fsz_workspace_create(&tmp, n); - if (s != FSZ_STATUS_OK) return s; - s = fsz_decompress_with_ws(d_out, d_cmp, n, eb, result, tmp, stream); - cudaError_t e = cudaStreamSynchronize(stream); - fsz_workspace_destroy(tmp); - if (s == FSZ_STATUS_OK && e != cudaSuccess) s = FSZ_STATUS_CUDA_ERROR; - return s; -} - -extern "C" fsz_status_t fsz_decompress( - float* d_out, const unsigned char* d_cmp, size_t n, float eb, - const fsz_compress_result_t* result, fsz_workspace_t* ws, cudaStream_t stream) -{ - return fsz_decompress_core(d_out, d_cmp, n, eb, result, ws, stream); -} - -extern "C" fsz_status_t fsz_decompress_f64( - double* d_out, const unsigned char* d_cmp, size_t n, double eb, - const fsz_compress_result_t* result, fsz_workspace_t* ws, cudaStream_t stream) -{ - return fsz_decompress_core(d_out, d_cmp, n, eb, result, ws, stream); -} - -template -static fsz_status_t fsz_compress_hostptr_core( - const FpT* in, - unsigned char* cmp, - size_t n, - FpT eb, - size_t* out_cmp_size) -{ - if (!in || !cmp || !out_cmp_size) return FSZ_STATUS_INVALID_ARGUMENT; - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - if (!(eb > (FpT)0.0f)) return FSZ_STATUS_INVALID_ARGUMENT; - - *out_cmp_size = 0; - - FpT* d_in = nullptr; - unsigned char* d_cmp = nullptr; - fsz_workspace_t* ws = nullptr; - - cudaError_t e = cudaMalloc(&d_in, n * sizeof(FpT)); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - e = cudaMalloc(&d_cmp, fsz_max_compressed_bytes(n)); - if (e != cudaSuccess) { - fsz_hostptr_release(d_in, nullptr, nullptr); - return FSZ_STATUS_CUDA_ERROR; - } - - fsz_status_t s = fsz_workspace_create(&ws, n); - if (s != FSZ_STATUS_OK) { - fsz_hostptr_release(d_in, d_cmp, nullptr); - return s; - } - - e = cudaMemcpy(d_in, in, n * sizeof(FpT), cudaMemcpyHostToDevice); - if (e != cudaSuccess) { - fsz_hostptr_release(d_in, d_cmp, ws); - return FSZ_STATUS_CUDA_ERROR; - } - - fsz_compress_result_t r; - s = fsz_compress_core(d_in, d_cmp, n, eb, ws, 0, &r); - if (s != FSZ_STATUS_OK) { - fsz_hostptr_release(d_in, d_cmp, ws); - return s; - } - - e = cudaMemcpy(cmp, d_cmp, r.cmp_size, cudaMemcpyDeviceToHost); - fsz_hostptr_release(d_in, d_cmp, ws); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - *out_cmp_size = r.cmp_size; - return FSZ_STATUS_OK; -} - -extern "C" fsz_status_t fsz_compress_hostptr( - const float* in, unsigned char* cmp, size_t n, float eb, size_t* out_cmp_size) -{ - return fsz_compress_hostptr_core(in, cmp, n, eb, out_cmp_size); -} - -extern "C" fsz_status_t fsz_compress_hostptr_f64( - const double* in, unsigned char* cmp, size_t n, double eb, size_t* out_cmp_size) -{ - return fsz_compress_hostptr_core(in, cmp, n, eb, out_cmp_size); -} - -template -static fsz_status_t fsz_decompress_hostptr_core( - FpT* out, - const unsigned char* cmp, - size_t cmp_size, - size_t n, - FpT eb) -{ - if (!out || !cmp) return FSZ_STATUS_INVALID_ARGUMENT; - if (n == 0) return FSZ_STATUS_INVALID_ARGUMENT; - if (!(eb > (FpT)0.0f)) return FSZ_STATUS_INVALID_ARGUMENT; - - size_t header = fsz_num_tiles(n) * FSZ_BLKS_PER_TILE; - if (cmp_size == 0 || cmp_size < header) return FSZ_STATUS_INVALID_ARGUMENT; - - FpT* d_out = nullptr; - unsigned char* d_cmp = nullptr; - fsz_workspace_t* ws = nullptr; - - cudaError_t e = cudaMalloc(&d_out, n * sizeof(FpT)); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - e = cudaMalloc(&d_cmp, cmp_size); - if (e != cudaSuccess) { - fsz_hostptr_release(d_out, nullptr, nullptr); - return FSZ_STATUS_CUDA_ERROR; - } - - fsz_status_t s = fsz_workspace_create(&ws, n); - if (s != FSZ_STATUS_OK) { - fsz_hostptr_release(d_out, d_cmp, nullptr); - return s; - } - - e = cudaMemcpy(d_cmp, cmp, cmp_size, cudaMemcpyHostToDevice); - if (e != cudaSuccess) { - fsz_hostptr_release(d_out, d_cmp, ws); - return FSZ_STATUS_CUDA_ERROR; - } - - fsz_compress_result_t r = fsz_make_result(n, cmp_size); - s = fsz_decompress_core(d_out, d_cmp, n, eb, &r, ws, 0); - if (s != FSZ_STATUS_OK) { - fsz_hostptr_release(d_out, d_cmp, ws); - return s; - } - - e = cudaStreamSynchronize(0); - if (e != cudaSuccess) { - fsz_hostptr_release(d_out, d_cmp, ws); - return FSZ_STATUS_CUDA_ERROR; - } - - e = cudaMemcpy(out, d_out, n * sizeof(FpT), cudaMemcpyDeviceToHost); - fsz_hostptr_release(d_out, d_cmp, ws); - if (e != cudaSuccess) return FSZ_STATUS_CUDA_ERROR; - - return FSZ_STATUS_OK; -} - -extern "C" fsz_status_t fsz_decompress_hostptr( - float* out, const unsigned char* cmp, size_t cmp_size, size_t n, float eb) -{ - return fsz_decompress_hostptr_core(out, cmp, cmp_size, n, eb); -} - -extern "C" fsz_status_t fsz_decompress_hostptr_f64( - double* out, const unsigned char* cmp, size_t cmp_size, size_t n, double eb) -{ - return fsz_decompress_hostptr_core(out, cmp, cmp_size, n, eb); -} diff --git a/src/szgpu.cu b/src/szgpu.cu new file mode 100644 index 0000000..75c5629 --- /dev/null +++ b/src/szgpu.cu @@ -0,0 +1,442 @@ +#include "szgpu/szgpu.h" + +#include "szgpu_internal.cuh" +#include "szgpu_kernels.cuh" + +#include +#include +#include +#include +#include + +using szgpu_detail::szgpu_compress_kernel_T; +using szgpu_detail::szgpu_decompress_kernel_T; + +struct szgpu_workspace { + unsigned int* d_workspace; + unsigned int* cmpOffset; + unsigned int* locOffset; + int* flag; + unsigned int* h_cmpSize; + int gsize_max; + int gsize_sm_max; + int cmpOffSize; + size_t max_n_elements; +}; + +static int szgpu_grid_size(size_t numTiles) { + return (int)((numTiles + SZGPU_TBLOCK_SZ * SZGPU_TILES_PER_THR - 1) + / (SZGPU_TBLOCK_SZ * SZGPU_TILES_PER_THR)); +} + +static int szgpu_grid_size_sm(size_t numTiles) { + return (int)((numTiles + SZGPU_TBLOCK_SZ * SZGPU_TILES_PER_THR_SM - 1) + / (SZGPU_TBLOCK_SZ * SZGPU_TILES_PER_THR_SM)); +} + +static void szgpu_hostptr_release(void* d_a, void* d_b, szgpu_workspace_t* ws) { + if (d_a) cudaFree(d_a); + if (d_b) cudaFree(d_b); + if (ws) szgpu_workspace_destroy(ws); +} + +extern "C" const char* szgpu_version(void) { + return SZGPU_VERSION_STRING; +} + +extern "C" const char* szgpu_status_string(szgpu_status_t s) { + switch (s) { + case SZGPU_STATUS_OK: return "OK"; + case SZGPU_STATUS_INVALID_ARGUMENT: return "invalid argument"; + case SZGPU_STATUS_OUT_OF_MEMORY: return "out of memory"; + case SZGPU_STATUS_WORKSPACE_TOO_SMALL: return "workspace too small"; + case SZGPU_STATUS_CUDA_ERROR: return "CUDA error"; + case SZGPU_STATUS_TOO_LARGE: return "compressed stream too large"; + } + return "unknown"; +} + +extern "C" size_t szgpu_num_tiles(size_t n) { + return (n + SZGPU_TILE_SIZE - 1) / SZGPU_TILE_SIZE; +} + +extern "C" size_t szgpu_max_compressed_bytes(size_t n) { + return szgpu_num_tiles(n) * SZGPU_MAX_TILE_BYTES; +} + +extern "C" szgpu_compress_result_t szgpu_make_result(size_t n, size_t cmp_size) { + size_t numTiles = szgpu_num_tiles(n); + szgpu_compress_result_t r; + r.cmp_size = cmp_size; + r.num_tiles = numTiles; + r.data_offset = numTiles * SZGPU_BLKS_PER_TILE; + return r; +} + +extern "C" szgpu_status_t szgpu_workspace_create(szgpu_workspace_t** out_ws, size_t max_n) { + if (!out_ws) return SZGPU_STATUS_INVALID_ARGUMENT; + *out_ws = nullptr; + if (max_n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + + szgpu_workspace* w = new(std::nothrow) szgpu_workspace(); + if (!w) return SZGPU_STATUS_OUT_OF_MEMORY; + + size_t numTiles = szgpu_num_tiles(max_n); + w->max_n_elements = max_n; + w->gsize_max = szgpu_grid_size(numTiles); + w->gsize_sm_max = szgpu_grid_size_sm(numTiles); + + int maxgs = w->gsize_max > w->gsize_sm_max ? w->gsize_max : w->gsize_sm_max; + w->cmpOffSize = maxgs + 1; + + cudaError_t e = cudaMalloc(&w->d_workspace, sizeof(unsigned int) * w->cmpOffSize * 3); + if (e != cudaSuccess) { delete w; return SZGPU_STATUS_CUDA_ERROR; } + + w->cmpOffset = w->d_workspace; + w->locOffset = w->d_workspace + w->cmpOffSize; + w->flag = (int*)(w->d_workspace + 2 * w->cmpOffSize); + + e = cudaHostAlloc(&w->h_cmpSize, sizeof(unsigned int), cudaHostAllocDefault); + if (e != cudaSuccess) { + cudaFree(w->d_workspace); + delete w; + return SZGPU_STATUS_CUDA_ERROR; + } + + *out_ws = w; + return SZGPU_STATUS_OK; +} + +extern "C" void szgpu_workspace_destroy(szgpu_workspace_t* ws) { + if (!ws) return; + if (ws->d_workspace) cudaFree(ws->d_workspace); + if (ws->h_cmpSize) cudaFreeHost(ws->h_cmpSize); + delete ws; +} + +extern "C" size_t szgpu_workspace_capacity(const szgpu_workspace_t* ws) { + return ws ? ws->max_n_elements : 0; +} + +template +static szgpu_status_t szgpu_compress_with_ws( + const FpT* d_in, + unsigned char* d_cmp, + size_t n, + FpT eb, + szgpu_workspace_t* ws, + cudaStream_t stream, + szgpu_compress_result_t* out_result) +{ + if (!d_in || !d_cmp || !ws || !out_result) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + if (!(eb > (FpT)0.0f)) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n > ws->max_n_elements) return SZGPU_STATUS_WORKSPACE_TOO_SMALL; + + size_t numTiles = szgpu_num_tiles(n); + size_t numBlocks = numTiles * SZGPU_BLKS_PER_TILE; + size_t dataOfs = numBlocks; + + bool use_small = (n <= SZGPU_SMALL_THRESHOLD); + int gs = use_small ? szgpu_grid_size_sm(numTiles) : szgpu_grid_size(numTiles); + + cudaError_t e = cudaMemsetAsync(ws->flag, 0, sizeof(int) * (gs + 1), stream); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + cudaGetLastError(); + + if (use_small) { + szgpu_compress_kernel_T<<>>( + d_in, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, + eb, n, numTiles, dataOfs); + } else { + szgpu_compress_kernel_T<<>>( + d_in, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, + eb, n, numTiles, dataOfs); + } + if (cudaGetLastError() != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + const unsigned int* d_total = (gs == 1) ? (ws->locOffset + 1) + : (ws->cmpOffset + gs); + e = cudaMemcpyAsync(ws->h_cmpSize, d_total, + sizeof(unsigned int), cudaMemcpyDeviceToHost, stream); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + e = cudaStreamSynchronize(stream); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + if ((unsigned long long)numTiles * (SZGPU_MAX_TILE_BYTES - SZGPU_BLKS_PER_TILE) + > 0xffffffffull) { + std::vector parts((size_t)gs); + e = cudaMemcpy(parts.data(), ws->locOffset + 1, + sizeof(unsigned int) * (size_t)gs, cudaMemcpyDeviceToHost); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + unsigned long long total = 0; + for (int i = 0; i < gs; ++i) total += parts[(size_t)i]; + if (total > 0xffffffffull) return SZGPU_STATUS_TOO_LARGE; + } + + out_result->cmp_size = dataOfs + *ws->h_cmpSize; + out_result->num_tiles = numTiles; + out_result->data_offset = dataOfs; + return SZGPU_STATUS_OK; +} + +template +static szgpu_status_t szgpu_compress_core( + const FpT* d_in, + unsigned char* d_cmp, + size_t n, + FpT eb, + szgpu_workspace_t* ws, + cudaStream_t stream, + szgpu_compress_result_t* out_result) +{ + if (ws) return szgpu_compress_with_ws(d_in, d_cmp, n, eb, ws, stream, out_result); + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + + szgpu_workspace_t* tmp = nullptr; + szgpu_status_t s = szgpu_workspace_create(&tmp, n); + if (s != SZGPU_STATUS_OK) return s; + s = szgpu_compress_with_ws(d_in, d_cmp, n, eb, tmp, stream, out_result); + szgpu_workspace_destroy(tmp); + return s; +} + +extern "C" szgpu_status_t szgpu_compress( + const float* d_in, unsigned char* d_cmp, size_t n, float eb, + szgpu_workspace_t* ws, cudaStream_t stream, szgpu_compress_result_t* out_result) +{ + return szgpu_compress_core(d_in, d_cmp, n, eb, ws, stream, out_result); +} + +extern "C" szgpu_status_t szgpu_compress_f64( + const double* d_in, unsigned char* d_cmp, size_t n, double eb, + szgpu_workspace_t* ws, cudaStream_t stream, szgpu_compress_result_t* out_result) +{ + return szgpu_compress_core(d_in, d_cmp, n, eb, ws, stream, out_result); +} + +template +static szgpu_status_t szgpu_decompress_with_ws( + FpT* d_out, + const unsigned char* d_cmp, + size_t n, + FpT eb, + const szgpu_compress_result_t* result, + szgpu_workspace_t* ws, + cudaStream_t stream) +{ + if (!d_out || !d_cmp || !ws || !result) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + if (!(eb > (FpT)0.0f)) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n > ws->max_n_elements) return SZGPU_STATUS_WORKSPACE_TOO_SMALL; + + const size_t expect_tiles = szgpu_num_tiles(n); + if (result->num_tiles != expect_tiles) return SZGPU_STATUS_INVALID_ARGUMENT; + if (result->data_offset != expect_tiles * SZGPU_BLKS_PER_TILE) + return SZGPU_STATUS_INVALID_ARGUMENT; + + bool use_small = (n <= SZGPU_SMALL_THRESHOLD); + int gs = use_small ? szgpu_grid_size_sm(result->num_tiles) + : szgpu_grid_size(result->num_tiles); + + cudaError_t e = cudaMemsetAsync(ws->flag, 0, sizeof(int) * (gs + 1), stream); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + size_t origSize = n * sizeof(FpT); + bool high_cr = (result->cmp_size > 0) && (origSize / result->cmp_size > 100); + if (high_cr) { + e = cudaMemsetAsync(d_out, 0, origSize, stream); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + } + + cudaGetLastError(); + + if (use_small) { + szgpu_decompress_kernel_T<<>>( + d_out, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, + eb, n, result->num_tiles, result->data_offset, high_cr); + } else { + szgpu_decompress_kernel_T<<>>( + d_out, d_cmp, ws->cmpOffset, ws->locOffset, ws->flag, + eb, n, result->num_tiles, result->data_offset, high_cr); + } + if (cudaGetLastError() != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + return SZGPU_STATUS_OK; +} + +template +static szgpu_status_t szgpu_decompress_core( + FpT* d_out, + const unsigned char* d_cmp, + size_t n, + FpT eb, + const szgpu_compress_result_t* result, + szgpu_workspace_t* ws, + cudaStream_t stream) +{ + if (ws) return szgpu_decompress_with_ws(d_out, d_cmp, n, eb, result, ws, stream); + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + + szgpu_workspace_t* tmp = nullptr; + szgpu_status_t s = szgpu_workspace_create(&tmp, n); + if (s != SZGPU_STATUS_OK) return s; + s = szgpu_decompress_with_ws(d_out, d_cmp, n, eb, result, tmp, stream); + cudaError_t e = cudaStreamSynchronize(stream); + szgpu_workspace_destroy(tmp); + if (s == SZGPU_STATUS_OK && e != cudaSuccess) s = SZGPU_STATUS_CUDA_ERROR; + return s; +} + +extern "C" szgpu_status_t szgpu_decompress( + float* d_out, const unsigned char* d_cmp, size_t n, float eb, + const szgpu_compress_result_t* result, szgpu_workspace_t* ws, cudaStream_t stream) +{ + return szgpu_decompress_core(d_out, d_cmp, n, eb, result, ws, stream); +} + +extern "C" szgpu_status_t szgpu_decompress_f64( + double* d_out, const unsigned char* d_cmp, size_t n, double eb, + const szgpu_compress_result_t* result, szgpu_workspace_t* ws, cudaStream_t stream) +{ + return szgpu_decompress_core(d_out, d_cmp, n, eb, result, ws, stream); +} + +template +static szgpu_status_t szgpu_compress_hostptr_core( + const FpT* in, + unsigned char* cmp, + size_t n, + FpT eb, + size_t* out_cmp_size) +{ + if (!in || !cmp || !out_cmp_size) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + if (!(eb > (FpT)0.0f)) return SZGPU_STATUS_INVALID_ARGUMENT; + + *out_cmp_size = 0; + + FpT* d_in = nullptr; + unsigned char* d_cmp = nullptr; + szgpu_workspace_t* ws = nullptr; + + cudaError_t e = cudaMalloc(&d_in, n * sizeof(FpT)); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + e = cudaMalloc(&d_cmp, szgpu_max_compressed_bytes(n)); + if (e != cudaSuccess) { + szgpu_hostptr_release(d_in, nullptr, nullptr); + return SZGPU_STATUS_CUDA_ERROR; + } + + szgpu_status_t s = szgpu_workspace_create(&ws, n); + if (s != SZGPU_STATUS_OK) { + szgpu_hostptr_release(d_in, d_cmp, nullptr); + return s; + } + + e = cudaMemcpy(d_in, in, n * sizeof(FpT), cudaMemcpyHostToDevice); + if (e != cudaSuccess) { + szgpu_hostptr_release(d_in, d_cmp, ws); + return SZGPU_STATUS_CUDA_ERROR; + } + + szgpu_compress_result_t r; + s = szgpu_compress_core(d_in, d_cmp, n, eb, ws, 0, &r); + if (s != SZGPU_STATUS_OK) { + szgpu_hostptr_release(d_in, d_cmp, ws); + return s; + } + + e = cudaMemcpy(cmp, d_cmp, r.cmp_size, cudaMemcpyDeviceToHost); + szgpu_hostptr_release(d_in, d_cmp, ws); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + *out_cmp_size = r.cmp_size; + return SZGPU_STATUS_OK; +} + +extern "C" szgpu_status_t szgpu_compress_hostptr( + const float* in, unsigned char* cmp, size_t n, float eb, size_t* out_cmp_size) +{ + return szgpu_compress_hostptr_core(in, cmp, n, eb, out_cmp_size); +} + +extern "C" szgpu_status_t szgpu_compress_hostptr_f64( + const double* in, unsigned char* cmp, size_t n, double eb, size_t* out_cmp_size) +{ + return szgpu_compress_hostptr_core(in, cmp, n, eb, out_cmp_size); +} + +template +static szgpu_status_t szgpu_decompress_hostptr_core( + FpT* out, + const unsigned char* cmp, + size_t cmp_size, + size_t n, + FpT eb) +{ + if (!out || !cmp) return SZGPU_STATUS_INVALID_ARGUMENT; + if (n == 0) return SZGPU_STATUS_INVALID_ARGUMENT; + if (!(eb > (FpT)0.0f)) return SZGPU_STATUS_INVALID_ARGUMENT; + + size_t header = szgpu_num_tiles(n) * SZGPU_BLKS_PER_TILE; + if (cmp_size == 0 || cmp_size < header) return SZGPU_STATUS_INVALID_ARGUMENT; + + FpT* d_out = nullptr; + unsigned char* d_cmp = nullptr; + szgpu_workspace_t* ws = nullptr; + + cudaError_t e = cudaMalloc(&d_out, n * sizeof(FpT)); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + e = cudaMalloc(&d_cmp, cmp_size); + if (e != cudaSuccess) { + szgpu_hostptr_release(d_out, nullptr, nullptr); + return SZGPU_STATUS_CUDA_ERROR; + } + + szgpu_status_t s = szgpu_workspace_create(&ws, n); + if (s != SZGPU_STATUS_OK) { + szgpu_hostptr_release(d_out, d_cmp, nullptr); + return s; + } + + e = cudaMemcpy(d_cmp, cmp, cmp_size, cudaMemcpyHostToDevice); + if (e != cudaSuccess) { + szgpu_hostptr_release(d_out, d_cmp, ws); + return SZGPU_STATUS_CUDA_ERROR; + } + + szgpu_compress_result_t r = szgpu_make_result(n, cmp_size); + s = szgpu_decompress_core(d_out, d_cmp, n, eb, &r, ws, 0); + if (s != SZGPU_STATUS_OK) { + szgpu_hostptr_release(d_out, d_cmp, ws); + return s; + } + + e = cudaStreamSynchronize(0); + if (e != cudaSuccess) { + szgpu_hostptr_release(d_out, d_cmp, ws); + return SZGPU_STATUS_CUDA_ERROR; + } + + e = cudaMemcpy(out, d_out, n * sizeof(FpT), cudaMemcpyDeviceToHost); + szgpu_hostptr_release(d_out, d_cmp, ws); + if (e != cudaSuccess) return SZGPU_STATUS_CUDA_ERROR; + + return SZGPU_STATUS_OK; +} + +extern "C" szgpu_status_t szgpu_decompress_hostptr( + float* out, const unsigned char* cmp, size_t cmp_size, size_t n, float eb) +{ + return szgpu_decompress_hostptr_core(out, cmp, cmp_size, n, eb); +} + +extern "C" szgpu_status_t szgpu_decompress_hostptr_f64( + double* out, const unsigned char* cmp, size_t cmp_size, size_t n, double eb) +{ + return szgpu_decompress_hostptr_core(out, cmp, cmp_size, n, eb); +} diff --git a/src/fsz_internal.cuh b/src/szgpu_internal.cuh similarity index 67% rename from src/fsz_internal.cuh rename to src/szgpu_internal.cuh index 0a1efe4..577c620 100644 --- a/src/fsz_internal.cuh +++ b/src/szgpu_internal.cuh @@ -1,30 +1,30 @@ -#ifndef FSZ_INTERNAL_CUH -#define FSZ_INTERNAL_CUH +#ifndef SZGPU_INTERNAL_CUH +#define SZGPU_INTERNAL_CUH #include #include -namespace fsz_detail { - -#define FSZ_BLK_DIM 32 -#define FSZ_TILE_SIZE 256 -#define FSZ_BLKS_PER_TILE 8 -#define FSZ_TILES_PER_THR 4 -#define FSZ_TILES_PER_THR_SM 1 -#define FSZ_THREAD_CHUNK 1024 -#define FSZ_WARP_SZ 32 -#define FSZ_TBLOCK_SZ 32 -#define FSZ_RATE_MASK 0x1F -#define FSZ_LZ2_FLAG 0x80 -#define FSZ_MEAN_FLAG 0x40 - -#define FSZ_MAX_TILE_BYTES (FSZ_BLKS_PER_TILE \ - + FSZ_BLKS_PER_TILE * (4 + FSZ_RATE_MASK * 4) \ +namespace szgpu_detail { + +#define SZGPU_BLK_DIM 32 +#define SZGPU_TILE_SIZE 256 +#define SZGPU_BLKS_PER_TILE 8 +#define SZGPU_TILES_PER_THR 4 +#define SZGPU_TILES_PER_THR_SM 1 +#define SZGPU_THREAD_CHUNK 1024 +#define SZGPU_WARP_SZ 32 +#define SZGPU_TBLOCK_SZ 32 +#define SZGPU_RATE_MASK 0x1F +#define SZGPU_LZ2_FLAG 0x80 +#define SZGPU_MEAN_FLAG 0x40 + +#define SZGPU_MAX_TILE_BYTES (SZGPU_BLKS_PER_TILE \ + + SZGPU_BLKS_PER_TILE * (4 + SZGPU_RATE_MASK * 4) \ + 8) -#define FSZ_SMALL_THRESHOLD 20000000 +#define SZGPU_SMALL_THRESHOLD 20000000 -__device__ __forceinline__ int fsz_quantize(float data, float recipPrec) { +__device__ __forceinline__ int szgpu_quantize(float data, float recipPrec) { int result; asm("{\n\t" ".reg .f32 dr, t;\n\t" @@ -40,11 +40,11 @@ __device__ __forceinline__ int fsz_quantize(float data, float recipPrec) { return result; } -__device__ __forceinline__ int fsz_bitnum(unsigned int x) { +__device__ __forceinline__ int szgpu_bitnum(unsigned int x) { int lz; asm("clz.b32 %0, %1;" : "=r"(lz) : "r"(x)); return 32 - lz; } -__device__ __forceinline__ float4 fsz_load4(const float* __restrict__ data, +__device__ __forceinline__ float4 szgpu_load4(const float* __restrict__ data, size_t gi, size_t nbEle) { float4 v; if (gi + 3 < nbEle) { @@ -58,16 +58,16 @@ __device__ __forceinline__ float4 fsz_load4(const float* __restrict__ data, return v; } -__device__ __forceinline__ uchar4 fsz_ldg_uchar4(const unsigned char* __restrict__ cmpData, +__device__ __forceinline__ uchar4 szgpu_ldg_uchar4(const unsigned char* __restrict__ cmpData, unsigned int byte_ofs) { return __ldg(reinterpret_cast(cmpData) + (byte_ofs >> 2)); } -__device__ __forceinline__ int fsz_quantize(double data, double recipPrec) { +__device__ __forceinline__ int szgpu_quantize(double data, double recipPrec) { return __double2int_rn(data * recipPrec); } -__device__ __forceinline__ double4 fsz_load4(const double* __restrict__ data, +__device__ __forceinline__ double4 szgpu_load4(const double* __restrict__ data, size_t gi, size_t nbEle) { double4 v; if (gi + 3 < nbEle) { @@ -81,14 +81,14 @@ __device__ __forceinline__ double4 fsz_load4(const double* __restrict__ data, return v; } -__device__ __forceinline__ float fsz_mean_value(int mean_q, float eb) { +__device__ __forceinline__ float szgpu_mean_value(int mean_q, float eb) { return (float)((double)mean_q * (double)(2.0f * eb)); } -__device__ __forceinline__ double fsz_mean_value(int mean_q, double eb) { +__device__ __forceinline__ double szgpu_mean_value(int mean_q, double eb) { return (double)mean_q * (2.0 * eb); } -__device__ __forceinline__ void fsz_write_mean(unsigned char* __restrict__ cmpData, +__device__ __forceinline__ void szgpu_write_mean(unsigned char* __restrict__ cmpData, unsigned int ofs, float mu) { unsigned int mu_bits; memcpy(&mu_bits, &mu, 4); @@ -99,7 +99,7 @@ __device__ __forceinline__ void fsz_write_mean(unsigned char* __restrict__ cmpDa mc.w = (mu_bits >> 24) & 0xff; reinterpret_cast(cmpData)[ofs >> 2] = mc; } -__device__ __forceinline__ void fsz_write_mean(unsigned char* __restrict__ cmpData, +__device__ __forceinline__ void szgpu_write_mean(unsigned char* __restrict__ cmpData, unsigned int ofs, double mu) { unsigned long long b; memcpy(&b, &mu, 8); @@ -117,12 +117,12 @@ __device__ __forceinline__ void fsz_write_mean(unsigned char* __restrict__ cmpDa } template -__device__ __forceinline__ FpT fsz_read_mean(const unsigned char* __restrict__ cmpData, +__device__ __forceinline__ FpT szgpu_read_mean(const unsigned char* __restrict__ cmpData, unsigned int ofs); template <> -__device__ __forceinline__ float fsz_read_mean( +__device__ __forceinline__ float szgpu_read_mean( const unsigned char* __restrict__ cmpData, unsigned int ofs) { - uchar4 mc = fsz_ldg_uchar4(cmpData, ofs); + uchar4 mc = szgpu_ldg_uchar4(cmpData, ofs); unsigned int mu_bits = (unsigned int)mc.x | ((unsigned int)mc.y << 8) | ((unsigned int)mc.z << 16) | ((unsigned int)mc.w << 24); float mu; @@ -130,10 +130,10 @@ __device__ __forceinline__ float fsz_read_mean( return mu; } template <> -__device__ __forceinline__ double fsz_read_mean( +__device__ __forceinline__ double szgpu_read_mean( const unsigned char* __restrict__ cmpData, unsigned int ofs) { - uchar4 lo = fsz_ldg_uchar4(cmpData, ofs); - uchar4 hi = fsz_ldg_uchar4(cmpData, ofs + 4); + uchar4 lo = szgpu_ldg_uchar4(cmpData, ofs); + uchar4 hi = szgpu_ldg_uchar4(cmpData, ofs + 4); unsigned long long b = (unsigned long long)lo.x | ((unsigned long long)lo.y << 8) | ((unsigned long long)lo.z << 16) | ((unsigned long long)lo.w << 24) | @@ -144,13 +144,13 @@ __device__ __forceinline__ double fsz_read_mean( return mu; } -__device__ __forceinline__ void fsz_store4(float* __restrict__ p, size_t gi, +__device__ __forceinline__ void szgpu_store4(float* __restrict__ p, size_t gi, float a, float b, float c, float d) { float4 o; o.x = a; o.y = b; o.z = c; o.w = d; reinterpret_cast(p)[gi >> 2] = o; } -__device__ __forceinline__ void fsz_store4(double* __restrict__ p, size_t gi, +__device__ __forceinline__ void szgpu_store4(double* __restrict__ p, size_t gi, double a, double b, double c, double d) { double4 o; o.x = a; o.y = b; o.z = c; o.w = d; diff --git a/src/fsz_kernels.cuh b/src/szgpu_kernels.cuh similarity index 80% rename from src/fsz_kernels.cuh rename to src/szgpu_kernels.cuh index 656a296..5b413a7 100644 --- a/src/fsz_kernels.cuh +++ b/src/szgpu_kernels.cuh @@ -1,12 +1,12 @@ -#ifndef FSZ_KERNELS_CUH -#define FSZ_KERNELS_CUH +#ifndef SZGPU_KERNELS_CUH +#define SZGPU_KERNELS_CUH -#include "fsz_internal.cuh" +#include "szgpu_internal.cuh" -namespace fsz_detail { +namespace szgpu_detail { template -__global__ void fsz_compress_kernel_T( +__global__ void szgpu_compress_kernel_T( const FpT* __restrict__ oriData, unsigned char* __restrict__ cmpData, volatile unsigned int* __restrict__ cmpOffset, @@ -18,7 +18,7 @@ __global__ void fsz_compress_kernel_T( __shared__ unsigned int excl_sum; __shared__ unsigned int base_idx; const int tid = threadIdx.x; - const int idx = blockIdx.x * FSZ_TBLOCK_SZ + tid; + const int idx = blockIdx.x * SZGPU_TBLOCK_SZ + tid; const int lane = idx & 0x1f; const int warp = idx >> 5; if (!tid) { excl_sum = 0; base_idx = 0; } @@ -26,8 +26,8 @@ __global__ void fsz_compress_kernel_T( const FpT recipPrec = (FpT)0.5f / eb; const int TPT = TILES_PER_THREAD; - const int BN = TPT * FSZ_BLKS_PER_TILE; - const int baseTile = (warp * FSZ_WARP_SZ + lane) * TPT; + const int BN = TPT * SZGPU_BLKS_PER_TILE; + const int baseTile = (warp * SZGPU_WARP_SZ + lane) * TPT; int fixed_rate[BN]; int tile_mean_q[TPT]; @@ -37,7 +37,7 @@ __global__ void fsz_compress_kernel_T( for (int t = 0; t < TPT; t++) { int tileId = baseTile + t; - size_t tBase = (size_t)tileId * FSZ_TILE_SIZE; + size_t tBase = (size_t)tileId * SZGPU_TILE_SIZE; if (tileId < (int)numTiles) { FpT fsum = (FpT)0.0f; @@ -48,23 +48,23 @@ __global__ void fsz_compress_kernel_T( int blk0_d1_nomean_lz2 = 0; int blk0_max_2_31_lz2 = 0; int blk0_max_1_31_lz1 = 0; - int lz1_rate_t[FSZ_BLKS_PER_TILE], lz2_rate_t[FSZ_BLKS_PER_TILE]; - int lz1_pb1_t[FSZ_BLKS_PER_TILE]; - int lz2_pb1_t[FSZ_BLKS_PER_TILE], lz2_pb2_t[FSZ_BLKS_PER_TILE]; + int lz1_rate_t[SZGPU_BLKS_PER_TILE], lz2_rate_t[SZGPU_BLKS_PER_TILE]; + int lz1_pb1_t[SZGPU_BLKS_PER_TILE]; + int lz2_pb1_t[SZGPU_BLKS_PER_TILE], lz2_pb2_t[SZGPU_BLKS_PER_TILE]; - for (int blk = 0; blk < FSZ_BLKS_PER_TILE; blk++) { - size_t bBase = tBase + blk * FSZ_BLK_DIM; + for (int blk = 0; blk < SZGPU_BLKS_PER_TILE; blk++) { + size_t bBase = tBase + blk * SZGPU_BLK_DIM; lz1_pb1_t[blk] = prev1_lz1; lz2_pb1_t[blk] = prev1_lz2; lz2_pb2_t[blk] = prev2_lz2; int lz1Max = 0, lz2Max = 0; - for (int k = 0; k < FSZ_BLK_DIM; k += 4) { - auto v4 = fsz_load4(oriData, bBase + k, nbEle); + for (int k = 0; k < SZGPU_BLK_DIM; k += 4) { + auto v4 = szgpu_load4(oriData, bBase + k, nbEle); fsum += v4.x + v4.y + v4.z + v4.w; FpT fv[4] = {v4.x, v4.y, v4.z, v4.w}; for (int kk = 0; kk < 4; kk++) { - int q = fsz_quantize(fv[kk], recipPrec); + int q = szgpu_quantize(fv[kk], recipPrec); int d1 = q - prev1_lz1; prev1_lz1 = q; int a1 = d1 < 0 ? -d1 : d1; @@ -83,19 +83,19 @@ __global__ void fsz_compress_kernel_T( } } } - int lz1r = fsz_bitnum((unsigned int)lz1Max); + int lz1r = szgpu_bitnum((unsigned int)lz1Max); if (lz1r > 31) lz1r = 31; lz1_rate_t[blk] = lz1r; - int lz2r = fsz_bitnum((unsigned int)lz2Max); + int lz2r = szgpu_bitnum((unsigned int)lz2Max); if (lz2r > 31) lz2r = 31; lz2_rate_t[blk] = lz2r; } - int mean_q = fsz_quantize(fsum / (FpT)FSZ_TILE_SIZE, recipPrec); + int mean_q = szgpu_quantize(fsum / (FpT)SZGPU_TILE_SIZE, recipPrec); int lz1_d0_mean = blk0_q0 - mean_q; int lz1_ad0_mean = lz1_d0_mean < 0 ? -lz1_d0_mean : lz1_d0_mean; int lz1_max0_mean = lz1_ad0_mean > blk0_max_1_31_lz1 ? lz1_ad0_mean : blk0_max_1_31_lz1; - int lz1_rate0_mean = fsz_bitnum((unsigned int)lz1_max0_mean); + int lz1_rate0_mean = szgpu_bitnum((unsigned int)lz1_max0_mean); if (lz1_rate0_mean > 31) lz1_rate0_mean = 31; int lz1_rate0_nomean = lz1_rate_t[0]; @@ -106,12 +106,12 @@ __global__ void fsz_compress_kernel_T( int lz2_max0_mean = lz2_ad0_mean; lz2_max0_mean = lz2_max0_mean > lz2_ad1_mean ? lz2_max0_mean : lz2_ad1_mean; lz2_max0_mean = lz2_max0_mean > blk0_max_2_31_lz2 ? lz2_max0_mean : blk0_max_2_31_lz2; - int lz2_rate0_mean = fsz_bitnum((unsigned int)lz2_max0_mean); + int lz2_rate0_mean = szgpu_bitnum((unsigned int)lz2_max0_mean); if (lz2_rate0_mean > 31) lz2_rate0_mean = 31; int lz2_rate0_nomean = lz2_rate_t[0]; int lz1_blks17 = 0, lz2_blks17 = 0; - for (int b = 1; b < FSZ_BLKS_PER_TILE; b++) { + for (int b = 1; b < SZGPU_BLKS_PER_TILE; b++) { lz1_blks17 += lz1_rate_t[b] ? (4 + lz1_rate_t[b] * 4) : 0; lz2_blks17 += lz2_rate_t[b] ? (4 + lz2_rate_t[b] * 4) : 0; } @@ -138,16 +138,16 @@ __global__ void fsz_compress_kernel_T( int chosen_rate0 = use_lz2 ? (use_mean ? lz2_rate0_mean : lz2_rate0_nomean) : (use_mean ? lz1_rate0_mean : lz1_rate0_nomean); - int rate_byte_0 = chosen_rate0 | (use_lz2 ? FSZ_LZ2_FLAG : 0) | (use_mean ? FSZ_MEAN_FLAG : 0); - int bIdx0 = t * FSZ_BLKS_PER_TILE; + int rate_byte_0 = chosen_rate0 | (use_lz2 ? SZGPU_LZ2_FLAG : 0) | (use_mean ? SZGPU_MEAN_FLAG : 0); + int bIdx0 = t * SZGPU_BLKS_PER_TILE; fixed_rate[bIdx0] = rate_byte_0; - cmpData[tileId * FSZ_BLKS_PER_TILE] = (unsigned char)rate_byte_0; + cmpData[tileId * SZGPU_BLKS_PER_TILE] = (unsigned char)rate_byte_0; - for (int b = 1; b < FSZ_BLKS_PER_TILE; b++) { - int bIdx = t * FSZ_BLKS_PER_TILE + b; + for (int b = 1; b < SZGPU_BLKS_PER_TILE; b++) { + int bIdx = t * SZGPU_BLKS_PER_TILE + b; int rate = use_lz2 ? lz2_rate_t[b] : lz1_rate_t[b]; fixed_rate[bIdx] = rate; - cmpData[tileId * FSZ_BLKS_PER_TILE + b] = (unsigned char)rate; + cmpData[tileId * SZGPU_BLKS_PER_TILE + b] = (unsigned char)rate; } tile_mean_q[t] = use_mean ? mean_q : 0; @@ -157,8 +157,8 @@ __global__ void fsz_compress_kernel_T( tile_use_lz2[t] = false; tile_use_mean[t] = false; tile_mean_q[t] = 0; - for (int blk = 0; blk < FSZ_BLKS_PER_TILE; blk++) { - int bIdx = t * FSZ_BLKS_PER_TILE + blk; + for (int blk = 0; blk < SZGPU_BLKS_PER_TILE; blk++) { + int bIdx = t * SZGPU_BLKS_PER_TILE + blk; fixed_rate[bIdx] = 0; } } @@ -211,8 +211,8 @@ __global__ void fsz_compress_kernel_T( int prev_tile = -1; for (int j = 0; j < BN; j++) { - int t = j / FSZ_BLKS_PER_TILE; - int blk = j % FSZ_BLKS_PER_TILE; + int t = j / SZGPU_BLKS_PER_TILE; + int blk = j % SZGPU_BLKS_PER_TILE; bool use_lz2 = tile_use_lz2[t]; bool mean_enabled = tile_use_mean[t]; int mean_q = tile_mean_q[t]; @@ -222,7 +222,7 @@ __global__ void fsz_compress_kernel_T( prev_tile = t; } - int rate = fixed_rate[j] & FSZ_RATE_MASK; + int rate = fixed_rate[j] & SZGPU_RATE_MASK; unsigned int tmp_byte = rate ? (4 + rate * 4) : 0; if (blk == 0 && mean_enabled) tmp_byte += (unsigned int)sizeof(FpT); @@ -238,23 +238,23 @@ __global__ void fsz_compress_kernel_T( if (rate || (blk == 0 && mean_enabled)) { if (blk == 0 && mean_enabled) { - FpT mu = fsz_mean_value(mean_q, eb); - fsz_write_mean(cmpData, cmp_ofs, mu); + FpT mu = szgpu_mean_value(mean_q, eb); + szgpu_write_mean(cmpData, cmp_ofs, mu); cmp_ofs += (unsigned int)sizeof(FpT); } if (rate) { - int absQ[FSZ_BLK_DIM]; + int absQ[SZGPU_BLK_DIM]; unsigned int sf = 0; - size_t bBase = (size_t)(baseTile + t) * FSZ_TILE_SIZE + blk * FSZ_BLK_DIM; + size_t bBase = (size_t)(baseTile + t) * SZGPU_TILE_SIZE + blk * SZGPU_BLK_DIM; int p1 = enc_p1; int p2 = enc_p2; - for (int k = 0; k < FSZ_BLK_DIM; k += 4) { - auto v4 = fsz_load4(oriData, bBase + k, nbEle); + for (int k = 0; k < SZGPU_BLK_DIM; k += 4) { + auto v4 = szgpu_load4(oriData, bBase + k, nbEle); FpT fv[4] = {v4.x, v4.y, v4.z, v4.w}; for (int kk = 0; kk < 4; kk++) { - int q = fsz_quantize(fv[kk], recipPrec) - mean_q; + int q = szgpu_quantize(fv[kk], recipPrec) - mean_q; int delta; if (use_lz2) { delta = q - 2 * p1 + p2; @@ -318,7 +318,7 @@ __global__ void fsz_compress_kernel_T( } template -__global__ void fsz_decompress_kernel_T( +__global__ void szgpu_decompress_kernel_T( FpT* __restrict__ decData, const unsigned char* __restrict__ cmpData, volatile unsigned int* __restrict__ cmpOffset, @@ -330,27 +330,27 @@ __global__ void fsz_decompress_kernel_T( __shared__ unsigned int excl_sum; __shared__ unsigned int base_idx; const int tid = threadIdx.x; - const int idx = blockIdx.x * FSZ_TBLOCK_SZ + tid; + const int idx = blockIdx.x * SZGPU_TBLOCK_SZ + tid; const int lane = idx & 0x1f; const int warp = idx >> 5; if (!tid) { excl_sum = 0; base_idx = 0; } __syncthreads(); const int TPT = TILES_PER_THREAD; - const int BN = TPT * FSZ_BLKS_PER_TILE; - const int baseTile = (warp * FSZ_WARP_SZ + lane) * TPT; + const int BN = TPT * SZGPU_BLKS_PER_TILE; + const int baseTile = (warp * SZGPU_WARP_SZ + lane) * TPT; int fixed_rate[BN]; unsigned int thread_ofs = 0; for (int j = 0; j < BN; j++) { - int t = j / FSZ_BLKS_PER_TILE, blk = j % FSZ_BLKS_PER_TILE; + int t = j / SZGPU_BLKS_PER_TILE, blk = j % SZGPU_BLKS_PER_TILE; int tileId = baseTile + t; if (tileId < (int)numTiles) { - int rate_byte = cmpData[tileId * FSZ_BLKS_PER_TILE + blk]; + int rate_byte = cmpData[tileId * SZGPU_BLKS_PER_TILE + blk]; fixed_rate[j] = rate_byte; - int rate = rate_byte & FSZ_RATE_MASK; + int rate = rate_byte & SZGPU_RATE_MASK; thread_ofs += rate ? (4 + rate * 4) : 0; - if (blk == 0 && (rate_byte & FSZ_MEAN_FLAG)) thread_ofs += (unsigned int)sizeof(FpT); + if (blk == 0 && (rate_byte & SZGPU_MEAN_FLAG)) thread_ofs += (unsigned int)sizeof(FpT); } else { fixed_rate[j] = 0; } @@ -406,20 +406,20 @@ __global__ void fsz_decompress_kernel_T( bool mean_enabled = false; for (int j = 0; j < BN; j++) { - int t = j / FSZ_BLKS_PER_TILE; - int blk = j % FSZ_BLKS_PER_TILE; + int t = j / SZGPU_BLKS_PER_TILE; + int blk = j % SZGPU_BLKS_PER_TILE; if (t != cur_tile) { cur_tile = t; p1 = 0; p2 = 0; mu = (FpT)0.0f; - int rate_byte_0 = fixed_rate[t * FSZ_BLKS_PER_TILE]; - use_lz2 = (rate_byte_0 & FSZ_LZ2_FLAG) != 0; - mean_enabled = (rate_byte_0 & FSZ_MEAN_FLAG) != 0; + int rate_byte_0 = fixed_rate[t * SZGPU_BLKS_PER_TILE]; + use_lz2 = (rate_byte_0 & SZGPU_LZ2_FLAG) != 0; + mean_enabled = (rate_byte_0 & SZGPU_MEAN_FLAG) != 0; } int rate_byte = fixed_rate[j]; - int rate = rate_byte & FSZ_RATE_MASK; + int rate = rate_byte & SZGPU_RATE_MASK; unsigned int tmp_byte = rate ? (4 + rate * 4) : 0; if (blk == 0 && mean_enabled) tmp_byte += (unsigned int)sizeof(FpT); @@ -435,7 +435,7 @@ __global__ void fsz_decompress_kernel_T( else cmp_ofs = base_cmp + cur_ofs + prev_thr; if (blk == 0 && mean_enabled) { - mu = fsz_read_mean(cmpData, cmp_ofs); + mu = szgpu_read_mean(cmpData, cmp_ofs); cmp_ofs += (unsigned int)sizeof(FpT); } @@ -443,15 +443,15 @@ __global__ void fsz_decompress_kernel_T( if (rate == 0) { if (tileId < (int)numTiles) { - size_t bBase = (size_t)tileId * FSZ_TILE_SIZE + blk * FSZ_BLK_DIM; + size_t bBase = (size_t)tileId * SZGPU_TILE_SIZE + blk * SZGPU_BLK_DIM; if (use_lz2) { FpT f_stride = (FpT)(p1 - p2) * step; FpT f_base = (FpT)(2*p1 - p2) * step + mu; if (!(skipZeroWrites && f_base == (FpT)0.0f && f_stride == (FpT)0.0f)) { size_t gi = bBase; - for (int k = 0; k < FSZ_BLK_DIM; k += 4) { + for (int k = 0; k < SZGPU_BLK_DIM; k += 4) { if (gi + 3 < nbEle) { - fsz_store4(decData, gi, + szgpu_store4(decData, gi, f_base + (FpT)(k) * f_stride, f_base + (FpT)(k+1) * f_stride, f_base + (FpT)(k+2) * f_stride, @@ -471,9 +471,9 @@ __global__ void fsz_decompress_kernel_T( FpT fill_val = (FpT)p1 * step + mu; if (!(skipZeroWrites && fill_val == (FpT)0.0f)) { size_t gi = bBase; - for (int k = 0; k < FSZ_BLK_DIM; k += 4) { + for (int k = 0; k < SZGPU_BLK_DIM; k += 4) { if (gi + 3 < nbEle) { - fsz_store4(decData, gi, + szgpu_store4(decData, gi, fill_val, fill_val, fill_val, fill_val); } else { for (int kk = 0; kk < 4 && gi+kk < nbEle; kk++) @@ -485,17 +485,17 @@ __global__ void fsz_decompress_kernel_T( } } } else { - uchar4 sc = fsz_ldg_uchar4(cmpData, cmp_ofs); + uchar4 sc = szgpu_ldg_uchar4(cmpData, cmp_ofs); unsigned int signs = ((unsigned int)sc.x << 24) | ((unsigned int)sc.y << 16) | ((unsigned int)sc.z << 8) | (unsigned int)sc.w; unsigned int cd = cmp_ofs + 4; - int mag[FSZ_BLK_DIM]; + int mag[SZGPU_BLK_DIM]; #pragma unroll - for (int k = 0; k < FSZ_BLK_DIM; k++) mag[k] = 0; + for (int k = 0; k < SZGPU_BLK_DIM; k++) mag[k] = 0; for (int b = 0; b < rate; b++) { - uchar4 bc = fsz_ldg_uchar4(cmpData, cd); + uchar4 bc = szgpu_ldg_uchar4(cmpData, cd); cd += 4; #pragma unroll for (int k = 0; k < 8; k++) mag[k] |= ((int)((bc.x >> (7-k)) & 1)) << b; @@ -507,25 +507,25 @@ __global__ void fsz_decompress_kernel_T( for (int k = 0; k < 8; k++) mag[24+k] |= ((int)((bc.w >> (7-k)) & 1)) << b; } - int vals[FSZ_BLK_DIM]; - for (int k = 0; k < FSZ_BLK_DIM; k++) + int vals[SZGPU_BLK_DIM]; + for (int k = 0; k < SZGPU_BLK_DIM; k++) vals[k] = (signs & (1u << (31-k))) ? -mag[k] : mag[k]; if (use_lz2) { - for (int k = 0; k < FSZ_BLK_DIM; k++) { + for (int k = 0; k < SZGPU_BLK_DIM; k++) { int q = 2*p1 - p2 + vals[k]; p2 = p1; p1 = q; vals[k] = q; } } else { - for (int k = 0; k < FSZ_BLK_DIM; k++) { + for (int k = 0; k < SZGPU_BLK_DIM; k++) { int q = p1 + vals[k]; p1 = q; vals[k] = q; } } if (tileId < (int)numTiles) { - size_t bBase = (size_t)tileId * FSZ_TILE_SIZE + blk * FSZ_BLK_DIM; - for (int k = 0; k < FSZ_BLK_DIM; k += 4) { + size_t bBase = (size_t)tileId * SZGPU_TILE_SIZE + blk * SZGPU_BLK_DIM; + for (int k = 0; k < SZGPU_BLK_DIM; k += 4) { size_t gi = bBase + k; if (gi + 3 < nbEle) { - fsz_store4(decData, gi, + szgpu_store4(decData, gi, (FpT)vals[k] * step + mu, (FpT)vals[k+1] * step + mu, (FpT)vals[k+2] * step + mu, diff --git a/tests/test_hostptr.cu b/tests/test_hostptr.cu index 1fc1136..b80f9d3 100644 --- a/tests/test_hostptr.cu +++ b/tests/test_hostptr.cu @@ -1,4 +1,4 @@ -#include "fsz/fsz.hpp" +#include "szgpu/szgpu.hpp" #include #include @@ -45,19 +45,19 @@ static void fill_near_constant(std::vector& v, float base) { static bool compress_device(const std::vector& h_in, float eb, std::vector& out_cmp, - fsz_compress_result_t* out_result) { + szgpu_compress_result_t* out_result) { std::size_t n = h_in.size(); float* d_in = nullptr; unsigned char* d_cmp = nullptr; if (cudaMalloc(&d_in, n * sizeof(float)) != cudaSuccess) return false; - if (cudaMalloc(&d_cmp, fsz::max_compressed_bytes(n)) != cudaSuccess) { + if (cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(n)) != cudaSuccess) { cudaFree(d_in); return false; } cudaMemcpy(d_in, h_in.data(), n * sizeof(float), cudaMemcpyHostToDevice); - fsz::Workspace ws(n); - *out_result = fsz::compress(d_in, d_cmp, n, eb, ws); + szgpu::Workspace ws(n); + *out_result = szgpu::compress(d_in, d_cmp, n, eb, ws); out_cmp.resize(out_result->cmp_size); cudaMemcpy(out_cmp.data(), d_cmp, out_result->cmp_size, cudaMemcpyDeviceToHost); @@ -69,19 +69,19 @@ static bool compress_device(const std::vector& h_in, float eb, static bool compress_device_f64(const std::vector& h_in, double eb, std::vector& out_cmp, - fsz_compress_result_t* out_result) { + szgpu_compress_result_t* out_result) { std::size_t n = h_in.size(); double* d_in = nullptr; unsigned char* d_cmp = nullptr; if (cudaMalloc(&d_in, n * sizeof(double)) != cudaSuccess) return false; - if (cudaMalloc(&d_cmp, fsz::max_compressed_bytes(n)) != cudaSuccess) { + if (cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(n)) != cudaSuccess) { cudaFree(d_in); return false; } cudaMemcpy(d_in, h_in.data(), n * sizeof(double), cudaMemcpyHostToDevice); - fsz::Workspace ws(n); - *out_result = fsz::compress(d_in, d_cmp, n, eb, ws); + szgpu::Workspace ws(n); + *out_result = szgpu::compress(d_in, d_cmp, n, eb, ws); out_cmp.resize(out_result->cmp_size); cudaMemcpy(out_cmp.data(), d_cmp, out_result->cmp_size, cudaMemcpyDeviceToHost); @@ -117,7 +117,7 @@ static double max_abs_diff_f64(const std::vector& a, const std::vector cmpA_dev; - fsz_compress_result_t rA; + szgpu_compress_result_t rA; if (!compress_device(inA, eb, cmpA_dev, &rA)) { std::printf("[FAIL] setup device-pointer compress of field A failed\n"); return 1; } - std::vector cmpA_host(fsz::max_compressed_bytes(nA)); + std::vector cmpA_host(szgpu::max_compressed_bytes(nA)); std::size_t cmpA_host_size = 0; - fsz_status_t st = fsz_compress_hostptr(inA.data(), cmpA_host.data(), nA, eb, + szgpu_status_t st = szgpu_compress_hostptr(inA.data(), cmpA_host.data(), nA, eb, &cmpA_host_size); - bool sizes_match = (st == FSZ_STATUS_OK) && (cmpA_host_size == rA.cmp_size); + bool sizes_match = (st == SZGPU_STATUS_OK) && (cmpA_host_size == rA.cmp_size); bool bytes_match = sizes_match && std::memcmp(cmpA_host.data(), cmpA_dev.data(), rA.cmp_size) == 0; std::snprintf(detail, sizeof(detail), "n=%zu status=%s device=%zu B hostptr=%zu B memcmp=%s", - nA, fsz_status_string(st), rA.cmp_size, cmpA_host_size, + nA, szgpu_status_string(st), rA.cmp_size, cmpA_host_size, bytes_match ? "equal" : "differs"); fails += report("parity", bytes_match, detail); std::vector outA(nA, 0.0f); - st = fsz_decompress_hostptr(outA.data(), cmpA_host.data(), cmpA_host_size, nA, eb); + st = szgpu_decompress_hostptr(outA.data(), cmpA_host.data(), cmpA_host_size, nA, eb); std::size_t badA = 0; double errA = max_abs_diff(inA, outA, &badA); - bool okA = (st == FSZ_STATUS_OK) && (badA == 0) && (errA <= 1.01 * (double)eb); + bool okA = (st == SZGPU_STATUS_OK) && (badA == 0) && (errA <= 1.01 * (double)eb); std::snprintf(detail, sizeof(detail), "n=%zu status=%s max_err=%.3e bound=%.3e nonfinite=%zu", - nA, fsz_status_string(st), errA, 1.01 * (double)eb, badA); + nA, szgpu_status_string(st), errA, 1.01 * (double)eb, badA); fails += report("roundtrip", okA, detail); const std::size_t nB = 3000000; std::vector inB(nB); fill_near_constant(inB, 1e-4f); - std::vector cmpB_host(fsz::max_compressed_bytes(nB)); + std::vector cmpB_host(szgpu::max_compressed_bytes(nB)); std::size_t cmpB_host_size = 0; - st = fsz_compress_hostptr(inB.data(), cmpB_host.data(), nB, eb, &cmpB_host_size); + st = szgpu_compress_hostptr(inB.data(), cmpB_host.data(), nB, eb, &cmpB_host_size); std::vector outB(nB, 1.0f); - fsz_status_t st_d = FSZ_STATUS_INVALID_ARGUMENT; - if (st == FSZ_STATUS_OK) - st_d = fsz_decompress_hostptr(outB.data(), cmpB_host.data(), cmpB_host_size, + szgpu_status_t st_d = SZGPU_STATUS_INVALID_ARGUMENT; + if (st == SZGPU_STATUS_OK) + st_d = szgpu_decompress_hostptr(outB.data(), cmpB_host.data(), cmpB_host_size, nB, eb); std::size_t badB = 0; double errB = max_abs_diff(inB, outB, &badB); double crB = cmpB_host_size ? (double)(nB * sizeof(float)) / (double)cmpB_host_size : 0.0; - bool okB = (st == FSZ_STATUS_OK) && (st_d == FSZ_STATUS_OK) && (crB > 100.0) + bool okB = (st == SZGPU_STATUS_OK) && (st_d == SZGPU_STATUS_OK) && (crB > 100.0) && (badB == 0) && (errB <= 1.01 * (double)eb); std::snprintf(detail, sizeof(detail), "n=%zu CR=%.2f cmp=%zu B max_err=%.3e nonfinite=%zu", @@ -194,17 +194,17 @@ int main() { std::vector inC(nC); for (std::size_t i = 0; i < nC; ++i) inC[i] = std::sin(0.05f * (float)i) + 0.2f * std::cos(0.31f * (float)i); - std::vector cmpC(fsz::max_compressed_bytes(nC)); + std::vector cmpC(szgpu::max_compressed_bytes(nC)); std::size_t cmpC_size = 0; - st = fsz_compress_hostptr(inC.data(), cmpC.data(), nC, eb, &cmpC_size); + st = szgpu_compress_hostptr(inC.data(), cmpC.data(), nC, eb, &cmpC_size); std::vector outC(nC, 0.0f); - st_d = FSZ_STATUS_INVALID_ARGUMENT; - if (st == FSZ_STATUS_OK) - st_d = fsz_decompress_hostptr(outC.data(), cmpC.data(), cmpC_size, nC, eb); + st_d = SZGPU_STATUS_INVALID_ARGUMENT; + if (st == SZGPU_STATUS_OK) + st_d = szgpu_decompress_hostptr(outC.data(), cmpC.data(), cmpC_size, nC, eb); std::size_t badC = 0; double errC = max_abs_diff(inC, outC, &badC); - std::size_t hdrC = fsz::num_tiles(nC) * 8; - bool okC = (st == FSZ_STATUS_OK) && (st_d == FSZ_STATUS_OK) && (cmpC_size > hdrC) + std::size_t hdrC = szgpu::num_tiles(nC) * 8; + bool okC = (st == SZGPU_STATUS_OK) && (st_d == SZGPU_STATUS_OK) && (cmpC_size > hdrC) && (badC == 0) && (errC <= 1.01 * (double)eb); std::snprintf(detail, sizeof(detail), "n=%zu cmp=%zu B (header %zu B) max_err=%.3e nonfinite=%zu", @@ -212,13 +212,13 @@ int main() { fails += report("small-array", okC, detail); std::vector cmpB_dev; - fsz_compress_result_t rB; + szgpu_compress_result_t rB; if (!compress_device(inB, eb, cmpB_dev, &rB)) { std::printf("[FAIL] setup device-pointer compress of field B failed\n"); return 1; } - fsz_compress_result_t mA = fsz::make_result(nA, rA.cmp_size); - fsz_compress_result_t mB = fsz::make_result(nB, rB.cmp_size); + szgpu_compress_result_t mA = szgpu::make_result(nA, rA.cmp_size); + szgpu_compress_result_t mB = szgpu::make_result(nB, rB.cmp_size); bool okM = same_result(mA, rA) && same_result(mB, rB); std::snprintf(detail, sizeof(detail), "A{%zu,%zu,%zu} B{%zu,%zu,%zu}", @@ -231,22 +231,22 @@ int main() { unsigned char* d_cmp = nullptr; float* d_out = nullptr; cudaMalloc(&d_in, nA * sizeof(float)); - cudaMalloc(&d_cmp, fsz::max_compressed_bytes(nA)); + cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(nA)); cudaMalloc(&d_out, nA * sizeof(float)); cudaMemcpy(d_in, inA.data(), nA * sizeof(float), cudaMemcpyHostToDevice); - fsz_compress_result_t rN{}; - fsz_status_t stN = fsz_compress(d_in, d_cmp, nA, eb, nullptr, 0, &rN); + szgpu_compress_result_t rN{}; + szgpu_status_t stN = szgpu_compress(d_in, d_cmp, nA, eb, nullptr, 0, &rN); std::vector cmpN(rN.cmp_size ? rN.cmp_size : 1); - if (stN == FSZ_STATUS_OK) + if (stN == SZGPU_STATUS_OK) cudaMemcpy(cmpN.data(), d_cmp, rN.cmp_size, cudaMemcpyDeviceToHost); - bool stream_match = (stN == FSZ_STATUS_OK) && (rN.cmp_size == rA.cmp_size) + bool stream_match = (stN == SZGPU_STATUS_OK) && (rN.cmp_size == rA.cmp_size) && std::memcmp(cmpN.data(), cmpA_dev.data(), rA.cmp_size) == 0; - fsz_status_t stD = fsz_decompress(d_out, d_cmp, nA, eb, &rN, nullptr, 0); + szgpu_status_t stD = szgpu_decompress(d_out, d_cmp, nA, eb, &rN, nullptr, 0); std::vector outN(nA, 0.0f); cudaMemcpy(outN.data(), d_out, nA * sizeof(float), cudaMemcpyDeviceToHost); std::size_t badN = 0; double errN = max_abs_diff(inA, outN, &badN); - bool okN = stream_match && (stD == FSZ_STATUS_OK) && (badN == 0) + bool okN = stream_match && (stD == SZGPU_STATUS_OK) && (badN == 0) && (errN <= 1.01 * (double)eb); std::snprintf(detail, sizeof(detail), "cmp=%zu B memcmp=%s max_err=%.3e nonfinite=%zu", @@ -267,19 +267,19 @@ int main() { std::vector inS(ns), outS(ns, 0.0f); for (std::size_t i = 0; i < ns; ++i) inS[i] = (i % 2) ? 4.0e9f : -4.0e9f; - std::vector cmpS(fsz::max_compressed_bytes(ns)); + std::vector cmpS(szgpu::max_compressed_bytes(ns)); std::size_t szS = 0; - fsz_status_t s1 = fsz_compress_hostptr(inS.data(), cmpS.data(), ns, 1.0f, &szS); - fsz_status_t s2 = (s1 == FSZ_STATUS_OK) - ? fsz_decompress_hostptr(outS.data(), cmpS.data(), szS, ns, 1.0f) + szgpu_status_t s1 = szgpu_compress_hostptr(inS.data(), cmpS.data(), ns, 1.0f, &szS); + szgpu_status_t s2 = (s1 == SZGPU_STATUS_OK) + ? szgpu_decompress_hostptr(outS.data(), cmpS.data(), szS, ns, 1.0f) : s1; std::size_t badS = 0; double errS = max_abs_diff(inS, outS, &badS); - if (!(s1 == FSZ_STATUS_OK && s2 == FSZ_STATUS_OK && szS <= cmpS.size() + if (!(s1 == SZGPU_STATUS_OK && s2 == SZGPU_STATUS_OK && szS <= cmpS.size() && badS == 0 && errS <= 1.01)) { okS = false; std::snprintf(worst, sizeof(worst), "n=%zu status=%s err=%.2e", - ns, fsz_status_string(s1 != FSZ_STATUS_OK ? s1 : s2), errS); + ns, szgpu_status_string(s1 != SZGPU_STATUS_OK ? s1 : s2), errS); } } std::snprintf(detail, sizeof(detail), "7 sizes, wide range: %s", @@ -293,52 +293,52 @@ int main() { fill_smooth_f64(inD); std::vector cmpD_dev; - fsz_compress_result_t rD; + szgpu_compress_result_t rD; if (!compress_device_f64(inD, ebd, cmpD_dev, &rD)) { std::printf("[FAIL] setup device-pointer compress of field D failed\n"); return 1; } - std::vector cmpD_host(fsz::max_compressed_bytes(nD)); + std::vector cmpD_host(szgpu::max_compressed_bytes(nD)); std::size_t cmpD_host_size = 0; - st = fsz_compress_hostptr_f64(inD.data(), cmpD_host.data(), nD, ebd, + st = szgpu_compress_hostptr_f64(inD.data(), cmpD_host.data(), nD, ebd, &cmpD_host_size); - bool sizes_match_f64 = (st == FSZ_STATUS_OK) && (cmpD_host_size == rD.cmp_size); + bool sizes_match_f64 = (st == SZGPU_STATUS_OK) && (cmpD_host_size == rD.cmp_size); bool bytes_match_f64 = sizes_match_f64 && std::memcmp(cmpD_host.data(), cmpD_dev.data(), rD.cmp_size) == 0; std::snprintf(detail, sizeof(detail), "n=%zu status=%s device=%zu B hostptr=%zu B memcmp=%s", - nD, fsz_status_string(st), rD.cmp_size, cmpD_host_size, + nD, szgpu_status_string(st), rD.cmp_size, cmpD_host_size, bytes_match_f64 ? "equal" : "differs"); fails += report("parity-f64", bytes_match_f64, detail); std::vector outD(nD, 0.0); - st = fsz_decompress_hostptr_f64(outD.data(), cmpD_host.data(), cmpD_host_size, + st = szgpu_decompress_hostptr_f64(outD.data(), cmpD_host.data(), cmpD_host_size, nD, ebd); std::size_t badD = 0; double errD = max_abs_diff_f64(inD, outD, &badD); - bool okD = (st == FSZ_STATUS_OK) && (badD == 0) && (errD <= 1.01 * ebd); + bool okD = (st == SZGPU_STATUS_OK) && (badD == 0) && (errD <= 1.01 * ebd); std::snprintf(detail, sizeof(detail), "n=%zu status=%s max_err=%.3e bound=%.3e nonfinite=%zu", - nD, fsz_status_string(st), errD, 1.01 * ebd, badD); + nD, szgpu_status_string(st), errD, 1.01 * ebd, badD); fails += report("roundtrip-f64", okD, detail); const std::size_t nE = 4096; std::vector inE(nE); for (std::size_t i = 0; i < nE; ++i) inE[i] = std::sin(0.05 * (double)i) + 0.2 * std::cos(0.31 * (double)i); - std::vector cmpE(fsz::max_compressed_bytes(nE)); + std::vector cmpE(szgpu::max_compressed_bytes(nE)); std::size_t cmpE_size = 0; - st = fsz_compress_hostptr_f64(inE.data(), cmpE.data(), nE, ebd, &cmpE_size); + st = szgpu_compress_hostptr_f64(inE.data(), cmpE.data(), nE, ebd, &cmpE_size); std::vector outE(nE, 0.0); - st_d = FSZ_STATUS_INVALID_ARGUMENT; - if (st == FSZ_STATUS_OK) - st_d = fsz_decompress_hostptr_f64(outE.data(), cmpE.data(), cmpE_size, nE, ebd); + st_d = SZGPU_STATUS_INVALID_ARGUMENT; + if (st == SZGPU_STATUS_OK) + st_d = szgpu_decompress_hostptr_f64(outE.data(), cmpE.data(), cmpE_size, nE, ebd); std::size_t badE = 0; double errE = max_abs_diff_f64(inE, outE, &badE); - std::size_t hdrE = fsz::num_tiles(nE) * 8; - bool okE = (st == FSZ_STATUS_OK) && (st_d == FSZ_STATUS_OK) && (cmpE_size > hdrE) + std::size_t hdrE = szgpu::num_tiles(nE) * 8; + bool okE = (st == SZGPU_STATUS_OK) && (st_d == SZGPU_STATUS_OK) && (cmpE_size > hdrE) && (badE == 0) && (errE <= 1.01 * ebd); std::snprintf(detail, sizeof(detail), "n=%zu cmp=%zu B (header %zu B) max_err=%.3e nonfinite=%zu", diff --git a/tests/test_roundtrip.cu b/tests/test_roundtrip.cu index 5a0acb8..c2f0f72 100644 --- a/tests/test_roundtrip.cu +++ b/tests/test_roundtrip.cu @@ -1,4 +1,4 @@ -#include "fsz/fsz.hpp" +#include "szgpu/szgpu.hpp" #include #include @@ -16,12 +16,12 @@ static int run_case(std::size_t n, float eb, const char* label) { unsigned char* d_cmp = nullptr; cudaMalloc(&d_in, n * sizeof(float)); cudaMalloc(&d_out, n * sizeof(float)); - cudaMalloc(&d_cmp, fsz::max_compressed_bytes(n)); + cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(n)); cudaMemcpy(d_in, h_in.data(), n * sizeof(float), cudaMemcpyHostToDevice); - fsz::Workspace ws(n); - auto r = fsz::compress(d_in, d_cmp, n, eb, ws); - fsz::decompress(d_out, d_cmp, n, eb, r, ws); + szgpu::Workspace ws(n); + auto r = szgpu::compress(d_in, d_cmp, n, eb, ws); + szgpu::decompress(d_out, d_cmp, n, eb, r, ws); cudaDeviceSynchronize(); std::vector h_out(n); @@ -52,12 +52,12 @@ static int run_case_f64(std::size_t n, double eb, const char* label) { unsigned char* d_cmp = nullptr; cudaMalloc(&d_in, n * sizeof(double)); cudaMalloc(&d_out, n * sizeof(double)); - cudaMalloc(&d_cmp, fsz::max_compressed_bytes(n)); + cudaMalloc(&d_cmp, szgpu::max_compressed_bytes(n)); cudaMemcpy(d_in, h_in.data(), n * sizeof(double), cudaMemcpyHostToDevice); - fsz::Workspace ws(n); - auto r = fsz::compress(d_in, d_cmp, n, eb, ws); - fsz::decompress(d_out, d_cmp, n, eb, r, ws); + szgpu::Workspace ws(n); + auto r = szgpu::compress(d_in, d_cmp, n, eb, ws); + szgpu::decompress(d_out, d_cmp, n, eb, r, ws); cudaDeviceSynchronize(); std::vector h_out(n); diff --git a/tools/fsz.cu b/tools/szgpu.cu similarity index 91% rename from tools/fsz.cu rename to tools/szgpu.cu index 1c0e528..fdeb0b9 100644 --- a/tools/fsz.cu +++ b/tools/szgpu.cu @@ -1,8 +1,8 @@ -// fsz: command line front end for the FSZ GPU error-bounded lossy compressor. -// Run 'fsz --help' for usage. +// szgpu: command line front end for the SZ-GPU error-bounded lossy compressor. +// Run 'szgpu --help' for usage. -#include "fsz/fsz.hpp" -#include "fsz_file_format.hpp" +#include "szgpu/szgpu.hpp" +#include "szgpu_file_format.hpp" #include #include @@ -50,7 +50,7 @@ std::string fmt(const char* f, ...) { [[noreturn]] void die_usage(const std::string& msg) { std::fprintf(stderr, "ERROR: %s\n", msg.c_str()); - std::fprintf(stderr, "Run 'fsz --help' for usage.\n"); + std::fprintf(stderr, "Run 'szgpu --help' for usage.\n"); std::exit(RC_ERROR); } @@ -77,23 +77,23 @@ size_t dtype_size(DType t) { return (t == DType::F64) ? sizeof(double) : si void usage(std::FILE* out) { std::fprintf(out, - "FSZ %s GPU error-bounded lossy compressor for float32 and float64 arrays\n" + "SZ-GPU %s error-bounded lossy compressor for float32 and float64 arrays\n" "\n" "Usage:\n" - " fsz -i data.f32 -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" + " szgpu -i data.f32 -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" " compress and decompress in memory, then report ratio,\n" " throughput and reconstruction error (writes nothing)\n" "\n" - " fsz -z -i data.f32 -o data.fsz -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" - " compress to a self-describing .fsz file\n" + " szgpu -z -i data.f32 -o data.szgpu -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" + " compress to a self-describing .szgpu file\n" "\n" - " fsz -t f64 -z -i data.d64 -o data.fsz -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" + " szgpu -t f64 -z -i data.d64 -o data.szgpu -eb rel 1e-3 [-d D1 [D2 [D3]]]\n" " the same for a float64 input array\n" "\n" - " fsz -x -i data.fsz -o data.out.f32\n" - " decompress a .fsz file, no other arguments needed\n" + " szgpu -x -i data.szgpu -o data.out.f32\n" + " decompress a .szgpu file, no other arguments needed\n" "\n" - " fsz -x -i stream.bin -o out.f32 -n N -eb abs V --bare\n" + " szgpu -x -i stream.bin -o out.f32 -n N -eb abs V --bare\n" " decompress a bare bitstream written with --bare\n" "\n" "Options:\n" @@ -127,10 +127,10 @@ void usage(std::FILE* out) { "there as it does for compression.\n" "\n" "With neither -z nor -x the tool inspects the first four bytes of the\n" - "input: an FSZ container is decompressed, anything else is reported.\n" + "input: an SZ-GPU container is decompressed, anything else is reported.\n" "\n" "Exit status: 0 pass, 1 error bound violated, 2 usage or file error.\n", - fsz::version()); + szgpu::version()); } bool parse_u64(const char* s, uint64_t& out) { @@ -295,8 +295,8 @@ void parse_dim_list(int argc, char** argv, int& i, std::vector& dims) die_usage("bad dimension list '" + first + "'; expected D1xD2xD3 with positive integers"); } - if (dims.size() == fsz_file::MAX_DIMS) { - die_usage("at most " + std::to_string(fsz_file::MAX_DIMS) + + if (dims.size() == szgpu_file::MAX_DIMS) { + die_usage("at most " + std::to_string(szgpu_file::MAX_DIMS) + " dimensions are supported, got '" + first + "'"); } dims.push_back(v); @@ -313,7 +313,7 @@ void parse_dim_list(int argc, char** argv, int& i, std::vector& dims) die_usage("bad dimension '" + first + "'; expected a positive integer"); } dims.push_back(v); - while (dims.size() < fsz_file::MAX_DIMS && i + 1 < argc && + while (dims.size() < szgpu_file::MAX_DIMS && i + 1 < argc && parse_u64(argv[i + 1], v)) { if (v == 0) die_usage("dimensions must be positive"); dims.push_back(v); @@ -501,19 +501,19 @@ int run_compress(const Options& o, bool write_output) { T* d_in = nullptr; T* d_out = nullptr; unsigned char* d_cmp = nullptr; - const size_t cmp_alloc = fsz::max_compressed_bytes(n); + const size_t cmp_alloc = szgpu::max_compressed_bytes(n); check_cuda(cudaMalloc(&d_in, n * sizeof(T)), "allocating the input buffer"); check_cuda(cudaMalloc(&d_out, n * sizeof(T)), "allocating the output buffer"); check_cuda(cudaMalloc(&d_cmp, cmp_alloc), "allocating the compressed buffer"); check_cuda(cudaMemcpy(d_in, h_in.data(), n * sizeof(T), cudaMemcpyHostToDevice), "copying the input to the device"); - fsz::Workspace ws(n); + szgpu::Workspace ws(n); - fsz_compress_result_t r{}; + szgpu_compress_result_t r{}; for (int i = 0; i < WARMUP_ITERS; ++i) { - r = fsz::compress(d_in, d_cmp, n, eb_abs, ws); - fsz::decompress(d_out, d_cmp, n, eb_abs, r, ws); + r = szgpu::compress(d_in, d_cmp, n, eb_abs, ws); + szgpu::decompress(d_out, d_cmp, n, eb_abs, r, ws); } check_cuda(cudaDeviceSynchronize(), "warmup"); @@ -521,13 +521,13 @@ int run_compress(const Options& o, bool write_output) { double c_ms_sum = 0.0; for (int i = 0; i < TIMED_ITERS; ++i) { t.start(); - r = fsz::compress(d_in, d_cmp, n, eb_abs, ws); + r = szgpu::compress(d_in, d_cmp, n, eb_abs, ws); c_ms_sum += t.stop(); } double d_ms_sum = 0.0; for (int i = 0; i < TIMED_ITERS; ++i) { t.start(); - fsz::decompress(d_out, d_cmp, n, eb_abs, r, ws); + szgpu::decompress(d_out, d_cmp, n, eb_abs, r, ws); d_ms_sum += t.stop(); } check_cuda(cudaGetLastError(), "the timed loop"); @@ -580,8 +580,8 @@ int run_compress(const Options& o, bool write_output) { write_all(o.out_path, h_cmp.data(), h_cmp.size()); total_bytes = (uint64_t)h_cmp.size(); } else { - const fsz_file::Header hdr = - fsz_file::make_header(dims, ndims, (double)eb_abs, is_f64, + const szgpu_file::Header hdr = + szgpu_file::make_header(dims, ndims, (double)eb_abs, is_f64, (uint64_t)r.cmp_size); std::FILE* f = std::fopen(o.out_path.c_str(), "wb"); if (!f) die("cannot open output " + o.out_path); @@ -601,7 +601,7 @@ int run_compress(const Options& o, bool write_output) { const size_t in_bytes = n * sizeof(T); const double cr = (double)n * (double)sizeof(T) / (double)r.cmp_size; - std::printf("FSZ %s\n", fsz::version()); + std::printf("SZ-GPU %s\n", szgpu::version()); std::printf(" input = %s (%.2f MB)\n", o.in_path.c_str(), (double)in_bytes / (1024.0 * 1024.0)); std::printf(" dtype = %s\n", TypeTraits::name); @@ -617,7 +617,7 @@ int run_compress(const Options& o, bool write_output) { } else { std::printf(" output = %s (%llu bytes total, %zu byte header)\n", o.out_path.c_str(), (unsigned long long)total_bytes, - sizeof(fsz_file::Header)); + sizeof(szgpu_file::Header)); } } std::printf(" compress = %.2f ms (%.2f GB/s)\n", c_ms, gbps(in_bytes, c_ms)); @@ -665,10 +665,10 @@ int run_decompress_typed(const Options& o, const DecompressPlan& p) { const size_t n = (size_t)p.n; const T eb_abs = (T)p.eb_abs; - fsz_compress_result_t r{}; + szgpu_compress_result_t r{}; r.cmp_size = p.payload_size; - r.num_tiles = fsz::num_tiles(n); - r.data_offset = fsz::num_tiles(n) * 8; + r.num_tiles = szgpu::num_tiles(n); + r.data_offset = szgpu::num_tiles(n) * 8; if (p.payload_size < r.data_offset) { die(fmt("the bitstream is %zu bytes but %zu elements need at least a " @@ -678,17 +678,17 @@ int run_decompress_typed(const Options& o, const DecompressPlan& p) { T* d_out = nullptr; unsigned char* d_cmp = nullptr; - size_t cmp_alloc = fsz::max_compressed_bytes(n); + size_t cmp_alloc = szgpu::max_compressed_bytes(n); if (cmp_alloc < p.payload_size) cmp_alloc = p.payload_size; check_cuda(cudaMalloc(&d_out, n * sizeof(T)), "allocating the output buffer"); check_cuda(cudaMalloc(&d_cmp, cmp_alloc), "allocating the compressed buffer"); check_cuda(cudaMemcpy(d_cmp, p.payload, p.payload_size, cudaMemcpyHostToDevice), "copying the bitstream to the device"); - fsz::Workspace ws(n); + szgpu::Workspace ws(n); for (int i = 0; i < WARMUP_ITERS; ++i) { - fsz::decompress(d_out, d_cmp, n, eb_abs, r, ws); + szgpu::decompress(d_out, d_cmp, n, eb_abs, r, ws); } check_cuda(cudaDeviceSynchronize(), "warmup"); @@ -696,7 +696,7 @@ int run_decompress_typed(const Options& o, const DecompressPlan& p) { double d_ms_sum = 0.0; for (int i = 0; i < TIMED_ITERS; ++i) { t.start(); - fsz::decompress(d_out, d_cmp, n, eb_abs, r, ws); + szgpu::decompress(d_out, d_cmp, n, eb_abs, r, ws); d_ms_sum += t.stop(); } check_cuda(cudaGetLastError(), "the timed loop"); @@ -714,7 +714,7 @@ int run_decompress_typed(const Options& o, const DecompressPlan& p) { const size_t out_bytes = n * sizeof(T); const double cr = (double)n * (double)sizeof(T) / (double)p.payload_size; - std::printf("FSZ %s\n", fsz::version()); + std::printf("SZ-GPU %s\n", szgpu::version()); std::printf(" input = %s (%zu bytes)\n", o.in_path.c_str(), p.raw_size); std::printf(" dtype = %s%s\n", TypeTraits::name, p.from_header ? " (from the file header)" : ""); @@ -787,26 +787,26 @@ int run_decompress(const Options& o) { payload = raw.data(); payload_size = raw.size(); } else { - if (raw.size() < fsz_file::HEADER_SIZE) { - die(fmt("%s is %zu bytes, too short to hold a %zu byte FSZ header; " + if (raw.size() < szgpu_file::HEADER_SIZE) { + die(fmt("%s is %zu bytes, too short to hold a %zu byte SZ-GPU header; " "if this is a bare bitstream, pass --bare with -n and -eb abs", - o.in_path.c_str(), raw.size(), fsz_file::HEADER_SIZE)); + o.in_path.c_str(), raw.size(), szgpu_file::HEADER_SIZE)); } - fsz_file::Header hdr{}; + szgpu_file::Header hdr{}; std::memcpy(&hdr, raw.data(), sizeof(hdr)); - if (!fsz_file::has_magic(hdr.magic)) { - die(o.in_path + " is not an FSZ container; if this is a bare bitstream, " + if (!szgpu_file::has_magic(hdr.magic)) { + die(o.in_path + " is not an SZ-GPU container; if this is a bare bitstream, " "pass --bare with -n and -eb abs"); } - if (hdr.version != fsz_file::FORMAT_VER) { + if (hdr.version != szgpu_file::FORMAT_VER) { die(fmt("%s uses container version %u, this build reads version %u", - o.in_path.c_str(), hdr.version, fsz_file::FORMAT_VER)); + o.in_path.c_str(), hdr.version, szgpu_file::FORMAT_VER)); } - if (hdr.ndims < 1 || hdr.ndims > fsz_file::MAX_DIMS) { + if (hdr.ndims < 1 || hdr.ndims > szgpu_file::MAX_DIMS) { die(fmt("%s records %u dimensions, which is out of range", o.in_path.c_str(), hdr.ndims)); } - const uint64_t have = (uint64_t)(raw.size() - fsz_file::HEADER_SIZE); + const uint64_t have = (uint64_t)(raw.size() - szgpu_file::HEADER_SIZE); if (hdr.payload_size > have) { die(fmt("%s claims a %llu byte bitstream but only %llu bytes follow " "the header; the file looks truncated", @@ -814,7 +814,7 @@ int run_decompress(const Options& o) { (unsigned long long)hdr.payload_size, (unsigned long long)have)); } - dtype = fsz_file::header_is_f64(hdr) ? DType::F64 : DType::F32; + dtype = szgpu_file::header_is_f64(hdr) ? DType::F64 : DType::F32; if (o.have_dtype && o.dtype != dtype) { die(fmt("%s records %s data but -t asks for %s; the container " "already carries its element type, so leave -t out", @@ -824,9 +824,9 @@ int run_decompress(const Options& o) { dims[0] = hdr.dims[0]; dims[1] = hdr.dims[1]; dims[2] = hdr.dims[2]; - n64 = fsz_file::element_count(hdr); - eb_abs = fsz_file::header_eb(hdr); - payload = raw.data() + fsz_file::HEADER_SIZE; + n64 = szgpu_file::element_count(hdr); + eb_abs = szgpu_file::header_eb(hdr); + payload = raw.data() + szgpu_file::HEADER_SIZE; payload_size = (size_t)hdr.payload_size; from_header = true; } @@ -851,13 +851,13 @@ int run_decompress(const Options& o) { bool looks_like_container(const std::string& path) { const long bytes = file_size_or_die(path); - if (bytes < (long)fsz_file::HEADER_SIZE) return false; + if (bytes < (long)szgpu_file::HEADER_SIZE) return false; std::FILE* f = std::fopen(path.c_str(), "rb"); if (!f) die("cannot open input " + path); char magic[4] = {0, 0, 0, 0}; const bool got = std::fread(magic, 1, 4, f) == 4; std::fclose(f); - return got && fsz_file::has_magic(magic); + return got && szgpu_file::has_magic(magic); } } // namespace @@ -877,7 +877,7 @@ int main(int argc, char** argv) { if (o.bare) die_usage("--bare applies to -z and -x; pick one"); if (looks_like_container(o.in_path)) return run_decompress(o); return f64 ? run_compress(o, false) : run_compress(o, false); - } catch (const fsz::Error& e) { + } catch (const szgpu::Error& e) { std::fprintf(stderr, "ERROR: %s\n", e.what()); return RC_ERROR; } catch (const std::exception& e) { diff --git a/tools/fsz_file_format.hpp b/tools/szgpu_file_format.hpp similarity index 91% rename from tools/fsz_file_format.hpp rename to tools/szgpu_file_format.hpp index d5941a9..3cbb704 100644 --- a/tools/fsz_file_format.hpp +++ b/tools/szgpu_file_format.hpp @@ -1,13 +1,13 @@ -#ifndef FSZ_TOOLS_FILE_FORMAT_HPP -#define FSZ_TOOLS_FILE_FORMAT_HPP +#ifndef SZGPU_TOOLS_FILE_FORMAT_HPP +#define SZGPU_TOOLS_FILE_FORMAT_HPP -// On-disk container written by the fsz command line tool. +// On-disk container written by the szgpu command line tool. #include #include #include -namespace fsz_file { +namespace szgpu_file { constexpr char MAGIC[4] = {'F', 'S', 'Z', '1'}; constexpr uint32_t FORMAT_VER = 1; @@ -29,7 +29,7 @@ struct Header { uint64_t payload_size; }; -static_assert(sizeof(Header) == HEADER_SIZE, "FSZ file header must be 56 bytes"); +static_assert(sizeof(Header) == HEADER_SIZE, "SZ-GPU file header must be 56 bytes"); static_assert(offsetof(Header, magic) == 0, "bad magic offset"); static_assert(offsetof(Header, version) == 4, "bad version offset"); static_assert(offsetof(Header, ndims) == 8, "bad ndims offset");