diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index db20bdbc1..b73a8eea9 100644 --- a/.github/workflows/build.yaml +++ b/.github/workflows/build.yaml @@ -28,7 +28,7 @@ concurrency: jobs: cpp-build: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -37,7 +37,7 @@ jobs: rust-build: needs: cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -50,7 +50,7 @@ jobs: python-build: needs: [cpp-build] secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -59,7 +59,7 @@ jobs: upload-conda: needs: [cpp-build, python-build] secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-upload-packages.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-upload-packages.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -70,7 +70,7 @@ jobs: if: github.ref_type == 'branch' needs: python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@python-3.12 with: arch: "amd64" branch: ${{ inputs.branch }} @@ -82,7 +82,7 @@ jobs: sha: ${{ inputs.sha }} wheel-build-cuvs: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -92,7 +92,7 @@ jobs: wheel-publish-cuvs: needs: wheel-build-cuvs secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@python-3.12 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 8ea2fa503..d34f74062 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -24,49 +24,49 @@ jobs: - wheel-tests-cuvs - devcontainer secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@python-3.12 checks: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@python-3.12 with: enable_check_generated_files: false conda-cpp-build: needs: checks secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@python-3.12 with: build_type: pull-request node_type: cpu16 conda-cpp-tests: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@python-3.12 with: build_type: pull-request conda-cpp-checks: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@python-3.12 with: build_type: pull-request enable_check_symbols: true - symbol_exclusions: (void (thrust::|cub::)|_ZN\d+raft_cutlass) + symbol_exclusions: (void (thrust::|cub::)|raft_cutlass) conda-python-build: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@python-3.12 with: build_type: pull-request conda-python-tests: needs: conda-python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@python-3.12 with: build_type: pull-request docs-build: needs: conda-python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@python-3.12 with: build_type: pull-request node_type: "gpu-v100-latest-1" @@ -76,7 +76,7 @@ jobs: rust-build: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@python-3.12 with: build_type: pull-request node_type: "gpu-v100-latest-1" @@ -86,20 +86,20 @@ jobs: wheel-build-cuvs: needs: checks secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@python-3.12 with: build_type: pull-request script: ci/build_wheel_cuvs.sh wheel-tests-cuvs: needs: wheel-build-cuvs secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@python-3.12 with: build_type: pull-request script: ci/test_wheel_cuvs.sh devcontainer: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@python-3.12 with: arch: '["amd64"]' cuda: '["12.5"]' diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index c0d07297b..f2daecbec 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -16,17 +16,17 @@ on: jobs: conda-cpp-checks: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-post-build-checks.yaml@python-3.12 with: build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} sha: ${{ inputs.sha }} enable_check_symbols: true - symbol_exclusions: (void (thrust::|cub::)|_ZN\d+raft_cutlass) + symbol_exclusions: (void (thrust::|cub::)|raft_cutlass) conda-cpp-tests: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@python-3.12 with: build_type: nightly branch: ${{ inputs.branch }} @@ -34,7 +34,7 @@ jobs: sha: ${{ inputs.sha }} conda-python-tests: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@python-3.12 with: build_type: nightly branch: ${{ inputs.branch }} @@ -42,7 +42,7 @@ jobs: sha: ${{ inputs.sha }} wheel-tests-cuvs: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.10 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@python-3.12 with: build_type: nightly branch: ${{ inputs.branch }} diff --git a/.gitignore b/.gitignore index 4b6f46320..fcbe0fa3a 100644 --- a/.gitignore +++ b/.gitignore @@ -28,6 +28,11 @@ bench/ann/data temporary_*.json rust/target/ rust/Cargo.lock +rmm_log.txt + +## example notebooks +notebooks/simplewiki-2020-11-01-nq-distilbert-base-v1.pt +notebooks/data/ ## scikit-build _skbuild diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 592693adb..3e3623f24 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -97,7 +97,7 @@ repos: hooks: - id: check-json - repo: https://github.com/rapidsai/pre-commit-hooks - rev: v0.3.1 + rev: v0.4.0 hooks: - id: verify-copyright files: | diff --git a/ci/build_wheel.sh b/ci/build_wheel.sh index 5d48bab22..d1030276f 100755 --- a/ci/build_wheel.sh +++ b/ci/build_wheel.sh @@ -16,10 +16,26 @@ rapids-generate-version > ./VERSION cd "${package_dir}" +case "${RAPIDS_CUDA_VERSION}" in + 12.*) + EXCLUDE_ARGS=( + --exclude "libcublas.so.12" + --exclude "libcublasLt.so.12" + --exclude "libcurand.so.10" + --exclude "libcusolver.so.11" + --exclude "libcusparse.so.12" + --exclude "libnvJitLink.so.12" + ) + ;; + 11.*) + EXCLUDE_ARGS=() + ;; +esac + # Hardcode the output dir python -m pip wheel . -w dist -vvv --no-deps --disable-pip-version-check mkdir -p final_dist -python -m auditwheel repair -w final_dist dist/* +python -m auditwheel repair -w final_dist "${EXCLUDE_ARGS[@]}" dist/* RAPIDS_PY_WHEEL_NAME="${underscore_package_name}_${RAPIDS_PY_CUDA_SUFFIX}" rapids-upload-wheels-to-s3 final_dist diff --git a/ci/build_wheel_cuvs.sh b/ci/build_wheel_cuvs.sh index 0fe28c2f7..e03da9f19 100755 --- a/ci/build_wheel_cuvs.sh +++ b/ci/build_wheel_cuvs.sh @@ -3,7 +3,16 @@ set -euo pipefail +case "${RAPIDS_CUDA_VERSION}" in + 12.*) + EXTRA_CMAKE_ARGS=";-DUSE_CUDA_MATH_WHEELS=ON" + ;; + 11.*) + EXTRA_CMAKE_ARGS=";-DUSE_CUDA_MATH_WHEELS=OFF" + ;; +esac + # Set up skbuild options. Enable sccache in skbuild config options -export SKBUILD_CMAKE_ARGS="-DDETECT_CONDA_ENV=OFF -DFIND_CUVS_CPP=OFF" +export SKBUILD_CMAKE_ARGS="-DDETECT_CONDA_ENV=OFF;-DFIND_CUVS_CPP=OFF${EXTRA_CMAKE_ARGS}" ci/build_wheel.sh cuvs python/cuvs diff --git a/conda/environments/all_cuda-118_arch-aarch64.yaml b/conda/environments/all_cuda-118_arch-aarch64.yaml index 54866d0e5..4bbdc3650 100644 --- a/conda/environments/all_cuda-118_arch-aarch64.yaml +++ b/conda/environments/all_cuda-118_arch-aarch64.yaml @@ -38,7 +38,7 @@ dependencies: - make - nccl>=2.9.9 - ninja -- numpy>=1.23,<2.0a0 +- numpy>=1.23,<3.0a0 - numpydoc - nvcc_linux-aarch64=11.8 - openblas diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index b29fe1b50..908421d08 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -38,7 +38,7 @@ dependencies: - make - nccl>=2.9.9 - ninja -- numpy>=1.23,<2.0a0 +- numpy>=1.23,<3.0a0 - numpydoc - nvcc_linux-64=11.8 - openblas diff --git a/conda/environments/all_cuda-125_arch-aarch64.yaml b/conda/environments/all_cuda-125_arch-aarch64.yaml index c6deb93f8..3131d0b77 100644 --- a/conda/environments/all_cuda-125_arch-aarch64.yaml +++ b/conda/environments/all_cuda-125_arch-aarch64.yaml @@ -35,7 +35,7 @@ dependencies: - make - nccl>=2.9.9 - ninja -- numpy>=1.23,<2.0a0 +- numpy>=1.23,<3.0a0 - numpydoc - openblas - pre-commit diff --git a/conda/environments/all_cuda-125_arch-x86_64.yaml b/conda/environments/all_cuda-125_arch-x86_64.yaml index 538fdf08b..2f107c4fb 100644 --- a/conda/environments/all_cuda-125_arch-x86_64.yaml +++ b/conda/environments/all_cuda-125_arch-x86_64.yaml @@ -35,7 +35,7 @@ dependencies: - make - nccl>=2.9.9 - ninja -- numpy>=1.23,<2.0a0 +- numpy>=1.23,<3.0a0 - numpydoc - openblas - pre-commit diff --git a/conda/environments/bench_ann_cuda-118_arch-aarch64.yaml b/conda/environments/bench_ann_cuda-118_arch-aarch64.yaml new file mode 100644 index 000000000..9b23faa67 --- /dev/null +++ b/conda/environments/bench_ann_cuda-118_arch-aarch64.yaml @@ -0,0 +1,47 @@ +# This file is generated by `rapids-dependency-file-generator`. +# To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. +channels: +- rapidsai +- rapidsai-nightly +- dask/label/dev +- conda-forge +- nvidia +dependencies: +- benchmark>=1.8.2 +- c-compiler +- clang-tools=16.0.6 +- clang==16.0.6 +- click +- cmake>=3.26.4,!=3.30.0 +- cuda-nvtx=11.8 +- cuda-profiler-api=11.8.86 +- cuda-python>=11.7.1,<12.0a0 +- cuda-version=11.8 +- cudatoolkit +- cxx-compiler +- cython>=3.0.0 +- dlpack>=0.8,<1.0 +- gcc_linux-aarch64=11.* +- glog>=0.6.0 +- h5py>=3.8.0 +- hnswlib=0.6.2 +- libcublas-dev=11.11.3.6 +- libcublas=11.11.3.6 +- libcurand-dev=10.3.0.86 +- libcurand=10.3.0.86 +- libcusolver-dev=11.4.1.48 +- libcusolver=11.4.1.48 +- libcusparse-dev=11.7.5.86 +- libcusparse=11.7.5.86 +- matplotlib +- nccl>=2.9.9 +- ninja +- nlohmann_json>=3.11.2 +- nvcc_linux-aarch64=11.8 +- openblas +- pandas +- pylibraft==24.10.*,>=0.0.0a0 +- pyyaml +- rmm==24.10.*,>=0.0.0a0 +- sysroot_linux-aarch64==2.17 +name: bench_ann_cuda-118_arch-aarch64 diff --git a/conda/environments/bench_ann_cuda-118_arch-x86_64.yaml b/conda/environments/bench_ann_cuda-118_arch-x86_64.yaml new file mode 100644 index 000000000..e73efd60c --- /dev/null +++ b/conda/environments/bench_ann_cuda-118_arch-x86_64.yaml @@ -0,0 +1,47 @@ +# This file is generated by `rapids-dependency-file-generator`. +# To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. +channels: +- rapidsai +- rapidsai-nightly +- dask/label/dev +- conda-forge +- nvidia +dependencies: +- benchmark>=1.8.2 +- c-compiler +- clang-tools=16.0.6 +- clang==16.0.6 +- click +- cmake>=3.26.4,!=3.30.0 +- cuda-nvtx=11.8 +- cuda-profiler-api=11.8.86 +- cuda-python>=11.7.1,<12.0a0 +- cuda-version=11.8 +- cudatoolkit +- cxx-compiler +- cython>=3.0.0 +- dlpack>=0.8,<1.0 +- gcc_linux-64=11.* +- glog>=0.6.0 +- h5py>=3.8.0 +- hnswlib=0.6.2 +- libcublas-dev=11.11.3.6 +- libcublas=11.11.3.6 +- libcurand-dev=10.3.0.86 +- libcurand=10.3.0.86 +- libcusolver-dev=11.4.1.48 +- libcusolver=11.4.1.48 +- libcusparse-dev=11.7.5.86 +- libcusparse=11.7.5.86 +- matplotlib +- nccl>=2.9.9 +- ninja +- nlohmann_json>=3.11.2 +- nvcc_linux-64=11.8 +- openblas +- pandas +- pylibraft==24.10.*,>=0.0.0a0 +- pyyaml +- rmm==24.10.*,>=0.0.0a0 +- sysroot_linux-64==2.17 +name: bench_ann_cuda-118_arch-x86_64 diff --git a/conda/environments/bench_ann_cuda-125_arch-aarch64.yaml b/conda/environments/bench_ann_cuda-125_arch-aarch64.yaml new file mode 100644 index 000000000..dec41007a --- /dev/null +++ b/conda/environments/bench_ann_cuda-125_arch-aarch64.yaml @@ -0,0 +1,43 @@ +# This file is generated by `rapids-dependency-file-generator`. +# To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. +channels: +- rapidsai +- rapidsai-nightly +- dask/label/dev +- conda-forge +- nvidia +dependencies: +- benchmark>=1.8.2 +- c-compiler +- clang-tools=16.0.6 +- clang==16.0.6 +- click +- cmake>=3.26.4,!=3.30.0 +- cuda-cudart-dev +- cuda-nvcc +- cuda-nvtx-dev +- cuda-profiler-api +- cuda-python>=12.0,<13.0a0 +- cuda-version=12.5 +- cxx-compiler +- cython>=3.0.0 +- dlpack>=0.8,<1.0 +- gcc_linux-aarch64=11.* +- glog>=0.6.0 +- h5py>=3.8.0 +- hnswlib=0.6.2 +- libcublas-dev +- libcurand-dev +- libcusolver-dev +- libcusparse-dev +- matplotlib +- nccl>=2.9.9 +- ninja +- nlohmann_json>=3.11.2 +- openblas +- pandas +- pylibraft==24.10.*,>=0.0.0a0 +- pyyaml +- rmm==24.10.*,>=0.0.0a0 +- sysroot_linux-aarch64==2.17 +name: bench_ann_cuda-125_arch-aarch64 diff --git a/conda/environments/bench_ann_cuda-125_arch-x86_64.yaml b/conda/environments/bench_ann_cuda-125_arch-x86_64.yaml new file mode 100644 index 000000000..f106644cd --- /dev/null +++ b/conda/environments/bench_ann_cuda-125_arch-x86_64.yaml @@ -0,0 +1,43 @@ +# This file is generated by `rapids-dependency-file-generator`. +# To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. +channels: +- rapidsai +- rapidsai-nightly +- dask/label/dev +- conda-forge +- nvidia +dependencies: +- benchmark>=1.8.2 +- c-compiler +- clang-tools=16.0.6 +- clang==16.0.6 +- click +- cmake>=3.26.4,!=3.30.0 +- cuda-cudart-dev +- cuda-nvcc +- cuda-nvtx-dev +- cuda-profiler-api +- cuda-python>=12.0,<13.0a0 +- cuda-version=12.5 +- cxx-compiler +- cython>=3.0.0 +- dlpack>=0.8,<1.0 +- gcc_linux-64=11.* +- glog>=0.6.0 +- h5py>=3.8.0 +- hnswlib=0.6.2 +- libcublas-dev +- libcurand-dev +- libcusolver-dev +- libcusparse-dev +- matplotlib +- nccl>=2.9.9 +- ninja +- nlohmann_json>=3.11.2 +- openblas +- pandas +- pylibraft==24.10.*,>=0.0.0a0 +- pyyaml +- rmm==24.10.*,>=0.0.0a0 +- sysroot_linux-64==2.17 +name: bench_ann_cuda-125_arch-x86_64 diff --git a/conda/recipes/cuvs/meta.yaml b/conda/recipes/cuvs/meta.yaml index df6cfeec6..2633b3db8 100644 --- a/conda/recipes/cuvs/meta.yaml +++ b/conda/recipes/cuvs/meta.yaml @@ -69,6 +69,8 @@ requirements: - libcuvs {{ version }} - python x.x - rmm ={{ minor_version }} + - cuda-python + - numpy >=1.23,<3.0a0 tests: requirements: diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 76fd09b58..ba46e60b4 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -424,6 +424,8 @@ add_library( src/selection/select_k_float_int64_t.cu src/selection/select_k_float_uint32_t.cu src/selection/select_k_half_uint32_t.cu + src/stats/silhouette_score.cu + src/stats/trustworthiness_score.cu ) target_compile_definitions(cuvs PRIVATE "CUVS_EXPLICIT_INSTANTIATE_ONLY") @@ -549,6 +551,7 @@ if(BUILD_C_LIBRARY) src/neighbors/ivf_flat_c.cpp src/neighbors/ivf_pq_c.cpp src/neighbors/cagra_c.cpp + src/neighbors/hnsw_c.cpp src/neighbors/refine/refine_c.cpp src/distance/pairwise_distance_c.cpp ) diff --git a/cpp/cmake/thirdparty/get_raft.cmake b/cpp/cmake/thirdparty/get_raft.cmake index fb25623ce..7640fbfa6 100644 --- a/cpp/cmake/thirdparty/get_raft.cmake +++ b/cpp/cmake/thirdparty/get_raft.cmake @@ -61,8 +61,8 @@ endfunction() # To use a different RAFT locally, set the CMake variable # CPM_raft_SOURCE=/path/to/local/raft find_and_configure_raft(VERSION ${RAFT_VERSION}.00 - FORK rhdong - PINNED_TAG half_knn + FORK ${RAFT_FORK} + PINNED_TAG ${RAFT_PINNED_TAG} ENABLE_MNMG_DEPENDENCIES OFF ENABLE_NVTX OFF USE_RAFT_STATIC ${CUVS_USE_RAFT_STATIC} diff --git a/cpp/include/cuvs/core/c_api.h b/cpp/include/cuvs/core/c_api.h index d931d6c13..4db7fd12c 100644 --- a/cpp/include/cuvs/core/c_api.h +++ b/cpp/include/cuvs/core/c_api.h @@ -127,6 +127,27 @@ cuvsError_t cuvsRMMAlloc(cuvsResources_t res, void** ptr, size_t bytes); */ cuvsError_t cuvsRMMFree(cuvsResources_t res, void* ptr, size_t bytes); +/** + * @brief Switches the working memory resource to use the RMM pool memory resource, which will + * bypass unnecessary synchronizations by allocating a chunk of device memory up front and carving + * that up for temporary memory allocations within algorithms. Be aware that this function will + * change the memory resource for the whole process and the new memory resource will be used until + * explicitly changed. + * + * @param[in] initial_pool_size_percent The initial pool size as a percentage of the total + * available memory + * @param[in] max_pool_size_percent The maximum pool size as a percentage of the total + * available memory + * @return cuvsError_t + */ +cuvsError_t cuvsRMMPoolMemoryResourceEnable(int initial_pool_size_percent, + int max_pool_size_percent); +/** + * @brief Resets the memory resource to use the default memory resource (cuda_memory_resource) + * @return cuvsError_t + */ +cuvsError_t cuvsRMMMemoryResourceReset(); + /** @} */ #ifdef __cplusplus diff --git a/cpp/include/cuvs/neighbors/cagra.h b/cpp/include/cuvs/neighbors/cagra.h index 87541f7f0..241f5d8b0 100644 --- a/cpp/include/cuvs/neighbors/cagra.h +++ b/cpp/include/cuvs/neighbors/cagra.h @@ -337,7 +337,10 @@ cuvsError_t cuvsCagraBuild(cuvsResources_t res, * It is also important to note that the CAGRA Index must have been built * with the same type of `queries`, such that `index.dtype.code == * queries.dl_tensor.dtype.code` Types for input are: - * 1. `queries`: `kDLDataType.code == kDLFloat` and `kDLDataType.bits = 32` + * 1. `queries`: + *` a. kDLDataType.code == kDLFloat` and `kDLDataType.bits = 32` + * b. `kDLDataType.code == kDLInt` and `kDLDataType.bits = 8` + * c. `kDLDataType.code == kDLUInt` and `kDLDataType.bits = 8` * 2. `neighbors`: `kDLDataType.code == kDLUInt` and `kDLDataType.bits = 32` * 3. `distances`: `kDLDataType.code == kDLFloat` and `kDLDataType.bits = 32` * @@ -394,7 +397,7 @@ cuvsError_t cuvsCagraSearch(cuvsResources_t res, * * Experimental, both the API and the serialization format are subject to change. * - * @code{.cpp} + * @code{.c} * #include * * // Create cuvsResources_t @@ -416,6 +419,34 @@ cuvsError_t cuvsCagraSerialize(cuvsResources_t res, cuvsCagraIndex_t index, bool include_dataset); +/** + * Save the CAGRA index to file in hnswlib format. + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. + * + * Experimental, both the API and the serialization format are subject to change. + * + * @code{.c} + * #include + * #include + * + * // Create cuvsResources_t + * cuvsResources_t res; + * cuvsError_t res_create_status = cuvsResourcesCreate(&res); + * + * // create an index with `cuvsCagraBuild` + * cuvsCagraSerializeHnswlib(res, "/path/to/index", index); + * @endcode + * + * @param[in] res cuvsResources_t opaque C handle + * @param[in] filename the file name for saving the index + * @param[in] index CAGRA index + * + */ +cuvsError_t cuvsCagraSerializeToHnswlib(cuvsResources_t res, + const char* filename, + cuvsCagraIndex_t index); + /** * Load index from file. * diff --git a/cpp/include/cuvs/neighbors/cagra.hpp b/cpp/include/cuvs/neighbors/cagra.hpp index f74eac711..5f77eb8a3 100644 --- a/cpp/include/cuvs/neighbors/cagra.hpp +++ b/cpp/include/cuvs/neighbors/cagra.hpp @@ -1345,6 +1345,8 @@ void deserialize(raft::resources const& handle, /** * Write the CAGRA built index as a base layer HNSW index to an output stream + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * @@ -1371,6 +1373,8 @@ void serialize_to_hnswlib(raft::resources const& handle, /** * Save a CAGRA build index in hnswlib base-layer-only serialized format + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * @@ -1398,6 +1402,8 @@ void serialize_to_hnswlib(raft::resources const& handle, /** * Write the CAGRA built index as a base layer HNSW index to an output stream + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * @@ -1424,6 +1430,8 @@ void serialize_to_hnswlib(raft::resources const& handle, /** * Save a CAGRA build index in hnswlib base-layer-only serialized format + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * @@ -1451,6 +1459,8 @@ void serialize_to_hnswlib(raft::resources const& handle, /** * Write the CAGRA built index as a base layer HNSW index to an output stream + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * @@ -1477,6 +1487,8 @@ void serialize_to_hnswlib(raft::resources const& handle, /** * Save a CAGRA build index in hnswlib base-layer-only serialized format + * NOTE: The saved index can only be read by the hnswlib wrapper in cuVS, + * as the serialization format is not compatible with the original hnswlib. * * Experimental, both the API and the serialization format are subject to change. * diff --git a/cpp/include/cuvs/neighbors/hnsw.h b/cpp/include/cuvs/neighbors/hnsw.h new file mode 100644 index 000000000..5e94de60a --- /dev/null +++ b/cpp/include/cuvs/neighbors/hnsw.h @@ -0,0 +1,207 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include +#include +#include + +#ifdef __cplusplus +extern "C" { +#endif + +/** + * @defgroup hnsw_c_search_params C API for hnswlib wrapper search params + * @{ + */ + +struct cuvsHnswSearchParams { + int32_t ef; + int32_t numThreads; +}; + +typedef struct cuvsHnswSearchParams* cuvsHnswSearchParams_t; + +/** + * @brief Allocate HNSW search params, and populate with default values + * + * @param[in] params cuvsHnswSearchParams_t to allocate + * @return cuvsError_t + */ +cuvsError_t cuvsHnswSearchParamsCreate(cuvsHnswSearchParams_t* params); + +/** + * @brief De-allocate HNSW search params + * + * @param[in] params cuvsHnswSearchParams_t to de-allocate + * @return cuvsError_t + */ +cuvsError_t cuvsHnswSearchParamsDestroy(cuvsHnswSearchParams_t params); + +/** + * @} + */ + +/** + * @defgroup hnsw_c_index C API for hnswlib wrapper index + * @{ + */ + +/** + * @brief Struct to hold address of cuvs::neighbors::Hnsw::index and its active trained dtype + * + */ +typedef struct { + uintptr_t addr; + DLDataType dtype; + +} cuvsHnswIndex; + +typedef cuvsHnswIndex* cuvsHnswIndex_t; + +/** + * @brief Allocate HNSW index + * + * @param[in] index cuvsHnswIndex_t to allocate + * @return HnswError_t + */ +cuvsError_t cuvsHnswIndexCreate(cuvsHnswIndex_t* index); + +/** + * @brief De-allocate HNSW index + * + * @param[in] index cuvsHnswIndex_t to de-allocate + */ +cuvsError_t cuvsHnswIndexDestroy(cuvsHnswIndex_t index); + +/** + * @} + */ + +/** + * @defgroup hnsw_c_index_search C API for CUDA ANN Graph-based nearest neighbor search + * @{ + */ +/** + * @brief Search a HNSW index with a `DLManagedTensor` which has underlying + * `DLDeviceType` equal to `kDLCPU`, `kDLCUDAHost`, or `kDLCUDAManaged`. + * It is also important to note that the HNSW Index must have been built + * with the same type of `queries`, such that `index.dtype.code == + * queries.dl_tensor.dtype.code` + * Supported types for input are: + * 1. `queries`: `kDLDataType.code == kDLFloat` or `kDLDataType.code == kDLInt` and + * `kDLDataType.bits = 32` + * 2. `neighbors`: `kDLDataType.code == kDLUInt` and `kDLDataType.bits = 64` + * 3. `distances`: `kDLDataType.code == kDLFloat` and `kDLDataType.bits = 32` + * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS, + * as the format is not compatible with the original hnswlib. + * + * @code {.c} + * #include + * #include + * + * // Create cuvsResources_t + * cuvsResources_t res; + * cuvsError_t res_create_status = cuvsResourcesCreate(&res); + * + * // Assume a populated `DLManagedTensor` type here + * DLManagedTensor dataset; + * DLManagedTensor queries; + * DLManagedTensor neighbors; + * + * // Create default search params + * cuvsHnswSearchParams_t params; + * cuvsError_t params_create_status = cuvsHnswSearchParamsCreate(¶ms); + * + * // Search the `index` built using `cuvsHnswBuild` + * cuvsError_t search_status = cuvsHnswSearch(res, params, index, &queries, &neighbors, + * &distances); + * + * // de-allocate `params` and `res` + * cuvsError_t params_destroy_status = cuvsHnswSearchParamsDestroy(params); + * cuvsError_t res_destroy_status = cuvsResourcesDestroy(res); + * @endcode + * + * @param[in] res cuvsResources_t opaque C handle + * @param[in] params cuvsHnswSearchParams_t used to search Hnsw index + * @param[in] index cuvsHnswIndex which has been returned by `cuvsHnswBuild` + * @param[in] queries DLManagedTensor* queries dataset to search + * @param[out] neighbors DLManagedTensor* output `k` neighbors for queries + * @param[out] distances DLManagedTensor* output `k` distances for queries + */ +cuvsError_t cuvsHnswSearch(cuvsResources_t res, + cuvsHnswSearchParams_t params, + cuvsHnswIndex_t index, + DLManagedTensor* queries, + DLManagedTensor* neighbors, + DLManagedTensor* distances); + +/** + * @} + */ + +/** + * @defgroup hnsw_c_serialize HNSW C-API serialize functions + * @{ + */ + +/** + * Load hnswlib index from file which was serialized from a HNSW index. + * NOTE: The loaded hnswlib index is immutable, and only be read by the + * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. + * Experimental, both the API and the serialization format are subject to change. + * + * @code{.c} + * #include + * #include + * #include + * + * // Create cuvsResources_t + * cuvsResources_t res; + * cuvsError_t res_create_status = cuvsResourcesCreate(&res); + * + * // create an index with `cuvsCagraBuild` + * cuvsCagraSerializeHnswlib(res, "/path/to/index", index); + * + * // Load the serialized CAGRA index from file as an hnswlib index + * // The index should have the same dtype as the one used to build CAGRA the index + * cuvsHnswIndex_t hnsw_index; + * cuvsHnswIndexCreate(&hnsw_index); + * hnsw_index->dtype = index->dtype; + * cuvsCagraDeserialize(res, "/path/to/index", hnsw_index); + * @endcode + * + * @param[in] res cuvsResources_t opaque C handle + * @param[in] filename the name of the file that stores the index + * @param[in] dim the dimension of the vectors in the index + * @param[in] metric the distance metric used to build the index + * @param[out] index HNSW index loaded disk + */ +cuvsError_t cuvsHnswDeserialize(cuvsResources_t res, + const char* filename, + int dim, + cuvsDistanceType metric, + cuvsHnswIndex_t index); +/** + * @} + */ + +#ifdef __cplusplus +} +#endif diff --git a/cpp/include/cuvs/neighbors/hnsw.hpp b/cpp/include/cuvs/neighbors/hnsw.hpp index 86f321564..007adef0d 100644 --- a/cpp/include/cuvs/neighbors/hnsw.hpp +++ b/cpp/include/cuvs/neighbors/hnsw.hpp @@ -34,7 +34,7 @@ namespace cuvs::neighbors::hnsw { /** - * @defgroup hnsw Build CAGRA index and search with hnswlib + * @defgroup hnsw_cpp_search_params Build CAGRA index and search with hnswlib * @{ */ @@ -44,6 +44,13 @@ struct search_params : cuvs::neighbors::search_params { // automatically maximizes parallelism }; +/**@}*/ + +/** + * @defgroup hnsw_cpp_index hnswlib index wrapper + * @{ + */ + template struct index : cuvs::neighbors::index { public: @@ -58,6 +65,8 @@ struct index : cuvs::neighbors::index { */ index(int dim, cuvs::distance::DistanceType metric) : dim_{dim}, metric_{metric} {} + virtual ~index() {} + /** @brief Get underlying index */ @@ -77,11 +86,19 @@ struct index : cuvs::neighbors::index { cuvs::distance::DistanceType metric_; }; +/**@}*/ + +/** + * @defgroup hnsw_cpp_index_load Load CAGRA index as hnswlib index + * @{ + */ + /** - * @brief Construct an hnswlib base-layer-only index from a CAGRA index - * NOTE: 1. This method uses the filesystem to write the CAGRA index in `/tmp/.bin` - * before reading it as an hnswlib index, then deleting the temporary file. - * 2. This function is only offered as a compiled symbol in `libraft.so` + * @brief Construct an immutable hnswlib base-layer-only index from a CAGRA index + * NOTE: This method uses the filesystem to write the CAGRA index in `/tmp/.bin` + * before reading it as an hnswlib index, then deleting the temporary file. The returned index + * is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not + * compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] cagra_index cagra index @@ -103,10 +120,11 @@ std::unique_ptr> from_cagra( raft::resources const& res, const cuvs::neighbors::cagra::index& cagra_index); /** - * @brief Construct an hnswlib base-layer-only index from a CAGRA index - * NOTE: 1. This method uses the filesystem to write the CAGRA index in `/tmp/.bin` - * before reading it as an hnswlib index, then deleting the temporary file. - * 2. This function is only offered as a compiled symbol in `libraft.so` + * @brief Construct an immutable hnswlib base-layer-only index from a CAGRA index + * NOTE: This method uses the filesystem to write the CAGRA index in `/tmp/.bin` + * before reading it as an hnswlib index, then deleting the temporary file. The returned index + * is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not + * compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] cagra_index cagra index @@ -128,10 +146,11 @@ std::unique_ptr> from_cagra( raft::resources const& res, const cuvs::neighbors::cagra::index& cagra_index); /** - * @brief Construct an hnswlib base-layer-only index from a CAGRA index - * NOTE: 1. This method uses the filesystem to write the CAGRA index in `/tmp/.bin` - * before reading it as an hnswlib index, then deleting the temporary file. - * 2. This function is only offered as a compiled symbol in `libraft.so` + * @brief Construct an immutable hnswlib base-layer-only index from a CAGRA index + * NOTE: This method uses the filesystem to write the CAGRA index in `/tmp/.bin` + * before reading it as an hnswlib index, then deleting the temporary file. The returned index + * is immutable and can only be searched by the hnswlib wrapper in cuVS, as the format is not + * compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] cagra_index cagra index @@ -152,8 +171,17 @@ std::unique_ptr> from_cagra( std::unique_ptr> from_cagra( raft::resources const& res, const cuvs::neighbors::cagra::index& cagra_index); +/**@}*/ + +/** + * @defgroup hnsw_cpp_index_search Search hnswlib index + * @{ + */ + /** * @brief Search hnswlib base-layer-only index constructed from a CAGRA index + * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS, + * as the format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] params configure the search @@ -195,6 +223,8 @@ void search(raft::resources const& res, /** * @brief Search hnswlib base-layer-only index constructed from a CAGRA index + * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS, + * as the format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] params configure the search @@ -236,6 +266,8 @@ void search(raft::resources const& res, /** * @brief Search hnswlib base-layer-only index constructed from a CAGRA index + * NOTE: The HNSW index can only be searched by the hnswlib wrapper in cuVS, + * as the format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] params configure the search @@ -275,8 +307,17 @@ void search(raft::resources const& res, raft::host_matrix_view neighbors, raft::host_matrix_view distances); +/**@}*/ + +/** + * @defgroup hnsw_cpp_index_deserialize Deserialize CAGRA index as hnswlib index + * @{ + */ + /** - * @brief De-serialize a CAGRA index saved to a file as an hnsw index + * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * NOTE: The loaded hnswlib index is immutable, and only be read by the + * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] filename path to the file containing the serialized CAGRA index @@ -310,7 +351,9 @@ void deserialize(raft::resources const& res, index** index); /** - * @brief De-serialize a CAGRA index saved to a file as an hnsw index + * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * NOTE: The loaded hnswlib index is immutable, and only be read by the + * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] filename path to the file containing the serialized CAGRA index @@ -344,7 +387,9 @@ void deserialize(raft::resources const& res, index** index); /** - * @brief De-serialize a CAGRA index saved to a file as an hnsw index + * @brief De-serialize a CAGRA index saved to a file as an hnswlib index + * NOTE: The loaded hnswlib index is immutable, and only be read by the + * hnswlib wrapper in cuVS, as the serialization format is not compatible with the original hnswlib. * * @param[in] res raft resources * @param[in] filename path to the file containing the serialized CAGRA index diff --git a/cpp/include/cuvs/stats/silhouette_score.hpp b/cpp/include/cuvs/stats/silhouette_score.hpp new file mode 100644 index 000000000..1771cc21c --- /dev/null +++ b/cpp/include/cuvs/stats/silhouette_score.hpp @@ -0,0 +1,121 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#pragma once + +#include +#include +#include + +namespace cuvs { +namespace stats { + +/** + * @defgroup stats_silhouette_score Silhouette Score + * @{ + */ +/** + * @brief main function that returns the average silhouette score for a given set of data and its + * clusterings + * @param[in] handle: raft handle for managing expensive resources + * @param[in] X_in: input matrix Data in row-major format (nRows x nCols) + * @param[in] labels: the pointer to the array containing labels for every data sample (length: + * nRows) + * @param[out] silhouette_score_per_sample: optional array populated with the silhouette score + * for every sample (length: nRows) + * @param[in] n_unique_labels: number of unique labels in the labels array + * @param[in] metric: Distance metric to use. Euclidean (L2) is used by default + * @return: The silhouette score. + */ +float silhouette_score( + raft::resources const& handle, + raft::device_matrix_view X_in, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded); + +/** + * @brief function that returns the average silhouette score for a given set of data and its + * clusterings + * @param[in] handle: raft handle for managing expensive resources + * @param[in] X: input matrix Data in row-major format (nRows x nCols) + * @param[in] labels: the pointer to the array containing labels for every data sample (length: + * nRows) + * @param[out] silhouette_score_per_sample: optional array populated with the silhouette score + * for every sample (length: nRows) + * @param[in] n_unique_labels: number of unique labels in the labels array + * @param[in] batch_size: number of samples per batch + * @param[in] metric: the numerical value that maps to the type of distance metric to be used in + * the calculations + * @return: The silhouette score. + */ +float silhouette_score_batched( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + int64_t batch_size, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded); + +/** + * @brief main function that returns the average silhouette score for a given set of data and its + * clusterings + * @param[in] handle: raft handle for managing expensive resources + * @param[in] X_in: input matrix Data in row-major format (nRows x nCols) + * @param[in] labels: the pointer to the array containing labels for every data sample (length: + * nRows) + * @param[out] silhouette_score_per_sample: optional array populated with the silhouette score + * for every sample (length: nRows) + * @param[in] n_unique_labels: number of unique labels in the labels array + * @param[in] metric: the numerical value that maps to the type of distance metric to be used in + * the calculations + * @return: The silhouette score. + */ +double silhouette_score( + raft::resources const& handle, + raft::device_matrix_view X_in, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded); + +/** + * @brief function that returns the average silhouette score for a given set of data and its + * clusterings + * @param[in] handle: raft handle for managing expensive resources + * @param[in] X: input matrix Data in row-major format (nRows x nCols) + * @param[in] labels: the pointer to the array containing labels for every data sample (length: + * nRows) + * @param[out] silhouette_score_per_sample: optional array populated with the silhouette score + * for every sample (length: nRows) + * @param[in] n_unique_labels: number of unique labels in the labels array + * @param[in] batch_size: number of samples per batch + * @param[in] metric: the numerical value that maps to the type of distance metric to be used in + * the calculations + * @return: The silhouette score. + */ +double silhouette_score_batched( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + int64_t batch_size, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded); + +} // namespace stats +} // namespace cuvs diff --git a/cpp/include/cuvs/stats/trustworthiness_score.hpp b/cpp/include/cuvs/stats/trustworthiness_score.hpp new file mode 100644 index 000000000..08a26fad5 --- /dev/null +++ b/cpp/include/cuvs/stats/trustworthiness_score.hpp @@ -0,0 +1,51 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once +#include +#include +#include + +namespace cuvs { +namespace stats { +/** + * @defgroup stats_trustworthiness Trustworthiness + * @{ + */ + +/** + * @brief Compute the trustworthiness score + * @param[in] handle the raft handle + * @param[in] X: Data in original dimension + * @param[in] X_embedded: Data in target dimension (embedding) + * @param[in] n_neighbors Number of neighbors considered by trustworthiness score + * @param[in] metric Distance metric to use. Euclidean (L2) is used by default + * @param[in] batch_size Batch size + * @return Trustworthiness score + * @note The constness of the data in X_embedded is currently casted away and the data is slightly + * modified. + */ +double trustworthiness_score( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_matrix_view X_embedded, + int n_neighbors, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2SqrtUnexpanded, + int batch_size = 512); + +/** @} */ // end group stats_trustworthiness +} // namespace stats +} // namespace cuvs diff --git a/cpp/src/core/c_api.cpp b/cpp/src/core/c_api.cpp index 96504a2fe..a75e5a1dd 100644 --- a/cpp/src/core/c_api.cpp +++ b/cpp/src/core/c_api.cpp @@ -22,6 +22,7 @@ #include #include #include +#include #include extern "C" cuvsError_t cuvsResourcesCreate(cuvsResources_t* res) @@ -82,6 +83,34 @@ extern "C" cuvsError_t cuvsRMMFree(cuvsResources_t res, void* ptr, size_t bytes) }); } +thread_local std::unique_ptr> pool_mr; + +extern "C" cuvsError_t cuvsRMMPoolMemoryResourceEnable(int initial_pool_size_percent, + int max_pool_size_percent) +{ + return cuvs::core::translate_exceptions([=] { + // Upstream memory resource needs to be a cuda_memory_resource + auto cuda_mr = rmm::mr::get_current_device_resource(); + auto* cuda_mr_casted = dynamic_cast(cuda_mr); + if (cuda_mr_casted == nullptr) { + throw std::runtime_error("Current memory resource is not a cuda_memory_resource"); + } + auto initial_size = rmm::percent_of_free_device_memory(initial_pool_size_percent); + auto max_size = rmm::percent_of_free_device_memory(max_pool_size_percent); + pool_mr = std::make_unique>( + cuda_mr_casted, initial_size, max_size); + rmm::mr::set_current_device_resource(pool_mr.get()); + }); +} + +extern "C" cuvsError_t cuvsRMMMemoryResourceReset() +{ + return cuvs::core::translate_exceptions([=] { + rmm::mr::set_current_device_resource(nullptr); + pool_mr.reset(); + }); +} + thread_local std::string last_error_text = ""; extern "C" const char* cuvsGetLastErrorText() diff --git a/cpp/src/neighbors/cagra_c.cpp b/cpp/src/neighbors/cagra_c.cpp index 868b3dec0..164448f2c 100644 --- a/cpp/src/neighbors/cagra_c.cpp +++ b/cpp/src/neighbors/cagra_c.cpp @@ -130,6 +130,14 @@ void _serialize(cuvsResources_t res, cuvs::neighbors::cagra::serialize(*res_ptr, std::string(filename), *index_ptr, include_dataset); } +template +void _serialize_to_hnswlib(cuvsResources_t res, const char* filename, cuvsCagraIndex_t index) +{ + auto res_ptr = reinterpret_cast(res); + auto index_ptr = reinterpret_cast*>(index->addr); + cuvs::neighbors::cagra::serialize_to_hnswlib(*res_ptr, std::string(filename), *index_ptr); +} + template void* _deserialize(cuvsResources_t res, const char* filename) { @@ -326,3 +334,20 @@ extern "C" cuvsError_t cuvsCagraSerialize(cuvsResources_t res, } }); } + +extern "C" cuvsError_t cuvsCagraSerializeToHnswlib(cuvsResources_t res, + const char* filename, + cuvsCagraIndex_t index) +{ + return cuvs::core::translate_exceptions([=] { + if (index->dtype.code == kDLFloat && index->dtype.bits == 32) { + _serialize_to_hnswlib(res, filename, index); + } else if (index->dtype.code == kDLInt && index->dtype.bits == 8) { + _serialize_to_hnswlib(res, filename, index); + } else if (index->dtype.code == kDLUInt && index->dtype.bits == 8) { + _serialize_to_hnswlib(res, filename, index); + } else { + RAFT_FAIL("Unsupported index dtype: %d and bits: %d", index->dtype.code, index->dtype.bits); + } + }); +} diff --git a/cpp/src/neighbors/cagra_optimize.cu b/cpp/src/neighbors/cagra_optimize.cu index 4c152cf12..cba66a5e9 100644 --- a/cpp/src/neighbors/cagra_optimize.cu +++ b/cpp/src/neighbors/cagra_optimize.cu @@ -23,13 +23,19 @@ void optimize(raft::resources const& handle, raft::device_matrix_view knn_graph, raft::host_matrix_view new_graph) { - cuvs::neighbors::cagra::optimize(handle, knn_graph, new_graph); + cuvs::neighbors::cagra::optimize< + uint32_t, + raft::host_device_accessor, + raft::memory_type::device>>(handle, knn_graph, new_graph); } void optimize(raft::resources const& handle, raft::host_matrix_view knn_graph, raft::host_matrix_view new_graph) { - cuvs::neighbors::cagra::optimize(handle, knn_graph, new_graph); + cuvs::neighbors::cagra::optimize< + uint32_t, + raft::host_device_accessor, + raft::memory_type::host>>(handle, knn_graph, new_graph); } } // namespace cuvs::neighbors::cagra \ No newline at end of file diff --git a/cpp/src/neighbors/hnsw_c.cpp b/cpp/src/neighbors/hnsw_c.cpp new file mode 100644 index 000000000..ab5268a6d --- /dev/null +++ b/cpp/src/neighbors/hnsw_c.cpp @@ -0,0 +1,166 @@ + +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "cuvs/distance/distance.h" +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +namespace { +template +void _search(cuvsResources_t res, + cuvsHnswSearchParams params, + cuvsHnswIndex index, + DLManagedTensor* queries_tensor, + DLManagedTensor* neighbors_tensor, + DLManagedTensor* distances_tensor) +{ + auto res_ptr = reinterpret_cast(res); + auto index_ptr = reinterpret_cast*>(index.addr); + + auto search_params = cuvs::neighbors::hnsw::search_params(); + search_params.ef = params.ef; + search_params.num_threads = params.numThreads; + + using queries_mdspan_type = raft::host_matrix_view; + using neighbors_mdspan_type = raft::host_matrix_view; + using distances_mdspan_type = raft::host_matrix_view; + auto queries_mds = cuvs::core::from_dlpack(queries_tensor); + auto neighbors_mds = cuvs::core::from_dlpack(neighbors_tensor); + auto distances_mds = cuvs::core::from_dlpack(distances_tensor); + cuvs::neighbors::hnsw::search( + *res_ptr, search_params, *index_ptr, queries_mds, neighbors_mds, distances_mds); +} + +template +void* _deserialize(cuvsResources_t res, const char* filename, int dim, cuvsDistanceType metric) +{ + auto res_ptr = reinterpret_cast(res); + cuvs::neighbors::hnsw::index* index = nullptr; + cuvs::neighbors::hnsw::deserialize(*res_ptr, std::string(filename), dim, metric, &index); + return index; +} +} // namespace + +extern "C" cuvsError_t cuvsHnswSearchParamsCreate(cuvsHnswSearchParams_t* params) +{ + return cuvs::core::translate_exceptions( + [=] { *params = new cuvsHnswSearchParams{.ef = 200, .numThreads = 0}; }); +} + +extern "C" cuvsError_t cuvsHnswSearchParamsDestroy(cuvsHnswSearchParams_t params) +{ + return cuvs::core::translate_exceptions([=] { delete params; }); +} + +extern "C" cuvsError_t cuvsHnswIndexCreate(cuvsHnswIndex_t* index) +{ + return cuvs::core::translate_exceptions([=] { *index = new cuvsHnswIndex{}; }); +} + +extern "C" cuvsError_t cuvsHnswIndexDestroy(cuvsHnswIndex_t index_c_ptr) +{ + return cuvs::core::translate_exceptions([=] { + auto index = *index_c_ptr; + + if (index.dtype.code == kDLFloat) { + auto index_ptr = reinterpret_cast*>(index.addr); + delete index_ptr; + } else if (index.dtype.code == kDLInt) { + auto index_ptr = reinterpret_cast*>(index.addr); + delete index_ptr; + } else if (index.dtype.code == kDLUInt) { + auto index_ptr = reinterpret_cast*>(index.addr); + delete index_ptr; + } + delete index_c_ptr; + }); +} + +extern "C" cuvsError_t cuvsHnswSearch(cuvsResources_t res, + cuvsHnswSearchParams_t params, + cuvsHnswIndex_t index_c_ptr, + DLManagedTensor* queries_tensor, + DLManagedTensor* neighbors_tensor, + DLManagedTensor* distances_tensor) +{ + return cuvs::core::translate_exceptions([=] { + auto queries = queries_tensor->dl_tensor; + auto neighbors = neighbors_tensor->dl_tensor; + auto distances = distances_tensor->dl_tensor; + + RAFT_EXPECTS(cuvs::core::is_dlpack_host_compatible(queries), + "queries should have host compatible memory"); + RAFT_EXPECTS(cuvs::core::is_dlpack_host_compatible(neighbors), + "neighbors should have host compatible memory"); + RAFT_EXPECTS(cuvs::core::is_dlpack_host_compatible(distances), + "distances should have host compatible memory"); + + RAFT_EXPECTS(neighbors.dtype.code == kDLUInt && neighbors.dtype.bits == 64, + "neighbors should be of type uint64_t"); + RAFT_EXPECTS(distances.dtype.code == kDLFloat && distances.dtype.bits == 32, + "distances should be of type float32"); + + auto index = *index_c_ptr; + RAFT_EXPECTS(queries.dtype.code == index.dtype.code, "type mismatch between index and queries"); + RAFT_EXPECTS(queries.dtype.bits == 32, "number of bits in queries dtype should be 32"); + + if (index.dtype.code == kDLFloat) { + _search( + res, *params, index, queries_tensor, neighbors_tensor, distances_tensor); + } else if (index.dtype.code == kDLUInt) { + _search( + res, *params, index, queries_tensor, neighbors_tensor, distances_tensor); + } else if (index.dtype.code == kDLInt) { + _search(res, *params, index, queries_tensor, neighbors_tensor, distances_tensor); + } else { + RAFT_FAIL("Unsupported index dtype: %d and bits: %d", queries.dtype.code, queries.dtype.bits); + } + }); +} + +extern "C" cuvsError_t cuvsHnswDeserialize(cuvsResources_t res, + const char* filename, + int dim, + cuvsDistanceType metric, + cuvsHnswIndex_t index) +{ + return cuvs::core::translate_exceptions([=] { + if (index->dtype.code == kDLFloat && index->dtype.bits == 32) { + index->addr = reinterpret_cast(_deserialize(res, filename, dim, metric)); + index->dtype.code = kDLFloat; + } else if (index->dtype.code == kDLUInt && index->dtype.bits == 8) { + index->addr = reinterpret_cast(_deserialize(res, filename, dim, metric)); + index->dtype.code = kDLInt; + } else if (index->dtype.code == kDLInt && index->dtype.bits == 8) { + index->addr = reinterpret_cast(_deserialize(res, filename, dim, metric)); + index->dtype.code = kDLUInt; + } else { + RAFT_FAIL("Unsupported dtype in file %s", filename); + } + }); +} diff --git a/cpp/src/stats/detail/batched/silhouette_score.cuh b/cpp/src/stats/detail/batched/silhouette_score.cuh new file mode 100644 index 000000000..aae14ea6c --- /dev/null +++ b/cpp/src/stats/detail/batched/silhouette_score.cuh @@ -0,0 +1,285 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "../silhouette_score.cuh" + +#include +#include +#include +#include +#include + +#include +#include + +#include +#include +#include + +namespace cuvs { +namespace stats { +namespace batched { +namespace detail { + +/** + * This kernel initializes matrix b (n_rows * n_labels) + * For each label that the corresponding row is not a part of is initialized as 0 + * If the corresponding row is the only sample in its label, again 0 + * Only if the there are > 1 samples in the label, row is initialized to max + */ +template +RAFT_KERNEL fill_b_kernel(value_t* b, + const label_idx* y, + value_idx n_rows, + label_idx n_labels, + const value_idx* cluster_counts) +{ + value_idx idx = threadIdx.x + blockIdx.x * blockDim.x; + label_idx idy = threadIdx.y + blockIdx.y * blockDim.y; + + if (idx >= n_rows || idy >= n_labels) { return; } + + auto row_cluster = y[idx]; + + auto col_cluster_count = cluster_counts[idy]; + + // b for own cluster should be max value + // so that it does not interfere with min operator + // b is also max if col cluster count is 0 + // however, b is 0 if self cluster count is 1 + if (row_cluster == idy || col_cluster_count == 0) { + if (cluster_counts[row_cluster] == 1) { + b[idx * n_labels + idy] = 0; + } else { + b[idx * n_labels + idy] = std::numeric_limits::max(); + } + } else { + b[idx * n_labels + idy] = 0; + } +} + +/** + * This kernel does an elementwise sweep of chunked pairwise distance matrix + * By knowing the offsets of the chunked pairwise distance matrix in the + * global pairwise distance matrix, we are able to calculate + * intermediate values of a and b for the rows and columns present in the + * current chunked pairwise distance matrix. + */ +template +RAFT_KERNEL compute_chunked_a_b_kernel(value_t* a, + value_t* b, + value_idx row_offset, + value_idx col_offset, + const label_idx* y, + label_idx n_labels, + const value_idx* cluster_counts, + const value_t* distances, + value_idx dist_rows, + value_idx dist_cols) +{ + value_idx row_id = threadIdx.x + blockIdx.x * blockDim.x; + value_idx col_id = threadIdx.y + blockIdx.y * blockDim.y; + + // these are global offsets of current element + // in the full pairwise distance matrix + value_idx pw_row_id = row_id + row_offset; + value_idx pw_col_id = col_id + col_offset; + + if (row_id >= dist_rows || col_id >= dist_cols || pw_row_id == pw_col_id) { return; } + + auto row_cluster = y[pw_row_id]; + if (cluster_counts[row_cluster] == 1) { return; } + + auto col_cluster = y[pw_col_id]; + auto col_cluster_counts = cluster_counts[col_cluster]; + + if (col_cluster == row_cluster) { + atomicAdd(&a[pw_row_id], distances[row_id * dist_cols + col_id] / (col_cluster_counts - 1)); + } else { + atomicAdd(&b[pw_row_id * n_labels + col_cluster], + distances[row_id * dist_cols + col_id] / col_cluster_counts); + } +} + +template +rmm::device_uvector get_cluster_counts(raft::resources const& handle, + const label_idx* y, + value_idx& n_rows, + label_idx& n_labels) +{ + auto stream = raft::resource::get_cuda_stream(handle); + + rmm::device_uvector cluster_counts(n_labels, stream); + + rmm::device_uvector workspace(1, stream); + + cuvs::stats::detail::countLabels(y, cluster_counts.data(), n_rows, n_labels, workspace, stream); + + return cluster_counts; +} + +template +rmm::device_uvector get_pairwise_distance(raft::resources const& handle, + const value_t* left_begin, + const value_t* right_begin, + value_idx& n_left_rows, + value_idx& n_right_rows, + value_idx& n_cols, + cuvs::distance::DistanceType metric, + cudaStream_t stream) +{ + rmm::device_uvector distances(n_left_rows * n_right_rows, stream); + + cuvs::distance::pairwise_distance( + handle, + raft::make_device_matrix_view(left_begin, n_left_rows, n_cols), + raft::make_device_matrix_view(right_begin, n_right_rows, n_cols), + raft::make_device_matrix_view(distances.data(), n_left_rows, n_right_rows), + metric); + + return distances; +} + +template +void compute_chunked_a_b(raft::resources const& handle, + value_t* a, + value_t* b, + value_idx& row_offset, + value_idx& col_offset, + const label_idx* y, + label_idx& n_labels, + const value_idx* cluster_counts, + const value_t* distances, + value_idx& dist_rows, + value_idx& dist_cols, + cudaStream_t stream) +{ + dim3 block_size(std::min(dist_rows, 32), std::min(dist_cols, 32)); + dim3 grid_size(raft::ceildiv(dist_rows, (value_idx)block_size.x), + raft::ceildiv(dist_cols, (value_idx)block_size.y)); + + detail::compute_chunked_a_b_kernel<<>>( + a, b, row_offset, col_offset, y, n_labels, cluster_counts, distances, dist_rows, dist_cols); +} + +template +value_t silhouette_score( + raft::resources const& handle, + const value_t* X, + value_idx n_rows, + value_idx n_cols, + const label_idx* y, + label_idx n_labels, + value_t* scores, + value_idx chunk, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded) +{ + ASSERT(n_labels >= 2 && n_labels <= (n_rows - 1), + "silhouette Score not defined for the given number of labels!"); + + rmm::device_uvector cluster_counts = get_cluster_counts(handle, y, n_rows, n_labels); + + auto stream = raft::resource::get_cuda_stream(handle); + auto policy = raft::resource::get_thrust_policy(handle); + + auto b_size = n_rows * n_labels; + + value_t *a_ptr, *b_ptr; + rmm::device_uvector a(0, stream); + rmm::device_uvector b(b_size, stream); + + b_ptr = b.data(); + + // since a and silhouette score per sample are same size, reusing + if (scores == nullptr || scores == NULL) { + a.resize(n_rows, stream); + a_ptr = a.data(); + } else { + a_ptr = scores; + } + + thrust::fill(policy, a_ptr, a_ptr + n_rows, 0); + + dim3 block_size(std::min(n_rows, 32), std::min(n_labels, 32)); + dim3 grid_size(raft::ceildiv(n_rows, (value_idx)block_size.x), + raft::ceildiv(n_labels, (label_idx)block_size.y)); + detail::fill_b_kernel<<>>( + b_ptr, y, n_rows, n_labels, cluster_counts.data()); + + raft::resource::wait_stream_pool_on_stream(handle); + + auto n_iters = 0; + + for (value_idx i = 0; i < n_rows; i += chunk) { + for (value_idx j = 0; j < n_rows; j += chunk) { + ++n_iters; + + auto chunk_stream = raft::resource::get_next_usable_stream(handle, i + chunk * j); + + const auto* left_begin = X + (i * n_cols); + const auto* right_begin = X + (j * n_cols); + + auto n_left_rows = (i + chunk) < n_rows ? chunk : (n_rows - i); + auto n_right_rows = (j + chunk) < n_rows ? chunk : (n_rows - j); + + rmm::device_uvector distances = get_pairwise_distance( + handle, left_begin, right_begin, n_left_rows, n_right_rows, n_cols, metric, chunk_stream); + + compute_chunked_a_b(handle, + a_ptr, + b_ptr, + i, + j, + y, + n_labels, + cluster_counts.data(), + distances.data(), + n_left_rows, + n_right_rows, + chunk_stream); + } + } + + raft::resource::sync_stream_pool(handle); + + // calculating row-wise minimum in b + // this prim only supports int indices for now + raft::linalg::reduce( + b_ptr, + b_ptr, + n_labels, + n_rows, + std::numeric_limits::max(), + true, + true, + stream, + false, + raft::identity_op(), + raft::min_op()); + + // calculating the silhouette score per sample + raft::linalg::binaryOp, value_t, value_idx>( + a_ptr, a_ptr, b_ptr, n_rows, cuvs::stats::detail::SilOp(), stream); + + return thrust::reduce(policy, a_ptr, a_ptr + n_rows, value_t(0)) / n_rows; +} + +} // namespace detail +} // namespace batched +} // namespace stats +} // namespace cuvs diff --git a/cpp/src/stats/detail/silhouette_score.cuh b/cpp/src/stats/detail/silhouette_score.cuh new file mode 100644 index 000000000..6c876ce6e --- /dev/null +++ b/cpp/src/stats/detail/silhouette_score.cuh @@ -0,0 +1,328 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include + +#include +#include +#include + +namespace cuvs { +namespace stats { +namespace detail { + +/** + * @brief kernel that calculates the average intra-cluster distance for every sample data point and + * updates the cluster distance to max value + * @tparam DataT: type of the data samples + * @tparam LabelT: type of the labels + * @param sampleToClusterSumOfDistances: the pointer to the 2D array that contains the sum of + * distances from every sample to every cluster (nRows x nLabels) + * @param binCountArray: pointer to the 1D array that contains the count of samples per cluster (1 x + * nLabels) + * @param d_aArray: the pointer to the array of average intra-cluster distances for every sample in + * device memory (1 x nRows) + * @param labels: the pointer to the array containing labels for every data sample (1 x nRows) + * @param nRows: number of data samples + * @param nLabels: number of Labels + * @param MAX_VAL: DataT specific upper limit + */ +template +RAFT_KERNEL populateAKernel(DataT* sampleToClusterSumOfDistances, + DataT* binCountArray, + DataT* d_aArray, + const LabelT* labels, + int nRows, + int nLabels, + const DataT MAX_VAL) +{ + // getting the current index + int sampleIndex = threadIdx.x + blockIdx.x * blockDim.x; + + if (sampleIndex >= nRows) return; + + // sampleDistanceVector is an array that stores that particular row of the distanceMatrix + DataT* sampleToClusterSumOfDistancesVector = + &sampleToClusterSumOfDistances[sampleIndex * nLabels]; + + LabelT sampleCluster = labels[sampleIndex]; + + int sampleClusterIndex = (int)sampleCluster; + + if (binCountArray[sampleClusterIndex] - 1 <= 0) { + d_aArray[sampleIndex] = -1; + return; + + } + + else { + d_aArray[sampleIndex] = (sampleToClusterSumOfDistancesVector[sampleClusterIndex]) / + (binCountArray[sampleClusterIndex] - 1); + + // modifying the sampleDistanceVector to give sample average distance + sampleToClusterSumOfDistancesVector[sampleClusterIndex] = MAX_VAL; + } +} + +/** + * @brief function to calculate the bincounts of number of samples in every label + * @tparam DataT: type of the data samples + * @tparam LabelT: type of the labels + * @param labels: the pointer to the array containing labels for every data sample (1 x nRows) + * @param binCountArray: pointer to the 1D array that contains the count of samples per cluster (1 x + * nLabels) + * @param nRows: number of data samples + * @param nUniqueLabels: number of Labels + * @param workspace: device buffer containing workspace memory + * @param stream: the cuda stream where to launch this kernel + */ +template +void countLabels(const LabelT* labels, + DataT* binCountArray, + int nRows, + int nUniqueLabels, + rmm::device_uvector& workspace, + cudaStream_t stream) +{ + int num_levels = nUniqueLabels + 1; + LabelT lower_level = 0; + LabelT upper_level = nUniqueLabels; + size_t temp_storage_bytes = 0; + + rmm::device_uvector countArray(nUniqueLabels, stream); + + RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(nullptr, + temp_storage_bytes, + labels, + binCountArray, + num_levels, + lower_level, + upper_level, + nRows, + stream)); + + workspace.resize(temp_storage_bytes, stream); + + RAFT_CUDA_TRY(cub::DeviceHistogram::HistogramEven(workspace.data(), + temp_storage_bytes, + labels, + binCountArray, + num_levels, + lower_level, + upper_level, + nRows, + stream)); +} + +/** + * @brief structure that defines the division Lambda for elementwise op + */ +template +struct DivOp { + HDI DataT operator()(DataT a, int b, int c) + { + if (b == 0) + return ULLONG_MAX; + else + return a / b; + } +}; + +/** + * @brief structure that defines the elementwise operation to calculate silhouette score using + * params 'a' and 'b' + */ +template +struct SilOp { + HDI DataT operator()(DataT a, DataT b) + { + if (a == 0 && b == 0 || a == b) + return 0; + else if (a == -1) + return 0; + else if (a > b) + return (b - a) / a; + else + return (b - a) / b; + } +}; + +/** + * @brief main function that returns the average silhouette score for a given set of data and its + * clusterings + * @tparam DataT: type of the data samples + * @tparam LabelT: type of the labels + * @param X_in: pointer to the input Data samples array (nRows x nCols) + * @param nRows: number of data samples + * @param nCols: number of features + * @param labels: the pointer to the array containing labels for every data sample (1 x nRows) + * @param nLabels: number of Labels + * @param silhouette_scorePerSample: pointer to the array that is optionally taken in as input and + * is populated with the silhouette score for every sample (1 x nRows) + * @param stream: the cuda stream where to launch this kernel + * @param metric: the numerical value that maps to the type of distance metric to be used in the + * calculations + */ +template +DataT silhouette_score( + raft::resources const& handle, + const DataT* X_in, + int nRows, + int nCols, + const LabelT* labels, + int nLabels, + DataT* silhouette_scorePerSample, + cudaStream_t stream, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded) +{ + ASSERT(nLabels >= 2 && nLabels <= (nRows - 1), + "silhouette Score not defined for the given number of labels!"); + + // compute the distance matrix + rmm::device_uvector distanceMatrix(nRows * nRows, stream); + rmm::device_uvector workspace(1, stream); + + auto X_in_view = raft::make_device_matrix_view(X_in, nRows, nCols); + + cuvs::distance::pairwise_distance( + handle, + X_in_view, + X_in_view, + raft::make_device_matrix_view(distanceMatrix.data(), nRows, nRows), + metric); + + // deciding on the array of silhouette scores for each dataPoint + rmm::device_uvector silhouette_scoreSamples(0, stream); + DataT* perSampleSilScore = nullptr; + if (silhouette_scorePerSample == nullptr) { + silhouette_scoreSamples.resize(nRows, stream); + perSampleSilScore = silhouette_scoreSamples.data(); + } else { + perSampleSilScore = silhouette_scorePerSample; + } + RAFT_CUDA_TRY(cudaMemsetAsync(perSampleSilScore, 0, nRows * sizeof(DataT), stream)); + + // getting the sample count per cluster + rmm::device_uvector binCountArray(nLabels, stream); + RAFT_CUDA_TRY(cudaMemsetAsync(binCountArray.data(), 0, nLabels * sizeof(DataT), stream)); + countLabels(labels, binCountArray.data(), nRows, nLabels, workspace, stream); + + // calculating the sample-cluster-distance-sum-array + rmm::device_uvector sampleToClusterSumOfDistances(nRows * nLabels, stream); + RAFT_CUDA_TRY(cudaMemsetAsync( + sampleToClusterSumOfDistances.data(), 0, nRows * nLabels * sizeof(DataT), stream)); + raft::linalg::reduce_cols_by_key(distanceMatrix.data(), + labels, + sampleToClusterSumOfDistances.data(), + nRows, + nRows, + nLabels, + stream); + + // creating the a array and b array + rmm::device_uvector d_aArray(nRows, stream); + rmm::device_uvector d_bArray(nRows, stream); + RAFT_CUDA_TRY(cudaMemsetAsync(d_aArray.data(), 0, nRows * sizeof(DataT), stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(d_bArray.data(), 0, nRows * sizeof(DataT), stream)); + + // kernel that populates the d_aArray + // kernel configuration + dim3 numThreadsPerBlock(32, 1, 1); + dim3 numBlocks(raft::ceildiv(nRows, numThreadsPerBlock.x), 1, 1); + + // calling the kernel + populateAKernel<<>>( + sampleToClusterSumOfDistances.data(), + binCountArray.data(), + d_aArray.data(), + labels, + nRows, + nLabels, + std::numeric_limits::max()); + + // elementwise dividing by bincounts + rmm::device_uvector averageDistanceBetweenSampleAndCluster(nRows * nLabels, stream); + RAFT_CUDA_TRY(cudaMemsetAsync( + averageDistanceBetweenSampleAndCluster.data(), 0, nRows * nLabels * sizeof(DataT), stream)); + + raft::linalg::matrixVectorOp(averageDistanceBetweenSampleAndCluster.data(), + sampleToClusterSumOfDistances.data(), + binCountArray.data(), + binCountArray.data(), + nLabels, + nRows, + true, + true, + DivOp(), + stream); + + // calculating row-wise minimum + raft::linalg::reduce( + d_bArray.data(), + averageDistanceBetweenSampleAndCluster.data(), + nLabels, + nRows, + std::numeric_limits::max(), + true, + true, + stream, + false, + raft::identity_op{}, + raft::min_op{}); + + // calculating the silhouette score per sample using the d_aArray and d_bArray + raft::linalg::binaryOp>( + perSampleSilScore, d_aArray.data(), d_bArray.data(), nRows, SilOp(), stream); + + // calculating the sum of all the silhouette score + rmm::device_scalar d_avgSilhouetteScore(stream); + RAFT_CUDA_TRY(cudaMemsetAsync(d_avgSilhouetteScore.data(), 0, sizeof(DataT), stream)); + + raft::linalg::mapThenSumReduce(d_avgSilhouetteScore.data(), + nRows, + raft::identity_op(), + stream, + perSampleSilScore, + perSampleSilScore); + + DataT avgSilhouetteScore = d_avgSilhouetteScore.value(stream); + + raft::resource::sync_stream(handle, stream); + + avgSilhouetteScore /= nRows; + + return avgSilhouetteScore; +} + +}; // namespace detail +}; // namespace stats +}; // namespace cuvs diff --git a/cpp/src/stats/detail/trustworthiness_score.cuh b/cpp/src/stats/detail/trustworthiness_score.cuh new file mode 100644 index 000000000..f4725a2e8 --- /dev/null +++ b/cpp/src/stats/detail/trustworthiness_score.cuh @@ -0,0 +1,220 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include + +#include +#include + +#define N_THREADS 512 + +namespace cuvs { +namespace stats { +namespace detail { + +/** + * @brief Build the lookup table + * @param[out] lookup_table: Lookup table giving nearest neighbor order + * of pairwise distance calculations given sample index + * @param[in] X_ind: Sorted indexes of pairwise distance calculations of X + * @param n: Number of samples + * @param work: Number of elements to consider + */ +RAFT_KERNEL build_lookup_table(int* lookup_table, const int* X_ind, int n, int work) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= work) return; + + int sample_idx = i / n; + int nn_idx = i % n; + + int idx = X_ind[i]; + lookup_table[(sample_idx * n) + idx] = nn_idx; +} + +/** + * @brief Compute a the rank of trustworthiness score + * @param[out] rank: Resulting rank + * @param[out] lookup_table: Lookup table giving nearest neighbor order + * of pairwise distance calculations given sample index + * @param[in] emb_ind: Indexes of KNN on embeddings + * @param n: Number of samples + * @param n_neighbors: Number of neighbors considered by trustworthiness score + * @param work: Batch to consider (to do it at once use n * n_neighbors) + */ +template +RAFT_KERNEL compute_rank(double* rank, + const int* lookup_table, + const knn_index_t* emb_ind, + int n, + int n_neighbors, + int work) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= work) return; + + int sample_idx = i / n_neighbors; + + knn_index_t emb_nn_ind = emb_ind[i]; + + int r = lookup_table[(sample_idx * n) + emb_nn_ind]; + int tmp = r - n_neighbors + 1; + if (tmp > 0) raft::myAtomicAdd(rank, tmp); +} + +/** + * @brief Compute a kNN and returns the indices of the nearest neighbors + * @param h Raft handle + * @param[in] input Input matrix containing the dataset + * @param n Number of samples + * @param d Number of features + * @param n_neighbors number of neighbors + * @param[out] indices KNN indexes + * @param[out] distances KNN distances + */ +template +void run_knn(const raft::resources& h, + math_t* input, + cuvs::distance::DistanceType metric, + int n, + int d, + int n_neighbors, + int64_t* indices, + math_t* distances) +{ + auto input_view = raft::make_device_matrix_view(input, n, d); + auto index = cuvs::neighbors::brute_force::build(h, input_view, metric); + + cuvs::neighbors::brute_force::search( + h, + index, + input_view, + raft::make_device_matrix_view(indices, n, n_neighbors), + raft::make_device_matrix_view(distances, n, n_neighbors), + std::nullopt); +} + +/** + * @brief Compute the trustworthiness score + * @param h Raft handle + * @param X[in]: Data in original dimension + * @param X_embedded[in]: Data in target dimension (embedding) + * @param n: Number of samples + * @param m: Number of features in high/original dimension + * @param d: Number of features in low/embedded dimension + * @param n_neighbors Number of neighbors considered by trustworthiness score + * @param batchSize Batch size + * @return Trustworthiness score + */ +template +double trustworthiness_score(const raft::resources& h, + const math_t* X, + math_t* X_embedded, + cuvs::distance::DistanceType metric, + int n, + int m, + int d, + int n_neighbors, + int batchSize = 512) +{ + cudaStream_t stream = raft::resource::get_cuda_stream(h); + + const int KNN_ALLOC = n * (n_neighbors + 1); + rmm::device_uvector emb_ind(KNN_ALLOC, stream); + rmm::device_uvector emb_dist(KNN_ALLOC, stream); + + run_knn(h, X_embedded, metric, n, d, n_neighbors + 1, emb_ind.data(), emb_dist.data()); + + const int PAIRWISE_ALLOC = batchSize * n; + rmm::device_uvector X_ind(PAIRWISE_ALLOC, stream); + rmm::device_uvector X_dist(PAIRWISE_ALLOC, stream); + rmm::device_uvector lookup_table(PAIRWISE_ALLOC, stream); + + double t = 0.0; + rmm::device_scalar t_dbuf(stream); + + int toDo = n; + while (toDo > 0) { + int curBatchSize = min(toDo, batchSize); + + // Takes at most batchSize vectors at a time + cuvs::distance::pairwise_distance( + h, + raft::make_device_matrix_view(&X[(n - toDo) * m], curBatchSize, m), + raft::make_device_matrix_view(X, n, m), + raft::make_device_matrix_view(X_dist.data(), curBatchSize, n), + metric); + + size_t colSortWorkspaceSize = 0; + bool bAllocWorkspace = false; + + raft::matrix::sort_cols_per_row(X_dist.data(), + X_ind.data(), + curBatchSize, + n, + bAllocWorkspace, + nullptr, + colSortWorkspaceSize, + stream); + + if (bAllocWorkspace) { + rmm::device_uvector sortColsWorkspace(colSortWorkspaceSize, stream); + + raft::matrix::sort_cols_per_row(X_dist.data(), + X_ind.data(), + curBatchSize, + n, + bAllocWorkspace, + sortColsWorkspace.data(), + colSortWorkspaceSize, + stream); + } + + int work = curBatchSize * n; + int n_blocks = raft::ceildiv(work, N_THREADS); + build_lookup_table<<>>( + lookup_table.data(), X_ind.data(), n, work); + + RAFT_CUDA_TRY(cudaMemsetAsync(t_dbuf.data(), 0, sizeof(double), stream)); + + work = curBatchSize * (n_neighbors + 1); + n_blocks = raft::ceildiv(work, N_THREADS); + compute_rank<<>>( + t_dbuf.data(), + lookup_table.data(), + &emb_ind.data()[(n - toDo) * (n_neighbors + 1)], + n, + n_neighbors + 1, + work); + RAFT_CUDA_TRY(cudaPeekAtLastError()); + + t += t_dbuf.value(stream); + + toDo -= curBatchSize; + } + + t = 1.0 - ((2.0 / ((n * n_neighbors) * ((2.0 * n) - (3.0 * n_neighbors) - 1.0))) * t); + + return t; +} + +} // namespace detail +} // namespace stats +} // namespace cuvs diff --git a/cpp/src/stats/silhouette_score.cu b/cpp/src/stats/silhouette_score.cu new file mode 100644 index 000000000..a3bd5f24d --- /dev/null +++ b/cpp/src/stats/silhouette_score.cu @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include "./detail/batched/silhouette_score.cuh" +#include "./detail/silhouette_score.cuh" + +namespace cuvs { +namespace stats { +namespace { +template +value_t _silhouette_score( + raft::resources const& handle, + raft::device_matrix_view X_in, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + idx_t n_unique_labels, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded) +{ + RAFT_EXPECTS(labels.extent(0) == X_in.extent(0), "Size mismatch between labels and data"); + + value_t* silhouette_score_per_sample_ptr = nullptr; + if (silhouette_score_per_sample.has_value()) { + silhouette_score_per_sample_ptr = silhouette_score_per_sample.value().data_handle(); + RAFT_EXPECTS(silhouette_score_per_sample.value().extent(0) == X_in.extent(0), + "Size mismatch between silhouette_score_per_sample and data"); + } + return detail::silhouette_score(handle, + X_in.data_handle(), + X_in.extent(0), + X_in.extent(1), + labels.data_handle(), + n_unique_labels, + silhouette_score_per_sample_ptr, + raft::resource::get_cuda_stream(handle), + metric); +} + +template +value_t _silhouette_score_batched( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + idx_t n_unique_labels, + idx_t batch_size, + cuvs::distance::DistanceType metric = cuvs::distance::DistanceType::L2Unexpanded) +{ + static_assert(std::is_integral_v, + "silhouette_score_batched: The index type " + "of each mdspan argument must be an integral type."); + static_assert(std::is_integral_v, + "silhouette_score_batched: The label type must be an integral type."); + RAFT_EXPECTS(labels.extent(0) == X.extent(0), "Size mismatch between labels and data"); + + value_t* scores_ptr = nullptr; + if (silhouette_score_per_sample.has_value()) { + scores_ptr = silhouette_score_per_sample.value().data_handle(); + RAFT_EXPECTS(silhouette_score_per_sample.value().extent(0) == X.extent(0), + "Size mismatch between silhouette_score_per_sample and data"); + } + return cuvs::stats::batched::detail::silhouette_score(handle, + X.data_handle(), + X.extent(0), + X.extent(1), + labels.data_handle(), + n_unique_labels, + scores_ptr, + batch_size, + metric); +} +} // namespace + +float silhouette_score( + raft::resources const& handle, + raft::device_matrix_view X_in, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + cuvs::distance::DistanceType metric) +{ + return _silhouette_score( + handle, X_in, labels, silhouette_score_per_sample, n_unique_labels, metric); +} + +double silhouette_score( + raft::resources const& handle, + raft::device_matrix_view X_in, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + cuvs::distance::DistanceType metric) +{ + return _silhouette_score( + handle, X_in, labels, silhouette_score_per_sample, n_unique_labels, metric); +} + +float silhouette_score_batched( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + int64_t batch_size, + cuvs::distance::DistanceType metric) +{ + return _silhouette_score_batched( + handle, X, labels, silhouette_score_per_sample, n_unique_labels, batch_size, metric); +} + +double silhouette_score_batched( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_vector_view labels, + std::optional> silhouette_score_per_sample, + int64_t n_unique_labels, + int64_t batch_size, + cuvs::distance::DistanceType metric) +{ + return _silhouette_score_batched( + handle, X, labels, silhouette_score_per_sample, n_unique_labels, batch_size, metric); +} +}; // namespace stats +}; // namespace cuvs diff --git a/cpp/src/stats/trustworthiness_score.cu b/cpp/src/stats/trustworthiness_score.cu new file mode 100644 index 000000000..9ab4cb706 --- /dev/null +++ b/cpp/src/stats/trustworthiness_score.cu @@ -0,0 +1,48 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include + +#include "./detail/trustworthiness_score.cuh" + +namespace cuvs { +namespace stats { + +double trustworthiness_score( + raft::resources const& handle, + raft::device_matrix_view X, + raft::device_matrix_view X_embedded, + int n_neighbors, + cuvs::distance::DistanceType metric, + int batch_size) +{ + RAFT_EXPECTS(X.extent(0) == X_embedded.extent(0), "Size mismatch between X and X_embedded"); + RAFT_EXPECTS(X.extent(0) <= std::numeric_limits::max(), "Index type not supported"); + + // TODO: Change the underlying implementation to remove the need to const_cast X_embedded. + return detail::trustworthiness_score(handle, + X.data_handle(), + const_cast(X_embedded.data_handle()), + metric, + static_cast(X.extent(0)), + static_cast(X.extent(1)), + static_cast(X_embedded.extent(1)), + n_neighbors, + batch_size); +} + +} // namespace stats +} // namespace cuvs diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index a30e2dec7..e04c39318 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -94,13 +94,13 @@ endfunction() if(BUILD_TESTS) ConfigureTest( - NAME NEIGHBORS_TEST PATH neighbors/brute_force.cu - neighbors/brute_force_prefiltered.cu neighbors/refine.cu GPUS 1 PERCENT 100 + NAME NEIGHBORS_TEST PATH neighbors/brute_force.cu neighbors/brute_force_prefiltered.cu + neighbors/refine.cu GPUS 1 PERCENT 100 ) ConfigureTest( - NAME CLUSTER_TEST PATH cluster/kmeans.cu cluster/kmeans_balanced.cu - cluster/kmeans_find_k.cu cluster/linkage.cu GPUS 1 PERCENT 100 + NAME CLUSTER_TEST PATH cluster/kmeans.cu cluster/kmeans_balanced.cu cluster/kmeans_find_k.cu + cluster/linkage.cu GPUS 1 PERCENT 100 ) ConfigureTest( @@ -130,15 +130,8 @@ if(BUILD_TESTS) ) ConfigureTest( - NAME - NEIGHBORS_ANN_BRUTE_FORCE_TEST - PATH - neighbors/ann_brute_force/test_float.cu - neighbors/ann_brute_force/test_half.cu - GPUS - 1 - PERCENT - 100 + NAME NEIGHBORS_ANN_BRUTE_FORCE_TEST PATH neighbors/ann_brute_force/test_float.cu + neighbors/ann_brute_force/test_half.cu GPUS 1 PERCENT 100 ) ConfigureTest( @@ -196,30 +189,33 @@ if(BUILD_TESTS) PERCENT 100 ) + ConfigureTest( + NAME STATS_TEST PATH stats/trustworthiness.cu stats/silhouette_score.cu GPUS 1 PERCENT 100 + ) endif() if(BUILD_C_TESTS) ConfigureTest(NAME INTEROP_TEST PATH core/interop.cu C_LIB) ConfigureTest( - NAME DISTANCE_C_TEST PATH distance/run_pairwise_distance_c.c - distance/pairwise_distance_c.cu C_LIB - ) - - ConfigureTest( - NAME BRUTEFORCE_C_TEST PATH neighbors/run_brute_force_c.c neighbors/brute_force_c.cu + NAME DISTANCE_C_TEST PATH distance/run_pairwise_distance_c.c distance/pairwise_distance_c.cu C_LIB ) ConfigureTest( - NAME IVF_FLAT_C_TEST PATH neighbors/run_ivf_flat_c.c neighbors/ann_ivf_flat_c.cu - C_LIB + NAME BRUTEFORCE_C_TEST PATH neighbors/run_brute_force_c.c neighbors/brute_force_c.cu C_LIB ) ConfigureTest( - NAME IVF_PQ_C_TEST PATH neighbors/run_ivf_pq_c.c neighbors/ann_ivf_pq_c.cu C_LIB + NAME IVF_FLAT_C_TEST PATH neighbors/run_ivf_flat_c.c neighbors/ann_ivf_flat_c.cu C_LIB ) + ConfigureTest(NAME IVF_PQ_C_TEST PATH neighbors/run_ivf_pq_c.c neighbors/ann_ivf_pq_c.cu C_LIB) + ConfigureTest(NAME CAGRA_C_TEST PATH neighbors/ann_cagra_c.cu C_LIB) + + if(BUILD_CAGRA_HNSWLIB) + ConfigureTest(NAME HNSW_C_TEST PATH neighbors/ann_hnsw_c.cu C_LIB) + endif() endif() # ################################################################################################## diff --git a/cpp/test/core/c_api.c b/cpp/test/core/c_api.c index a5b73d8fb..27973c2dd 100644 --- a/cpp/test/core/c_api.c +++ b/cpp/test/core/c_api.c @@ -31,6 +31,33 @@ int main() cuvsError_t stream_error = cuvsStreamSet(res, stream); if (stream_error == CUVS_ERROR) { exit(EXIT_FAILURE); } + // Allocate memory + void* ptr; + size_t bytes = 1024; + cuvsError_t alloc_error = cuvsRMMAlloc(res, &ptr, bytes); + if (alloc_error == CUVS_ERROR) { exit(EXIT_FAILURE); } + + // Free memory + cuvsError_t free_error = cuvsRMMFree(res, ptr, bytes); + if (free_error == CUVS_ERROR) { exit(EXIT_FAILURE); } + + // Enable pool memory resource + cuvsError_t pool_error = cuvsRMMPoolMemoryResourceEnable(10, 100); + if (pool_error == CUVS_ERROR) { exit(EXIT_FAILURE); } + + // Allocate memory again + void* ptr2; + cuvsError_t alloc_error_pool = cuvsRMMAlloc(res, &ptr2, 1024); + if (alloc_error_pool == CUVS_ERROR) { exit(EXIT_FAILURE); } + + // Free memory + cuvsError_t free_error_pool = cuvsRMMFree(res, ptr2, 1024); + if (free_error_pool == CUVS_ERROR) { exit(EXIT_FAILURE); } + + // Reset pool memory resource + cuvsError_t reset_error = cuvsRMMMemoryResourceReset(); + if (reset_error == CUVS_ERROR) { exit(EXIT_FAILURE); } + // Destroy resources cuvsError_t destroy_error = cuvsResourcesDestroy(res); if (destroy_error == CUVS_ERROR) { exit(EXIT_FAILURE); } diff --git a/cpp/test/neighbors/ann_hnsw_c.cu b/cpp/test/neighbors/ann_hnsw_c.cu new file mode 100644 index 000000000..fc740b924 --- /dev/null +++ b/cpp/test/neighbors/ann_hnsw_c.cu @@ -0,0 +1,131 @@ +/* + * Copyright (c) 2023-2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "../test_utils.cuh" +#include +#include +#include + +#include +#include +#include + +#include +#include +#include +#include + +float dataset[4][2] = {{0.74021935, 0.9209938}, + {0.03902049, 0.9689629}, + {0.92514056, 0.4463501}, + {0.6673192, 0.10993068}}; +float queries[4][2] = {{0.48216683, 0.0428398}, + {0.5084142, 0.6545497}, + {0.51260436, 0.2643005}, + {0.05198065, 0.5789965}}; + +std::vector neighbors_exp = {3, 0, 3, 1}; +std::vector distances_exp = {0.03878258, 0.12472608, 0.04776672, 0.15224178}; + +TEST(CagraHnswC, BuildSearch) +{ + // create cuvsResources_t + cuvsResources_t res; + cuvsResourcesCreate(&res); + + // create dataset DLTensor + DLManagedTensor dataset_tensor; + dataset_tensor.dl_tensor.data = dataset; + dataset_tensor.dl_tensor.device.device_type = kDLCPU; + dataset_tensor.dl_tensor.ndim = 2; + dataset_tensor.dl_tensor.dtype.code = kDLFloat; + dataset_tensor.dl_tensor.dtype.bits = 32; + dataset_tensor.dl_tensor.dtype.lanes = 1; + int64_t dataset_shape[2] = {4, 2}; + dataset_tensor.dl_tensor.shape = dataset_shape; + dataset_tensor.dl_tensor.strides = nullptr; + + // create index + cuvsCagraIndex_t index; + cuvsCagraIndexCreate(&index); + + // build index + cuvsCagraIndexParams_t build_params; + cuvsCagraIndexParamsCreate(&build_params); + cuvsCagraBuild(res, build_params, &dataset_tensor, index); + cuvsCagraSerializeToHnswlib(res, "/tmp/cagra_hnswlib.index", index); + + DLManagedTensor queries_tensor; + queries_tensor.dl_tensor.data = queries; + queries_tensor.dl_tensor.device.device_type = kDLCPU; + queries_tensor.dl_tensor.ndim = 2; + queries_tensor.dl_tensor.dtype.code = kDLFloat; + queries_tensor.dl_tensor.dtype.bits = 32; + queries_tensor.dl_tensor.dtype.lanes = 1; + int64_t queries_shape[2] = {4, 2}; + queries_tensor.dl_tensor.shape = queries_shape; + queries_tensor.dl_tensor.strides = nullptr; + + // create neighbors DLTensor + std::vector neighbors(4); + + DLManagedTensor neighbors_tensor; + neighbors_tensor.dl_tensor.data = neighbors.data(); + neighbors_tensor.dl_tensor.device.device_type = kDLCPU; + neighbors_tensor.dl_tensor.ndim = 2; + neighbors_tensor.dl_tensor.dtype.code = kDLUInt; + neighbors_tensor.dl_tensor.dtype.bits = 64; + neighbors_tensor.dl_tensor.dtype.lanes = 1; + int64_t neighbors_shape[2] = {4, 1}; + neighbors_tensor.dl_tensor.shape = neighbors_shape; + neighbors_tensor.dl_tensor.strides = nullptr; + + // create distances DLTensor + std::vector distances(4); + + DLManagedTensor distances_tensor; + distances_tensor.dl_tensor.data = distances.data(); + distances_tensor.dl_tensor.device.device_type = kDLCPU; + distances_tensor.dl_tensor.ndim = 2; + distances_tensor.dl_tensor.dtype.code = kDLFloat; + distances_tensor.dl_tensor.dtype.bits = 32; + distances_tensor.dl_tensor.dtype.lanes = 1; + int64_t distances_shape[2] = {4, 1}; + distances_tensor.dl_tensor.shape = distances_shape; + distances_tensor.dl_tensor.strides = nullptr; + + // create hnsw index + cuvsHnswIndex_t hnsw_index; + cuvsHnswIndexCreate(&hnsw_index); + hnsw_index->dtype = index->dtype; + cuvsHnswDeserialize(res, "/tmp/cagra_hnswlib.index", 2, L2Expanded, hnsw_index); + + // search index + cuvsHnswSearchParams_t search_params; + cuvsHnswSearchParamsCreate(&search_params); + cuvsHnswSearch( + res, search_params, hnsw_index, &queries_tensor, &neighbors_tensor, &distances_tensor); + + // verify output + ASSERT_TRUE(cuvs::hostVecMatch(neighbors_exp, neighbors, cuvs::Compare())); + ASSERT_TRUE(cuvs::hostVecMatch(distances_exp, distances, cuvs::CompareApprox(0.001f))); + + cuvsCagraIndexParamsDestroy(build_params); + cuvsCagraIndexDestroy(index); + cuvsHnswSearchParamsDestroy(search_params); + cuvsHnswIndexDestroy(hnsw_index); + cuvsResourcesDestroy(res); +} diff --git a/cpp/test/stats/silhouette_score.cu b/cpp/test/stats/silhouette_score.cu new file mode 100644 index 000000000..7662763fd --- /dev/null +++ b/cpp/test/stats/silhouette_score.cu @@ -0,0 +1,235 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include "../test_utils.cuh" + +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include + +namespace cuvs { +namespace stats { + +// parameter structure definition +struct silhouetteScoreParam { + int nRows; + int nCols; + int nLabels; + cuvs::distance::DistanceType metric; + int chunk; + double tolerance; +}; + +// test fixture class +template +class silhouetteScoreTest : public ::testing::TestWithParam { + protected: + silhouetteScoreTest() + : d_X(0, raft::resource::get_cuda_stream(handle)), + sampleSilScore(0, raft::resource::get_cuda_stream(handle)), + d_labels(0, raft::resource::get_cuda_stream(handle)) + { + } + + void host_silhouette_score() + { + // generating random value test input + std::vector h_X(nElements, 0.0); + std::vector h_labels(nRows, 0); + std::random_device rd; + std::default_random_engine dre(nElements * nLabels); + std::uniform_int_distribution intGenerator(0, nLabels - 1); + std::uniform_real_distribution realGenerator(0, 100); + + std::generate(h_X.begin(), h_X.end(), [&]() { return realGenerator(dre); }); + std::generate(h_labels.begin(), h_labels.end(), [&]() { return intGenerator(dre); }); + + // allocating and initializing memory to the GPU + auto stream = raft::resource::get_cuda_stream(handle); + d_X.resize(nElements, stream); + d_labels.resize(nElements, stream); + RAFT_CUDA_TRY(cudaMemsetAsync(d_X.data(), 0, d_X.size() * sizeof(DataT), stream)); + RAFT_CUDA_TRY(cudaMemsetAsync(d_labels.data(), 0, d_labels.size() * sizeof(LabelT), stream)); + sampleSilScore.resize(nElements, stream); + + raft::update_device(d_X.data(), &h_X[0], (int)nElements, stream); + raft::update_device(d_labels.data(), &h_labels[0], (int)nElements, stream); + + // finding the distance matrix + + rmm::device_uvector d_distanceMatrix(nRows * nRows, stream); + double* h_distanceMatrix = (double*)malloc(nRows * nRows * sizeof(double*)); + + auto d_X_view = raft::make_device_matrix_view(d_X.data(), nRows, nCols); + cuvs::distance::pairwise_distance( + handle, + d_X_view, + d_X_view, + raft::make_device_matrix_view(d_distanceMatrix.data(), nRows, nRows), + params.metric); + + raft::resource::sync_stream(handle, stream); + + raft::update_host(h_distanceMatrix, d_distanceMatrix.data(), nRows * nRows, stream); + + // finding the bincount array + + double* binCountArray = (double*)malloc(nLabels * sizeof(double*)); + memset(binCountArray, 0, nLabels * sizeof(double)); + + for (int i = 0; i < nRows; ++i) { + binCountArray[h_labels[i]] += 1; + } + + // finding the average intra cluster distance for every element + + double* a = (double*)malloc(nRows * sizeof(double*)); + + for (int i = 0; i < nRows; ++i) { + int myLabel = h_labels[i]; + double sumOfIntraClusterD = 0; + + for (int j = 0; j < nRows; ++j) { + if (h_labels[j] == myLabel) { sumOfIntraClusterD += h_distanceMatrix[i * nRows + j]; } + } + + if (binCountArray[myLabel] <= 1) + a[i] = -1; + else + a[i] = sumOfIntraClusterD / (binCountArray[myLabel] - 1); + } + + // finding the average inter cluster distance for every element + + double* b = (double*)malloc(nRows * sizeof(double*)); + + for (int i = 0; i < nRows; ++i) { + int myLabel = h_labels[i]; + double minAvgInterCD = ULLONG_MAX; + + for (int j = 0; j < nLabels; ++j) { + int curClLabel = j; + if (curClLabel == myLabel) continue; + double avgInterCD = 0; + + for (int k = 0; k < nRows; ++k) { + if (h_labels[k] == curClLabel) { avgInterCD += h_distanceMatrix[i * nRows + k]; } + } + + if (binCountArray[curClLabel]) + avgInterCD /= binCountArray[curClLabel]; + else + avgInterCD = ULLONG_MAX; + minAvgInterCD = min(minAvgInterCD, avgInterCD); + } + + b[i] = minAvgInterCD; + } + + // finding the silhouette score for every element + + double* truthSampleSilScore = (double*)malloc(nRows * sizeof(double*)); + for (int i = 0; i < nRows; ++i) { + if (a[i] == -1) + truthSampleSilScore[i] = 0; + else if (a[i] == 0 && b[i] == 0) + truthSampleSilScore[i] = 0; + else + truthSampleSilScore[i] = (b[i] - a[i]) / max(a[i], b[i]); + truthSilhouetteScore += truthSampleSilScore[i]; + } + + truthSilhouetteScore /= nRows; + } + + // the constructor + void SetUp() override + { + // getting the parameters + params = ::testing::TestWithParam::GetParam(); + + nRows = params.nRows; + nCols = params.nCols; + nLabels = params.nLabels; + chunk = params.chunk; + nElements = nRows * nCols; + + host_silhouette_score(); + + // calling the silhouette_score CUDA implementation + computedSilhouetteScore = cuvs::stats::silhouette_score( + handle, + raft::make_device_matrix_view(d_X.data(), nRows, nCols), + raft::make_device_vector_view(d_labels.data(), nRows), + std::make_optional(raft::make_device_vector_view(sampleSilScore.data(), nRows)), + nLabels, + params.metric); + + batchedSilhouetteScore = cuvs::stats::silhouette_score_batched( + handle, + raft::make_device_matrix_view(d_X.data(), nRows, nCols), + raft::make_device_vector_view(d_labels.data(), nRows), + std::make_optional(raft::make_device_vector_view(sampleSilScore.data(), nRows)), + nLabels, + chunk, + params.metric); + } + + // declaring the data values + raft::resources handle; + silhouetteScoreParam params; + int nLabels; + rmm::device_uvector d_X; + rmm::device_uvector sampleSilScore; + rmm::device_uvector d_labels; + int nRows; + int nCols; + int nElements; + double truthSilhouetteScore = 0; + double computedSilhouetteScore = 0; + double batchedSilhouetteScore = 0; + int chunk; +}; + +// setting test parameter values +const std::vector inputs = { + {4, 2, 3, cuvs::distance::DistanceType::L2Expanded, 4, 0.00001}, + {4, 2, 2, cuvs::distance::DistanceType::L2SqrtUnexpanded, 2, 0.00001}, + {8, 8, 3, cuvs::distance::DistanceType::L2Unexpanded, 4, 0.00001}, + {11, 2, 5, cuvs::distance::DistanceType::L2Expanded, 3, 0.00001}, + {40, 2, 8, cuvs::distance::DistanceType::L2Expanded, 10, 0.00001}, + {12, 7, 3, cuvs::distance::DistanceType::CosineExpanded, 8, 0.00001}, + {7, 5, 5, cuvs::distance::DistanceType::L1, 2, 0.00001}}; + +// writing the test suite +typedef silhouetteScoreTest silhouetteScoreTestClass; +TEST_P(silhouetteScoreTestClass, Result) +{ + ASSERT_NEAR(computedSilhouetteScore, truthSilhouetteScore, params.tolerance); + ASSERT_NEAR(batchedSilhouetteScore, truthSilhouetteScore, params.tolerance); +} +INSTANTIATE_TEST_CASE_P(silhouetteScore, silhouetteScoreTestClass, ::testing::ValuesIn(inputs)); + +} // end namespace stats +} // end namespace cuvs diff --git a/cpp/test/stats/trustworthiness.cu b/cpp/test/stats/trustworthiness.cu new file mode 100644 index 000000000..e2ed04a03 --- /dev/null +++ b/cpp/test/stats/trustworthiness.cu @@ -0,0 +1,356 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "../test_utils.cuh" + +#include +#include +#include +#include + +#include + +#include +#include + +namespace cuvs { +namespace stats { + +class TrustworthinessScoreTest : public ::testing::Test { + public: + TrustworthinessScoreTest() + : d_X(0, raft::resource::get_cuda_stream(handle)), + d_X_embedded(0, raft::resource::get_cuda_stream(handle)) + { + } + + protected: + void basicTest() + { + std::vector X = { + 5.6142087, 8.59787, -4.382763, -3.6452143, -5.8816037, -0.6330313, 4.6920023, + -0.79210913, 0.6106314, 2.1210914, 5.919943, -8.43784, -6.4819884, 0.41001374, + -6.1052523, -4.0825715, -5.314755, -2.834671, 5.751696, -6.5012555, -0.4719201, + -7.53353, 7.6789393, -1.4959852, -5.5977287, -9.564147, 1.2902534, 3.559834, + -6.7659483, 8.265964, 4.595404, 9.133477, -6.1553917, -6.319754, -2.9039452, + 4.4150834, -3.094395, -4.426273, 9.584571, -5.64133, 6.6209483, 7.4044604, + 3.9620576, 5.639907, 10.33007, -0.8792053, 5.143776, -7.464049, 1.2448754, + -5.6300974, 5.4518576, 4.119535, 6.749645, 7.627064, -7.2298336, 1.9681473, + -6.9083176, 6.404673, 0.07186685, 9.0994835, 8.51037, -8.986389, 0.40534487, + 2.115397, 4.086756, 1.2284287, -2.6272132, 0.06527536, -9.587425, -7.206078, + 7.864875, 7.4397306, -6.9233336, -2.6643622, 3.3466153, 7.0408177, -3.6069896, + -9.971769, 4.4075623, 7.9063697, 2.559074, 4.323717, 1.6867131, -1.1576937, + -9.893141, -3.251416, -7.4889135, -4.0588717, -2.73338, -7.4852257, 3.4460473, + 9.759119, -5.4680476, -4.722435, -8.032619, -1.4598992, 4.227361, 3.135568, + 1.1950601, 1.1982028, 6.998856, -6.131138, -6.6921015, 0.5361224, -7.1213965, + -5.6104236, -7.2212887, -2.2710054, 8.544764, -6.0254574, 1.4582269, -5.5587835, + 8.031556, -0.26328218, -5.2591386, -9.262641, 2.8691363, 5.299787, -9.209455, + 8.523085, 5.180329, 10.655528, -5.7171874, -6.7739563, -3.6306462, 4.067106, + -1.5912259, -3.2345476, 8.042973, -3.6364832, 4.1242137, 9.886953, 5.4743724, + 6.3058076, 9.369645, -0.5175337, 4.9859877, -7.879498, 1.358422, -4.147944, + 3.8984218, 5.894656, 6.4903927, 8.702036, -8.023722, 2.802145, -7.748032, + 5.8461113, -0.34215945, 11.298865, 1.4107164, -9.949621, -1.6257563, -10.655836, + 2.4528909, 1.1570255, 5.170669, 2.8398793, 7.1838694, 9.088459, 2.631155, + 3.964414, 2.8769252, 0.04198391, -0.16993195, 3.6747139, -2.8377378, 6.1782537, + 10.759618, -4.5642614, -8.522967, 0.8614642, 6.623416, -1.029324, 5.5488334, + -7.804511, 2.128833, 7.9042315, 7.789576, -2.7944536, 0.72271067, -10.511495, + -0.78634536, -10.661714, 2.9376361, 1.9148129, 6.22859, 0.26264945, 8.028384, + 6.8743043, 0.9351067, 7.0690722, 4.2846055, 1.4134506, -0.18144785, 5.2778087, + -1.7140163, 9.217541, 8.602799, -2.6537218, -7.8377395, 1.1244944, 5.4540544, + -0.38506773, 3.9885726, -10.76455, 1.4440702, 9.136163, 6.664117, -5.7046547, + 8.038592, -9.229767, -0.2799413, 3.6064725, 4.187257, 1.0516582, -2.0707326, + -0.7615968, -8.561018, -3.7831352, 10.300297, 5.332594, -6.5880876, -4.2508664, + 1.7985519, 5.7226253, -4.1223383, -9.6697855, 1.4885283, 7.524974, 1.7206005, + 4.890457, 3.7264557, 0.4428284, -9.922455, -4.250455, -6.4410596, -2.107994, + -1.4109765, -6.1325397, 0.32883006, 6.0489736, 7.7257385, -8.281174, 1.0129383, + -10.792166, 8.378851, 10.802716, 9.848448, -9.188757, 1.3151443, 1.9971865, + -2.521849, 4.3268294, -7.775683, -2.2902298, 3.0824065, -7.17559, 9.6100855, + 7.3965735, -10.476525, 5.895973, -3.6974669, -7.6688933, 1.7354839, -7.4045196, + -1.7992063, -4.0394845, 5.2471714, -2.250571, 2.528036, -8.343515, -2.2374575, + -10.019771, 0.73371273, 3.1853926, 2.7994921, 2.6637669, 7.620401, 7.515571, + 0.68636256, 5.834537, 4.650282, -1.0362619, 0.4461701, 3.7870514, -4.1340904, + 7.202998, 9.736904, -3.005512, -8.920467, 1.1228397, 6.2598724, 1.2812365, + 4.5442104, -8.791537, 0.92113096, 8.464749, 8.359035, -4.3923397, 1.2252625, + -10.1986475, -1.4409319, -10.013967, 3.9071581, 1.683064, 4.877419, 1.6570637, + 9.559105, 7.3546534, 0.36635467, 5.220211, 4.6303267, 0.6601065, 0.16149978, + 3.8818731, -3.4438233, 8.42085, 8.659159, -3.0935583, -8.039611, 2.3060374, + 5.134666, 1.0458113, 6.0190983, -9.143728, 0.99048865, 9.210842, 6.670241, + -5.9614363, 0.8747396, 7.078824, 8.067469, -10.314754, 0.45977542, -9.28306, + 9.1838665, 9.318644, 7.189082, -11.092555, 1.0320464, 3.882163, 0.10953151, + 7.9029684, -6.9068265, -1.3526366, 5.3996363, -8.430931, 11.452577, 6.39663, + -11.090514, 4.6662245, -3.1268113, -8.357452, 2.2276728, -10.357126, -0.9291848, + -3.4193344, 3.1289792, -2.5030103, 6.772719, 11.457757, -4.2125936, -6.684548, + -4.7611327, 3.6960156, -2.3030636, -3.0591488, 10.452471, -4.1267314, 5.66614, + 7.501461, 5.072407, 6.636537, 8.990381, -0.2559256, 4.737867, -6.2149944, + 2.535682, -5.5484023, 5.7113924, 3.4742818, 7.9915137, 7.0052586, -7.156467, + 1.4354781, -8.286235, 5.7523417, -2.4175215, 9.678009, 0.05066403, -9.645226, + -2.2658763, -9.518178, 4.493372, 2.3232365, 2.1659086, 0.42507997, 8.360246, + 8.23535, 2.6878164, 5.236947, 3.4924245, -0.6089895, 0.8884741, 4.359464, + -4.6073823, 7.83441, 8.958755, -3.4690795, -9.182282, 1.2478025, 5.6311107, + -1.2408862, 3.6316886, -8.684654, 2.1078515, 7.2813864, 7.9265943, -3.6135032, + 0.4571511, 8.493568, 10.496853, -7.432897, 0.8625995, -9.607528, 7.2899456, + 8.83158, 8.908199, -10.300263, 1.1451302, 3.7871468, -0.97040755, 5.7664757, + -8.9688, -2.146672, 5.9641485, -6.2908535, 10.126465, 6.1553903, -12.066902, + 6.301596, -5.0419583, -8.228695, 2.4879954, -8.918582, -3.7434099, -4.1593685, + 3.7431836, -1.1704745, 0.5524103, 9.109399, 9.571567, -11.209955, 1.2462777, + -9.554555, 9.091726, 11.477966, 7.630937, -10.450911, 1.9205878, 5.358983, + -0.44546837, 6.7611346, -9.74753, -0.5939732, 3.8892255, -6.437991, 10.294727, + 5.6723895, -10.7883, 6.192348, -5.293862, -10.811491, 1.0194173, -7.074576, + -3.192368, -2.5231771, 4.2791643, -0.53309685, 0.501366, 9.636625, 7.710316, + -6.4219728, 1.0975566, -8.218886, 6.9011984, 9.873679, 8.903804, -9.316832, + 1.2404599, 4.9039655, 1.2272617, 4.541515, -5.2753224, -3.2196746, 3.1303136, + -7.285681, 9.041425, 5.6417427, -9.93667, 5.7548947, -5.113397, -8.544622, + 4.182665, -7.7709813, -3.2810235, -3.312072, 3.8900535, -2.0604856, 6.709082, + -8.461194, 1.2666026, 4.8770437, 2.6955879, 3.0340345, -1.1614609, -3.536341, + -7.090382, -5.36146, 9.072544, 6.4554095, -4.4728956, -1.88395, 3.1095037, + 8.782348, -3.316743, -8.65248, 1.6802986, 8.186188, 2.1783829, 4.931278, + 4.158475, 1.4033595, -11.320101, -3.7084908, -6.740436, -2.5555193, -1.0451177, + -6.5569925, 0.82810307, 8.505919, 8.332857, -9.488569, -0.21588463, -8.056692, + 8.493993, 7.6401625, 8.812983, -9.377281, 2.4369764, 3.1766508, 0.6300803, + 5.6666765, -7.913654, -0.42301777, 4.506412, -7.8954244, 10.904591, 5.042256, + -9.626183, 8.347351, -3.605006, -7.923387, 1.1024277, -8.705793, -2.5151258, + -2.5066147, 4.0515003, -2.060757, 6.2635093, 8.286584, -6.0509276, -6.76452, + -3.1158175, 1.6578803, -1.4608748, -1.24211, 8.151246, -4.2970877, 6.093071, + 7.4911637, 4.51018, 4.8425875, 9.211085, -2.4386222, 4.5830803, -5.6079445, + 2.3713675, -4.0707507, 3.1787417, 5.462342, 6.915912, 6.3928423, -7.2970796, + 5.0112796, -9.140893, 4.9990606, 0.38391754, 7.7088532, 1.9340848, 8.18833, + 8.16617, -9.42086, -0.3388326, -9.659727, 8.243045, 8.099073, 8.439428, + -7.038694, 2.1077902, 3.3866816, -1.9975324, 7.4972878, -7.2525196, -1.553731, + 4.08758, -6.6922374, 9.50525, 4.026735, -9.243538, 7.2740564, -3.9319072, + -6.3228955, 1.6693478, -7.923119, -3.7423058, -2.2813146, 5.3469067, -1.8285407, + 3.3118162, 8.826356, -4.4641976, -6.4751124, -9.200089, -2.519147, 4.225298, + 2.4105988, -0.4344186, 0.53441775, 5.2836394, -8.2816105, -4.996147, -1.6870759, + -7.8543897, -3.9788852, -7.0346904, -3.1289773, 7.4567637, -5.6227813, 1.0709786, + -8.866012, 8.427324, -1.1755563, -5.789216, -8.197835, 5.3342214, 6.0646234, + -6.8975716, 7.717031, 3.480355, 8.312151, -3.6645212, -3.0976524, -8.090359, + -1.9176173, 2.4257212, 1.9700835, 0.4098958, 2.1341088, 7.652741, -9.9595585, + -5.989757, 0.10119354, -7.935407, -5.792786, -5.22783, -4.318978, 5.414037, + -6.4621663, 1.670883, -6.9224787, 8.696932, -2.0214002, -6.6681314, -8.326418, + 4.9049683, 5.4442496, -6.403739, 7.5822453, 7.0972915, -9.072851, -0.23897195, + 1.7662339, 5.3096304, 1.983179, -2.222645, -0.34700772, -9.094717, -6.107907, + 9.525174, 8.1550665, -5.6940084, -4.1636486, 1.7360662, 8.528821, -3.7299833, + -9.341266, 2.608542, 9.108706, 0.7978509, 4.2488184, 2.454484, 0.9446999, + -10.106636, -3.8973773, -6.6566644, -4.5647273, -0.99837756, -6.568582, 9.324853, + -7.9020953, 2.0910501, 2.2896829, 1.6790711, 1.3159255, -3.5258796, 1.8898442, + -8.105812, -4.924962, 8.771129, 7.1202874, -5.991957, -3.4106019, 2.4450088, + 7.796387, -3.055946, -7.8971434, 1.9856719, 9.001636, 1.8511922, 3.019749, + 3.1227696, 0.4822102, -10.021213, -3.530504, -6.225959, -3.0029628, -1.7881511, + -7.3879776, 1.3925704, 9.499782, -3.7318087, -3.7074296, -7.7466836, -1.5284524, + 4.0535855, 3.112011, 0.10340207, -0.5429599, 6.67026, -9.155924, -4.924038, + 0.64248866, -10.0103655, -3.2742946, -4.850029, -3.6707063, 8.586258, -5.855605, + 4.906918, -6.7813993, 7.9938135, -2.5473144, -5.688948, -7.822478, 2.1421318, + 4.66659, -9.701272, 9.549149, 0.8998125, -8.651497, -0.56899565, -8.639817, + 2.3088377, 2.1264515, 3.2764478, 2.341989, 8.594338, 8.630639, 2.8440373, + 6.2043204, 4.433932, 0.6320018, -1.8179281, 5.09452, -1.5741565, 8.153934, + 8.744339, -3.6945698, -8.883078, 1.5329908, 5.2745943, 0.44716078, 4.8809066, + -7.9594903, 1.134374, 9.233994, 6.5528665, -4.520542, 9.477355, -8.622195, + -0.23191702, 2.0485356, 3.9379985, 1.5916302, -1.4516805, -0.0843819, -7.8554378, + -5.88308, 7.999766, 6.2572145, -5.585321, -4.0097756, 0.42382592, 6.160884, + -3.631315, -8.333449, 2.770595, 7.8495173, 3.3331623, 4.940415, 3.6207345, + -0.037517, -11.034698, -3.185103, -6.614664, -3.2177854, -2.0792234, -6.8879867, + 7.821685, -8.455084, 1.0784642, 4.0033927, 2.7343264, 2.6052725, -4.1224284, + -0.89305353, -6.8267674, -4.9715133, 8.880253, 5.6994023, -5.9695024, -4.9181266, + 1.3017995, 7.972617, -3.9452884, -10.424556, 2.4504194, 6.21529, 0.93840516, + 4.2070026, 6.159839, 0.91979957, -8.706724, -4.317946, -6.6823545, -3.0388, + -2.464262, -7.3716645, 1.3926703, 6.544412, -5.6251183, -5.122411, -8.622049, + -2.3905911, 3.9138813, 1.9779967, -0.05011125, 0.13310997, 7.229751, -9.742043, + -8.08724, 1.2426697, -7.9230795, -3.3162494, -7.129571, -3.5488048, 7.4701195, + -5.2357526, 0.5917681, -6.272206, 6.342328, -2.909731, -4.991607, -8.845513, + 3.3228495, 7.033246, -7.8180246, 8.214469, 6.3910093, 9.185153, -6.20472, + -7.713809, -3.8481297, 3.5579286, 0.7078448, -3.2893546, 7.384514, -4.448121, + 3.0104196, 9.492943, 8.024847, 4.9114385, 9.965594, -3.014036, 5.182494, + -5.8806014, 2.5312455, -5.9926524, 4.474469, 6.3717875, 6.993105, 6.493093, + -8.935534, 3.004074, -8.055647, 8.315765, -1.3026813, 8.250377, 0.02606229, + 6.8508425, 9.655665, -7.0116496, -0.41060972, -10.049198, 7.897801, 6.7791023, + 8.3362, -9.821014, 2.491157, 3.5160472, -1.6228812, 7.398063, -8.769123, + -3.1743705, 3.2827861, -6.497855, 10.831924, 5.2761307, -9.704417, 4.3817043, + -3.9841619, -8.111647, 1.1883026, -8.115312, -2.9240117, -5.8879666, 4.20928, + -0.3587938, 6.935672, -10.177582, 0.48819053, 3.1250648, 2.9306343, 3.082544, + -3.477687, -1.3768549, -7.4922366, -3.756631, 10.039836, 3.6670392, -5.9761434, + -4.4728765, 3.244255, 7.027899, -2.3806512, -10.4100685, 1.605716, 7.7953773, + 0.5408159, 1.7156523, 3.824097, -1.0604783, -10.142124, -5.246805, -6.5283823, + -4.579547, -2.42714, -6.709197, 2.7782338, 7.33353, -6.454507, -2.9929368, + -7.8362985, -2.695445, 2.4900775, 1.6682367, 0.4641757, -1.0495365, 6.9631333, + -9.291356, -8.23837, -0.34263706, -8.275113, -2.8454232, -5.0864096, -2.681942, + 7.5450225, -6.2517986, 0.06810654, -6.470652, 4.9042645, -1.8369255, -6.6937943, + -7.9625087, 2.8510258, 6.180508, -8.282598, 7.919079, 1.4897474, 6.7217417, + -4.2459426, -4.114431, -8.375707, -2.143264, 5.6972933, 1.5574739, 0.39375135, + 1.7930849, 5.1737595, -7.826241, -5.160268, -0.80433255, -7.839536, -5.2620406, + -5.4643164, -3.185536, 6.620315, -7.065227, 1.0524757, -6.125088, 5.7126627, + -1.6161644, -3.852159, -9.164279, 2.7005782, 5.946544, -8.468236, 8.2145405, + 1.1035942, 6.590157, -4.0461283, -4.8090615, -7.6702685, -2.1121511, 5.1147075, + 1.6128504, 2.0064135, 1.0544407, 6.0038295, -7.8282537, -4.801278, 0.32349443, + -8.0649805, -4.372714, -5.61336, -5.21394, 8.176595, -5.4753284, 1.7800134, + -8.267283, 7.2133374, -0.16594432, -6.317046, -9.490406, 4.1261597, 5.473317, + -7.7551675, 7.007468, 7.478628, -8.801905, 0.10975724, 3.5478222, 4.797803, + 1.3825226, -3.357369, 0.99262005, -6.94877, -5.4781394, 9.632604, 5.7492557, + -5.9014316, -3.1632116, 2.340859, 8.708098, -3.1255999, -8.848661, 4.5612836, + 8.455157, 0.73460823, 4.112301, 4.392744, -0.30759293, -6.8036823, -3.0331545, + -8.269506, -2.82415, -0.9411246, -5.993506, 2.1618164, -8.716055, -0.7432543, + -10.255819, 3.095418, 2.5131428, 4.752442, 0.9907621, 7.8279433, 7.85814, + 0.50430876, 5.2840405, 4.457291, 0.03330028, -0.40692952, 3.9244103, -2.117118, + 7.6977615, 8.759009, -4.2157164, -9.136053, 3.247858, 4.668686, 0.76162136, + 5.3833632, -9.231471, 0.44309422, 8.380872, 6.7211227, -3.091507, 2.173508, + -9.038242, -1.3666698, -9.819077, 0.37825826, 2.3898845, 4.2440815, 1.9161536, + 7.24787, 6.9124637, 1.6238527, 5.1140285, 3.1935842, 1.02845, -1.1273454, + 5.638998, -2.497932, 8.342559, 8.586319, -2.9069402, -7.6387944, 3.5975037, + 4.4115705, 0.41506064, 4.9078383, -9.68327, 1.8159529, 9.744613, 8.40622, + -4.495336, 9.244892, -8.789869, 1.3158468, 4.018167, 3.3922846, 2.652022, + -2.7495477, 0.2528986, -8.268324, -6.004913, 10.428784, 6.6580734, -5.537176, + -1.7177434, 2.7504628, 6.7735, -2.4454272, -9.998361, 2.9483433, 6.8266654, + 2.3787718, 4.472637, 2.5871701, 0.7355365, -7.7027745, -4.1879907, -7.172832, + -4.1843605, -0.03646783, -5.419406, 6.958486, 11.011111, -7.1821184, -7.956423, + -3.408451, 4.6850276, -2.348787, -4.398289, 6.9787564, -3.8324208, 5.967827, + 8.433518, 4.660108, 5.5657144, 9.964243, -1.3515275, 6.404833, -6.4805903, + 2.4379845, -6.0816774, 1.752272, 5.3771873, 6.9613523, 6.9788294, -6.3894596, + 3.7521114, -6.8034263, 6.4458385, -0.7233525, 10.512529, 4.362273, 9.231461, + -6.3382263, -7.659, -3.461823, 4.71463, 0.17817476, -3.685746, 7.2962036, + -4.6489477, 5.218017, 11.546999, 4.7218375, 6.8498397, 9.281103, -3.900459, + 6.844054, -7.0886965, -0.05019227, -8.233724, 5.5808983, 6.374517, 8.321048, + 7.969449, -7.3478637, 1.4917561, -8.003144, 4.780668, -1.1981848, 7.753739, + 2.0260844, -8.880096, -3.4258451, -7.141975, 1.9637157, 1.814725, 5.311151, + 1.4831505, 7.8483663, 7.257948, 1.395786, 6.417756, 5.376912, 0.59505713, + 0.00062552, 3.6634305, -4.159713, 7.3571978, 10.966816, -2.5419605, -8.466229, + 1.904205, 5.6338267, -0.52567476, 5.59736, -8.361799, 0.5009981, 8.460681, + 7.3891273, -3.5272243, 5.0552278, 9.921456, -7.69693, -7.286378, -1.9198836, + 3.1666567, -2.5832257, -2.2445817, 9.888111, -5.076563, 5.677401, 7.497946, + 5.662994, 5.414262, 8.566503, -2.5530663, 7.1032815, -6.0612082, 1.3419591, + -4.9595256, 4.3377542, 4.3790717, 6.793512, 8.383502, -7.1278043, 3.3240774, + -9.379446, 6.838661, -0.81241214, 8.694813, 0.79141915, 7.632467, 8.575382, + -8.533798, 0.28954387, -7.5675836, 5.8653326, 8.97235, 7.1649346, -10.575289, + 0.9359381, 5.02381, -0.5609511, 5.543464, -7.69131, -2.1792977, 2.4729247, + -6.1917787, 10.373678, 7.6549597, -8.809486, 5.5657206, -3.3169382, -8.042887, + 2.0874746, -7.079005, -3.33398, -3.6843317, 4.0172358, -2.0754814, 1.1726758, + 7.4618697, 6.9483604, -8.469206, 0.7401797, -10.318176, 8.384557, 10.5476265, + 9.146971, -9.250223, 0.6290606, 4.4941425, -0.7514017, 7.2271705, -8.309598, + -1.4761636, 4.0140634, -6.021102, 9.132852, 5.6610966, -11.249811, 8.359293, + -1.9445792, -7.7393436, -0.3931331, -8.824441, -2.5995944, -2.5714035, 4.140213, + -3.6863053, 5.517265, 9.020411, -4.9286127, -7.871219, -3.7446704, 2.5179656, + -1.4543481, -2.2703636, 7.010597, -3.6436229, 6.753862, 7.4129915, 7.1406755, + 5.653706, 9.5445175, 0.15698843, 4.761813, -7.698002, 1.6870106, -4.5410123, + 4.171763, 5.3747005, 6.341021, 7.456738, -8.231657, 2.763487, -9.208167, + 6.676799, -1.1957736, 10.062605, 4.0975976, 7.312957, -2.4981596, -2.9658387, + -8.150425, -2.1075552, 2.64375, 1.6636052, 1.1483809, 0.09276015, 5.8556347, + -7.8481026, -5.9913163, -0.02840613, -9.937289, -1.0486673, -5.2340155, -3.83912, + 7.7165728, -8.409944, 0.80863273, -6.9119215, 7.5712357, 0.36031485, -6.056131, + -8.470033, 1.8678337, 3.0121377, -7.3096333, 8.205484, 5.262654, 8.774514, + -4.7603083, -7.2096143, -4.437014, 3.6080024, -1.624254, -4.2787876, 8.880863, + -4.8984556, 5.1782074, 9.944454, 3.911282, 3.5396595, 8.867042, -1.2006199, + 5.393288, -5.6455317, 0.7829499, -4.0338907, 2.479272, 6.5080743, 8.582535, + 7.0097537, -6.9823785, 3.984318, -7.225381, 5.3135114, -1.0391048, 8.951443, + -0.70119005, -8.510742, -0.42949116, -10.9224825, 2.8176029, 1.6800792, 5.778404, + 1.7269998, 7.1975236, 7.7258267, 2.7632928, 5.3399253, 3.4650044, 0.01971426, + -1.6468811, 4.114996, -1.5110453, 6.8689218, 8.269899, -3.1568048, -7.0344677, + 1.2911975, 5.950357, 0.19028673, 4.657226, -8.199647, 2.246055, 8.989509, + 5.3101015, -4.2400866}; + + std::vector X_embedded = { + -0.41849962, -0.53906363, 0.46958843, -0.35832694, -0.23779503, -0.29751351, -0.01072748, + -0.21353109, -0.54769957, -0.55086273, 0.37093949, -0.12714292, -0.06639574, -0.36098689, + -0.13060696, -0.07362658, -1.01205945, -0.39285606, 0.2864089, -0.32031146, -0.19595343, + 0.08900568, -0.04813879, -0.06563424, -0.42655188, -0.69014251, 0.51459783, -0.1942696, + -0.07767916, -0.6119386, 0.04813685, -0.22557008, -0.56890118, -0.60293794, 0.43429622, + -0.09240723, -0.00624062, -0.25800395, -0.1886092, 0.01655941, -0.01961523, -0.14147359, + 0.41414487, -0.8512944, -0.61199242, -0.18586016, 0.14024924, -0.41635606, -0.02890144, + 0.1065347, 0.39700791, -1.14060664, -0.95313865, 0.14416681, 0.17306046, -0.53189689, + -0.98987544, -0.67918193, 0.41787854, -0.20878236, -0.06612862, 0.03502904, -0.03765266, + -0.0980606, -0.00971657, 0.29432917, 0.36575687, -1.1645509, -0.89094597, 0.03718805, + 0.2310573, -0.38345811, -0.10401925, -0.10653082, 0.38469055, -0.88302094, -0.80197543, + 0.03548668, 0.02775662, -0.54374295, 0.03379983, 0.00923623, 0.29320273, -1.05263519, + -0.93360096, 0.03778313, 0.12360487, -0.56437284, 0.0644429, 0.33432651, 0.36450726, + -1.22978747, -0.83822101, -0.18796451, 0.34888434, -0.3801491, -0.45327303, -0.59747899, + 0.39697698, -0.15616602, -0.06159166, -0.40301991, -0.11725303, -0.11913263, -0.12406619, + -0.11227967, 0.43083835, -0.90535849, -0.81646025, 0.10012121, -0.0141237, -0.63747931, + 0.04805023, 0.34190539, 0.50725192, -1.17861414, -0.74641538, -0.09333111, 0.27992678, + -0.56214809, 0.04970971, 0.36249384, 0.57705611, -1.16913795, -0.69849908, 0.10957897, + 0.27983218, -0.62088525, 0.0410459, 0.23973398, 0.40960434, -1.14183664, -0.83321381, + 0.02149482, 0.21720445, -0.49869928, -0.95655465, -0.51680422, 0.45761383, -0.08351214, + -0.12151554, 0.00819737, -0.20813803, -0.01055793, 0.25319234, 0.36154974, 0.1822421, + -1.15837133, -0.92209691, -0.0501582, 0.08535917, -0.54003763, -1.08675635, -1.04009593, + 0.09408128, 0.07009826, -0.01762833, -0.19180447, -0.18029785, -0.20342001, 0.04034991, + 0.1814747, 0.36906669, -1.13532007, -0.8852452, 0.0782818, 0.16825101, -0.50301319, + -0.29128098, -0.65341312, 0.51484352, -0.38758236, -0.22531103, -0.55021971, 0.10804344, + -0.3521522, -0.38849035, -0.74110794, 0.53761131, -0.25142813, -0.1118066, -0.47453368, + 0.06347904, -0.23796193, -1.02682328, -0.47594091, 0.39515916, -0.2782529, -0.16566519, + 0.08063579, 0.00810116, -0.06213913, -1.059654, -0.62496334, 0.53698546, -0.11806234, + 0.00356161, 0.11513405, -0.14213292, 0.04102662, -0.36622161, -0.73686272, 0.48323864, + -0.27338892, -0.14203401, -0.41736352, 0.03332564, -0.21907479, -0.06396769, 0.01831361, + 0.46263444, -1.01878166, -0.86486858, 0.17622118, -0.01249686, -0.74530888, -0.9354887, + -0.5027945, 0.38170099, -0.15547098, 0.00677824, -0.04677663, -0.13541745, 0.07253501, + -0.97933143, -0.58001202, 0.48235369, -0.18836913, -0.02430783, 0.07572441, -0.08101331, + 0.00630076, -0.16881248, -0.67989182, 0.46083611, -0.43910736, -0.29321918, -0.38735861, + 0.07669903, -0.29749861, -0.40047669, -0.56722462, 0.33168188, -0.13118173, -0.06672747, + -0.56856316, -0.26269144, -0.14236671, 0.10651901, 0.4962585, 0.38848072, -1.06653547, + -0.64079332, -0.47378591, 0.43195483, -0.04856951, -0.9840439, -0.70610428, 0.34028092, + -0.2089237, -0.05382041, 0.01625874, -0.02080803, -0.12535211, -0.04146428, -1.24533033, + 0.48944879, 0.0578458, 0.26708388, -0.90321028, 0.35377088, -0.36791429, -0.35382384, + -0.52748734, 0.42854419, -0.31744713, -0.19174226, -0.39073724, -0.03258846, -0.19978228, + -0.36185205, -0.57412046, 0.43681973, -0.25414538, -0.12904905, -0.46334973, -0.03123853, + -0.11303604, -0.87073672, -0.45441297, 0.41825858, -0.25303507, -0.21845073, 0.10248682, + -0.11045569, -0.10002795, -0.00572806, 0.16519061, 0.42651513, -1.11417019, -0.83789682, + 0.02995787, 0.16843079, -0.53874511, 0.03056994, 0.17877036, 0.49632853, -1.03276777, + -0.74778616, -0.03971953, 0.10907949, -0.67385727, -0.9523471, -0.56550741, 0.40409449, + -0.2703723, -0.10175014, 0.13605487, -0.06306008, -0.01768126, -0.4749442, -0.56964815, + 0.39389887, -0.19248079, -0.04161081, -0.38728487, -0.20341556, -0.12656988, -0.35949609, + -0.46137866, 0.28798422, -0.06603147, -0.04363992, -0.60343552, -0.23565227, -0.10242701, + -0.06792886, 0.09689897, 0.33259571, -0.98854214, -0.84444433, 0.00673901, 0.13457057, + -0.43145794, -0.51500046, -0.50821936, 0.38000089, 0.0132636, 0.0580942, -0.40157595, + -0.11967677, 0.02549113, -0.10350953, 0.22918226, 0.40411913, -1.05619383, -0.71218503, + -0.02197581, 0.26422262, -0.34765676, 0.06601537, 0.21712676, 0.34723559, -1.20982027, + -0.95646334, 0.00793948, 0.27620381, -0.43475035, -0.67326003, -0.6137197, 0.43724492, + -0.17666136, -0.06591748, -0.18937394, -0.07400128, -0.06881691, -0.5201112, -0.61088628, + 0.4225319, -0.18969463, -0.06921366, -0.33993208, -0.06990873, -0.10288513, -0.70659858, + -0.56003648, 0.46628812, -0.16090363, -0.0185108, -0.1431348, -0.1128775, -0.0078648, + -0.02323332, 0.04292452, 0.39291084, -0.94897962, -0.63863206, -0.16546988, 0.23698957, + -0.30633628}; + + auto stream = raft::resource::get_cuda_stream(handle); + + d_X.resize(X.size(), stream); + d_X_embedded.resize(X_embedded.size(), stream); + raft::update_device(d_X.data(), X.data(), X.size(), stream); + raft::update_device(d_X_embedded.data(), X_embedded.data(), X_embedded.size(), stream); + auto n_sample = 50; + auto n_features_origin = 30; + auto n_features_embedded = 8; + + // euclidean test + score = cuvs::stats::trustworthiness_score( + handle, + raft::make_device_matrix_view(d_X.data(), n_sample, n_features_origin), + raft::make_device_matrix_view( + d_X_embedded.data(), n_sample, n_features_embedded), + 5, + cuvs::distance::DistanceType::L2SqrtUnexpanded); + } + + void SetUp() override { basicTest(); } + + void TearDown() override {} + + protected: + raft::resources handle; + + rmm::device_uvector d_X; + rmm::device_uvector d_X_embedded; + + double score; +}; + +typedef TrustworthinessScoreTest TrustworthinessScoreTestF; +TEST_F(TrustworthinessScoreTestF, Result) { ASSERT_TRUE(0.9375 < score && score < 0.9379); } +}; // namespace stats +}; // namespace cuvs diff --git a/dependencies.yaml b/dependencies.yaml index 98c9043aa..a53aef0f0 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -21,6 +21,19 @@ files: - test_py_cuvs - cupy - rust + bench_ann: + output: conda + matrix: + cuda: ["11.8", "12.5"] + arch: [x86_64, aarch64] + includes: + - rapids_build + - build_py_cuvs + - cuda + - cuda_version + - develop + - bench + - bench_python test_cpp: output: none includes: @@ -78,6 +91,7 @@ files: extras: table: project includes: + - cuda_wheels - run_py_cuvs py_test_py_cuvs: output: pyproject @@ -89,6 +103,20 @@ files: - test_python_common - test_py_cuvs - cupy + py_build_cuvs_bench: + output: pyproject + pyproject_dir: python/cuvs_bench + extras: + table: build-system + includes: + - rapids_build_setuptools + py_run_cuvs_bench: + output: pyproject + pyproject_dir: python/cuvs_bench + extras: + table: project + includes: + - bench_python channels: - rapidsai - rapidsai-nightly @@ -153,6 +181,14 @@ dependencies: - matrix: {cuda: "11.2", arch: aarch64} packages: [nvcc_linux-aarch64=11.2] + rapids_build_setuptools: + common: + - output_types: [requirements, pyproject] + packages: + - &rapids_build_backend rapids-build-backend>=0.3.0,<0.4.0.dev0 + - setuptools + - wheel + build_py_cuvs: common: - output_types: [conda] @@ -310,6 +346,36 @@ dependencies: - *libcusolver114 - *libcusparse_dev114 - *libcusparse114 + cuda_wheels: + specific: + - output_types: pyproject + matrices: + - matrix: + cuda: "12.*" + use_cuda_wheels: "true" + packages: + - nvidia-cublas-cu12 + - nvidia-curand-cu12 + - nvidia-cusolver-cu12 + - nvidia-cusparse-cu12 + # CUDA 11 does not provide wheels, so use the system libraries instead + - matrix: + cuda: "11.*" + use_cuda_wheels: "true" + packages: + # if use_cuda_wheels=false is provided, do not add dependencies on any CUDA wheels + # (e.g. for DLFW and pip devcontainers) + - matrix: + use_cuda_wheels: "false" + packages: + # if no matching matrix selectors passed, list the unsuffixed packages + # (just as a source of documentation, as this populates pyproject.toml in source control) + - matrix: + packages: + - nvidia-cublas + - nvidia-curand + - nvidia-cusolver + - nvidia-cusparse cupy: common: @@ -365,10 +431,6 @@ dependencies: specific: - output_types: conda matrices: - - matrix: - py: "3.9" - packages: - - python=3.9 - matrix: py: "3.10" packages: @@ -377,14 +439,18 @@ dependencies: py: "3.11" packages: - python=3.11 + - matrix: + py: "3.12" + packages: + - python=3.12 - matrix: packages: - - python>=3.9,<3.12 + - python>=3.10,<3.13 run_py_cuvs: common: - output_types: [conda, pyproject] packages: - - &numpy numpy>=1.23,<2.0a0 + - &numpy numpy>=1.23,<3.0a0 - output_types: [conda] packages: - *rmm_unsuffixed @@ -429,3 +495,23 @@ dependencies: - output_types: [conda, requirements, pyproject] packages: - scikit-learn + bench: + common: + - output_types: [conda, pyproject, requirements] + packages: + - hnswlib=0.6.2 + - nlohmann_json>=3.11.2 + - glog>=0.6.0 + - h5py>=3.8.0 + - benchmark>=1.8.2 + - openblas + - *rmm_unsuffixed + bench_python: + common: + - output_types: [conda] + packages: + - matplotlib + - pandas + - pyyaml + - pandas + - click diff --git a/docs/source/c_api/neighbors_cagra_c.rst b/docs/source/c_api/neighbors_cagra_c.rst index eb40d5578..a5ffc45b9 100644 --- a/docs/source/c_api/neighbors_cagra_c.rst +++ b/docs/source/c_api/neighbors_cagra_c.rst @@ -50,4 +50,10 @@ Index search :members: :content-only: +Index serialize +------------ +.. doxygengroup:: cagra_c_index_serialize + :project: cuvs + :members: + :content-only: diff --git a/docs/source/c_api/neighbors_hnsw_c.rst b/docs/source/c_api/neighbors_hnsw_c.rst new file mode 100644 index 000000000..4d83cd3e3 --- /dev/null +++ b/docs/source/c_api/neighbors_hnsw_c.rst @@ -0,0 +1,43 @@ +HNSW +==== + +This is a wrapper for hnswlib, to load a CAGRA index as an immutable HNSW index. The loaded HNSW index is only compatible in cuVS, and can be searched using wrapper functions. + + +.. role:: py(code) + :language: c + :class: highlight + +``#include `` + +Index search parameters +----------------------- + +.. doxygengroup:: hnsw_c_search_params + :project: cuvs + :members: + :content-only: + +Index +----- + +.. doxygengroup:: hnsw_c_index + :project: cuvs + :members: + :content-only: + +Index search +------------ + +.. doxygengroup:: cagra_c_index_search + :project: cuvs + :members: + :content-only: + +Index serialize +------------ + +.. doxygengroup:: hnsw_c_index_serialize + :project: cuvs + :members: + :content-only: diff --git a/docs/source/cpp_api.rst b/docs/source/cpp_api.rst index 5b8b4841a..49732dc92 100644 --- a/docs/source/cpp_api.rst +++ b/docs/source/cpp_api.rst @@ -11,3 +11,4 @@ C++ API Documentation cpp_api/distance.rst cpp_api/neighbors.rst cpp_api/selection.rst + cpp_api/stats.rst diff --git a/docs/source/cpp_api/neighbors_cagra.rst b/docs/source/cpp_api/neighbors_cagra.rst index 0e07406d1..d9f503871 100644 --- a/docs/source/cpp_api/neighbors_cagra.rst +++ b/docs/source/cpp_api/neighbors_cagra.rst @@ -28,7 +28,7 @@ Index search parameters :content-only: Index extend parameters ----------------------- +----------------------- .. doxygengroup:: cagra_cpp_extend_params :project: cuvs @@ -36,7 +36,7 @@ Index extend parameters :content-only: Index extend memory buffers ----------------------- +--------------------------- .. doxygengroup:: cagra_cpp_extend_memory_buffers :project: cuvs @@ -68,10 +68,17 @@ Index search :content-only: Index extend ------------ +------------ .. doxygengroup:: cagra_cpp_index_extend :project: cuvs :members: :content-only: +Index serialize +--------------- + +.. doxygengroup:: cagra_cpp_serialize + :project: cuvs + :members: + :content-only: diff --git a/docs/source/cpp_api/neighbors_hnsw.rst b/docs/source/cpp_api/neighbors_hnsw.rst new file mode 100644 index 000000000..b0af88af0 --- /dev/null +++ b/docs/source/cpp_api/neighbors_hnsw.rst @@ -0,0 +1,52 @@ +HNSW +==== + +This is a wrapper for hnswlib, to load a CAGRA index as an immutable HNSW index. The loaded HNSW index is only compatible in cuVS, and can be searched using wrapper functions. + +.. role:: py(code) + :language: c++ + :class: highlight + +``#include `` + +namespace *cuvs::neighbors::hnsw* + +Index search parameters +----------------------- + +.. doxygengroup:: hnsw_cpp_search_params + :project: cuvs + :members: + :content-only: + +Index +----- + +.. doxygengroup:: hnsw_cpp_index + :project: cuvs + :members: + :content-only: + +Index load +------------ + +.. doxygengroup:: hnsw_cpp_index_search + :project: cuvs + :members: + :content-only: + +Index search +------------ + +.. doxygengroup:: hnsw_cpp_index_search + :project: cuvs + :members: + :content-only: + +Index deserialize +--------------- + +.. doxygengroup:: hnsw_cpp_index_deserialize + :project: cuvs + :members: + :content-only: diff --git a/docs/source/cpp_api/stats.rst b/docs/source/cpp_api/stats.rst new file mode 100644 index 000000000..80d6c65fc --- /dev/null +++ b/docs/source/cpp_api/stats.rst @@ -0,0 +1,35 @@ +Stats +===== + + +This page provides C++ class references for the publicly-exposed elements of the `cuvs/stats` +package. + +.. role:: py(code) + :language: c++ + :class: highlight + +Silhouette Score +---------------- + +``#include `` + +namespace *cuvs::stats* + +.. doxygengroup:: stats_silhouette_score + :project: cuvs + :members: + :content-only: + +Trustworthiness Score +--------------------- + +``#include `` + +namespace *cuvs::stats* + +.. doxygengroup:: stats_trustworthiness + :project: cuvs + :members: + :content-only: + diff --git a/examples/c/CMakeLists.txt b/examples/c/CMakeLists.txt index fb508728f..d47cd4f1c 100644 --- a/examples/c/CMakeLists.txt +++ b/examples/c/CMakeLists.txt @@ -35,3 +35,7 @@ include(../cmake/thirdparty/get_cuvs.cmake) add_executable(CAGRA_C_EXAMPLE src/cagra_c_example.c) target_include_directories(CAGRA_C_EXAMPLE PUBLIC "$") target_link_libraries(CAGRA_C_EXAMPLE PRIVATE cuvs::c_api $) + +add_executable(L2_C_EXAMPLE src/L2_c_example.c) +target_include_directories(L2_C_EXAMPLE PUBLIC "$") +target_link_libraries(L2_C_EXAMPLE PRIVATE cuvs::c_api $) diff --git a/examples/c/src/L2_c_example.c b/examples/c/src/L2_c_example.c new file mode 100644 index 000000000..73ddf6103 --- /dev/null +++ b/examples/c/src/L2_c_example.c @@ -0,0 +1,123 @@ +/* + * Copyright (c) 2024, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include + +#include +#include +#include + +#define DIM 4 +#define N_ROWS 1 + +float PointA[N_ROWS][DIM] = {1.0,2.0,3.0,4.0}; +float PointB[N_ROWS][DIM] = {2.0,3.0,4.0,5.0}; + +cuvsResources_t res; + +void outputVector(float * Vec) { + printf("Vector is "); + for (int i = 0; i < DIM; ++i){ + printf(" %f",Vec[i]); + } + printf("\n"); +} + +/** + * @brief Initialize Tensor. + * + * @param[in] x_d Pointer to a vector + * @param[in] x_shape[] Two-dimensional array, which stores the number of rows and columns of vectors. + * @param[out] x_tensor Stores the initialized DLManagedTensor. + */ +void tensor_initialize(float* x_d, int64_t x_shape[2], DLManagedTensor* x_tensor) { + x_tensor->dl_tensor.data = x_d; + x_tensor->dl_tensor.device.device_type = kDLCUDA; + x_tensor->dl_tensor.ndim = 2; + x_tensor->dl_tensor.dtype.code = kDLFloat; + x_tensor->dl_tensor.dtype.bits = 32; + x_tensor->dl_tensor.dtype.lanes = 1; + x_tensor->dl_tensor.shape = x_shape; + x_tensor->dl_tensor.strides = NULL; +} + +/** + * @brief Calculate the euclidean distance between two arrays. + * + * @param[in] n_cols array length，also the dimension of the vector + * @param[in] x[] Pointer to a vector + * @param[in] y[] Pointer to another vector + * @param[out] ret will store the result about the euclidean distance + */ +void l2_distance_calc(int64_t n_cols,float x[], float y[], float *ret) { + float *x_d, *y_d; + float *distance_d; + cuvsRMMAlloc(res, (void**) &x_d, sizeof(float) * N_ROWS * n_cols); + cuvsRMMAlloc(res, (void**) &y_d, sizeof(float) * N_ROWS * n_cols); + cuvsRMMAlloc(res, (void**) &distance_d, sizeof(float) * N_ROWS * N_ROWS); + + // Use DLPack to represent x[] and y[] as tensors + cudaMemcpy(x_d, x, sizeof(float) * N_ROWS * n_cols, cudaMemcpyDefault); + cudaMemcpy(y_d, y, sizeof(float) * N_ROWS * n_cols, cudaMemcpyDefault); + + DLManagedTensor x_tensor; + int64_t x_shape[2] = {N_ROWS, n_cols}; + tensor_initialize(x_d, x_shape, &x_tensor); + + DLManagedTensor y_tensor; + int64_t y_shape[2] = {N_ROWS, n_cols}; + tensor_initialize(y_d, y_shape, &y_tensor); + + DLManagedTensor dist_tensor; + int64_t distances_shape[2] = {N_ROWS, N_ROWS}; + tensor_initialize(distance_d, distances_shape, &dist_tensor); + + // metric_arg default value is 2.0,used for Minkowski distance + cuvsPairwiseDistance(res, &x_tensor, &y_tensor, &dist_tensor, L2SqrtUnexpanded, 2.0); + + cudaMemcpy(ret, distance_d, sizeof(float) * N_ROWS * N_ROWS, cudaMemcpyDefault); + + cuvsRMMFree(res, distance_d, sizeof(float) * N_ROWS * N_ROWS); + cuvsRMMFree(res, x_d, sizeof(float) * N_ROWS * n_cols); + cuvsRMMFree(res, y_d, sizeof(float) * N_ROWS * n_cols); + +} + +int euclidean_distance_calculation_example() { + // Create a cuvsResources_t object + cuvsResourcesCreate(&res); + + outputVector((float *)PointA); + outputVector((float *)PointB); + + float ret; + + l2_distance_calc(DIM, (float *)PointA, (float *)PointB, &ret); + printf("L2 distance is %f.\n", ret); + + cuvsResourcesDestroy(res); + + return 0; +} + +int main() { + euclidean_distance_calculation_example(); + return 0; +} diff --git a/notebooks/VectorSearch_QuestionRetrieval.ipynb b/notebooks/VectorSearch_QuestionRetrieval.ipynb index 4023a1821..21d59975b 100644 --- a/notebooks/VectorSearch_QuestionRetrieval.ipynb +++ b/notebooks/VectorSearch_QuestionRetrieval.ipynb @@ -344,7 +344,7 @@ "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", - "version": "3.10.13" + "version": "3.11.9" } }, "nbformat": 4, diff --git a/notebooks/ivf_flat_example.ipynb b/notebooks/ivf_flat_example.ipynb index 2d9c5fb58..ce3586683 100644 --- a/notebooks/ivf_flat_example.ipynb +++ b/notebooks/ivf_flat_example.ipynb @@ -520,6 +520,30 @@ "metadata": {}, "outputs": [], "source": [] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "23010fbc-8f5a-4403-a112-33f190a85498", + "metadata": {}, + "outputs": [], + "source": [] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "774848e8-fa45-4223-bd2a-e8585650531e", + "metadata": {}, + "outputs": [], + "source": [] + }, + { + "cell_type": "code", + "execution_count": null, + "id": "6309b8a7-f4eb-4976-a824-cd4499a0000d", + "metadata": {}, + "outputs": [], + "source": [] } ], "metadata": { @@ -538,7 +562,7 @@ "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", - "version": "3.10.13" + "version": "3.11.9" } }, "nbformat": 4, diff --git a/notebooks/tutorial_ivf_pq.ipynb b/notebooks/tutorial_ivf_pq.ipynb index cc0fe4142..9d59daea2 100644 --- a/notebooks/tutorial_ivf_pq.ipynb +++ b/notebooks/tutorial_ivf_pq.ipynb @@ -124,6 +124,7 @@ "outputs": [], "source": [ "DATASET_URL = \"http://ann-benchmarks.com/sift-128-euclidean.hdf5\"\n", + "DATASET_NAME = \"SIFT-128\"\n", "f = load_dataset(DATASET_URL)" ] }, @@ -206,7 +207,7 @@ "# This function takes a row-major either numpy or cupy (GPU) array.\n", "# Generally, it's a bit faster with GPU inputs, but the CPU version may come in handy\n", "# if the whole dataset cannot fit into GPU memory.\n", - "index = ivf_pq.build(index_params, dataset, handle=resources)\n", + "index = ivf_pq.build(index_params, dataset, resources=resources)\n", "# This function is asynchronous so we need to explicitly synchronize the GPU before we can measure the execution time\n", "resources.sync()\n", "index" @@ -262,7 +263,7 @@ "outputs": [], "source": [ "%%time\n", - "distances, neighbors = ivf_pq.search(search_params, index, queries, k, handle=resources)\n", + "distances, neighbors = ivf_pq.search(search_params, index, queries, k, resources=resources)\n", "# Sync the GPU to make sure we've got the timing right\n", "resources.sync()" ] @@ -303,8 +304,8 @@ "source": [ "%%time\n", "\n", - "candidates = ivf_pq.search(search_params, index, queries, k * 2, handle=resources)[1]\n", - "distances, neighbors = refine(dataset, queries, candidates, k, handle=resources)\n", + "candidates = ivf_pq.search(search_params, index, queries, k * 2, resources=resources)[1]\n", + "distances, neighbors = refine(dataset, queries, candidates, k, resources=resources)\n", "resources.sync()" ] }, @@ -349,7 +350,7 @@ "bench_avg = np.zeros_like(bench_k, dtype=np.float32)\n", "bench_std = np.zeros_like(bench_k, dtype=np.float32)\n", "for i, k in enumerate(bench_k):\n", - " r = %timeit -o ivf_pq.search(search_params, index, queries, k, handle=resources); resources.sync()\n", + " r = %timeit -o ivf_pq.search(search_params, index, queries, k, resources=resources); resources.sync()\n", " bench_avg[i] = (queries.shape[0] * r.loops / np.array(r.all_runs)).mean()\n", " bench_std[i] = (queries.shape[0] * r.loops / np.array(r.all_runs)).std()\n", "\n", @@ -387,9 +388,9 @@ "k = 100\n", "for i, n_probes in enumerate(bench_probes):\n", " sp = ivf_pq.SearchParams(n_probes=n_probes)\n", - " r = %timeit -o ivf_pq.search(sp, index, queries, k, handle=resources); resources.sync()\n", + " r = %timeit -o ivf_pq.search(sp, index, queries, k, resources=resources); resources.sync()\n", " bench_qps[i] = (queries.shape[0] * r.loops / np.array(r.all_runs)).mean()\n", - " bench_recall[i] = calc_recall(ivf_pq.search(sp, index, queries, k, handle=resources)[1], gt_neighbors)\n", + " bench_recall[i] = calc_recall(ivf_pq.search(sp, index, queries, k, resources=resources)[1], gt_neighbors)\n", " " ] }, @@ -492,9 +493,9 @@ "bench_names = ['32/32', '32/16', '32/8', '16/16', '16/8']\n", "\n", "for i, sp in enumerate(search_ps):\n", - " r = %timeit -o ivf_pq.search(sp, index, queries, k, handle=resources); resources.sync()\n", + " r = %timeit -o ivf_pq.search(sp, index, queries, k, resources=resources); resources.sync()\n", " bench_qps_s1[i] = (queries.shape[0] * r.loops / np.array(r.all_runs)).mean()\n", - " bench_recall_s1[i] = calc_recall(ivf_pq.search(sp, index, queries, k, handle=resources)[1], gt_neighbors)" + " bench_recall_s1[i] = calc_recall(ivf_pq.search(sp, index, queries, k, resources=resources)[1], gt_neighbors)" ] }, { @@ -505,7 +506,7 @@ "source": [ "fig, ax = plt.subplots(1, 1, figsize=plt.figaspect(1/2))\n", "fig.suptitle(\n", - " f'Effects of search parameters on QPS/recall trade-off ({DATASET_FILENAME})\\n' + \\\n", + " f'Effects of search parameters on QPS/recall trade-off ({DATASET_NAME})\\n' + \\\n", " f'k = {k}, n_probes = {n_probes}, pq_dim = {pq_dim}')\n", "ax.plot(bench_recall_s1, bench_qps_s1, 'o')\n", "ax.set_xlabel('recall')\n", @@ -553,8 +554,8 @@ "source": [ "def search_refine(ps, ratio):\n", " k_search = k * ratio\n", - " candidates = ivf_pq.search(ps, index, queries, k_search, handle=resources)[1]\n", - " return candidates if ratio == 1 else refine(dataset, queries, candidates, k, handle=resources)[1]\n", + " candidates = ivf_pq.search(ps, index, queries, k_search, resources=resources)[1]\n", + " return candidates if ratio == 1 else refine(dataset, queries, candidates, k, resources=resources)[1]\n", "\n", "ratios = [1, 2, 4]\n", "bench_qps_sr = np.zeros((len(ratios), len(search_ps)), dtype=np.float32)\n", @@ -575,7 +576,7 @@ "source": [ "fig, ax = plt.subplots(1, 1, figsize=plt.figaspect(1/2))\n", "fig.suptitle(\n", - " f'Effects of search parameters on QPS/recall trade-off ({DATASET_FILENAME})\\n' + \\\n", + " f'Effects of search parameters on QPS/recall trade-off ({DATASET_NAME})\\n' + \\\n", " f'k = {k}, n_probes = {n_probes}, pq_dim = {pq_dim}')\n", "labels = []\n", "for j, ratio in enumerate(ratios):\n", @@ -629,8 +630,8 @@ " n_probes=n_probes,\n", " internal_distance_dtype=internal_distance_dtype,\n", " lut_dtype=lut_dtype)\n", - " candidates = ivf_pq.search(ps, index, queries, k_search, handle=resources)[1]\n", - " return candidates if ratio == 1 else refine(dataset, queries, candidates, k, handle=resources)[1]\n", + " candidates = ivf_pq.search(ps, index, queries, k_search, resources=resources)[1]\n", + " return candidates if ratio == 1 else refine(dataset, queries, candidates, k, resources=resources)[1]\n", "\n", "search_configs = [\n", " lambda n_probes: search_refine(np.float16, np.float16, 1, n_probes),\n", @@ -703,12 +704,13 @@ "\n", "for i, n_lists in enumerate(n_list_variants):\n", " index_params = ivf_pq.IndexParams(n_lists=n_lists, metric=metric, pq_dim=pq_dim)\n", - " index = ivf_pq.build(index_params, dataset, handle=resources)\n", + " index = ivf_pq.build(index_params, dataset, resources=resources)\n", " for j, pl_ratio in enumerate(pl_ratio_variants):\n", " n_probes = max(1, n_lists // pl_ratio)\n", " r = %timeit -o search_fun(n_probes); resources.sync()\n", " bench_qps_nl[i, j] = (queries.shape[0] * r.loops / np.array(r.all_runs)).mean()\n", - " bench_recall_nl[i, j] = calc_recall(search_fun(n_probes), gt_neighbors)" + " bench_recall_nl[i, j] = calc_recall(search_fun(n_probes), gt_neighbors)\n", + " del index" ] }, { @@ -719,7 +721,7 @@ "source": [ "fig, ax = plt.subplots(1, 1, figsize=plt.figaspect(1/2))\n", "fig.suptitle(\n", - " f'Effects of n_list on QPS/recall trade-off ({DATASET_FILENAME})\\n' + \\\n", + " f'Effects of n_list on QPS/recall trade-off ({DATASET_NAME})\\n' + \\\n", " f'k = {k}, pq_dim = {pq_dim}, search = {search_label}')\n", "labels = []\n", "for i, n_lists in enumerate(n_list_variants):\n", @@ -875,7 +877,7 @@ "bench_recall_ip = np.zeros_like(bench_qps_ip, dtype=np.float32)\n", "\n", "for i, index_params in enumerate(build_configs.values()):\n", - " index = ivf_pq.build(index_params, dataset, handle=resources)\n", + " index = ivf_pq.build(index_params, dataset, resources=resources)\n", " for l, search_fun in enumerate(search_configs):\n", " for j, n_probes in enumerate(n_probes_variants):\n", " r = %timeit -o search_fun(n_probes); resources.sync()\n", @@ -891,7 +893,7 @@ "source": [ "fig, ax = plt.subplots(len(search_config_names), 1, figsize=(16, len(search_config_names)*8))\n", "fig.suptitle(\n", - " f'Effects of index parameters on QPS/recall trade-off ({DATASET_FILENAME})\\n' + \\\n", + " f'Effects of index parameters on QPS/recall trade-off ({DATASET_NAME})\\n' + \\\n", " f'k = {k}, n_lists = {n_lists}')\n", "\n", "for j, search_label in enumerate(search_config_names):\n", @@ -932,7 +934,7 @@ "name": "python", "nbconvert_exporter": "python", "pygments_lexer": "ipython3", - "version": "3.10.13" + "version": "3.11.9" }, "vscode": { "interpreter": { diff --git a/pyproject.toml b/pyproject.toml index 2982db2a2..fbf4cf41f 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,6 +1,6 @@ [tool.black] line-length = 79 -target-version = ["py39"] +target-version = ["py310"] include = '\.py?$' force-exclude = ''' /( diff --git a/python/cuvs/CMakeLists.txt b/python/cuvs/CMakeLists.txt index 4d24682cf..7d2f8dcf9 100644 --- a/python/cuvs/CMakeLists.txt +++ b/python/cuvs/CMakeLists.txt @@ -37,6 +37,7 @@ project( option(FIND_CUVS_CPP "Search for existing CUVS C++ installations before defaulting to local files" OFF ) +option(USE_CUDA_MATH_WHEELS "Use the CUDA math wheels instead of the system libraries" OFF) message( "CUVS_PY: Searching for existing cuVS C/C++ installations before defaulting to local files: ${FIND_CUVS_CPP}" @@ -62,6 +63,8 @@ else() endif() if(NOT cuvs_FOUND) + find_package(CUDAToolkit REQUIRED) + set(BUILD_TESTS OFF) set(BUILD_C_LIBRARY ON) @@ -70,8 +73,26 @@ if(NOT cuvs_FOUND) set(CUDA_STATIC_MATH_LIBRARIES ON) set(CUVS_USE_RAFT_STATIC ON) + if(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 12.0) + set(CUDA_STATIC_MATH_LIBRARIES OFF) + elseif(USE_CUDA_MATH_WHEELS) + message(FATAL_ERROR "Cannot use CUDA math wheels with CUDA < 12.0") + endif() + add_subdirectory(../../cpp cuvs-cpp EXCLUDE_FROM_ALL) + if(NOT CUDA_STATIC_MATH_LIBRARIES AND USE_CUDA_MATH_WHEELS) + set(rpaths + "$ORIGIN/../nvidia/cublas/lib" + "$ORIGIN/../nvidia/curand/lib" + "$ORIGIN/../nvidia/cusolver/lib" + "$ORIGIN/../nvidia/cusparse/lib" + "$ORIGIN/../nvidia/nvjitlink/lib" + ) + set_property(TARGET cuvs PROPERTY INSTALL_RPATH ${rpaths} APPEND) + set_property(TARGET cuvs_c PROPERTY INSTALL_RPATH ${rpaths} APPEND) + endif() + set(cython_lib_dir cuvs) install(TARGETS cuvs cuvs_c DESTINATION ${cython_lib_dir}) endif() diff --git a/python/cuvs/pyproject.toml b/python/cuvs/pyproject.toml index cc2aa6a7f..0eb98d601 100644 --- a/python/cuvs/pyproject.toml +++ b/python/cuvs/pyproject.toml @@ -29,17 +29,22 @@ authors = [ { name = "NVIDIA Corporation" }, ] license = { text = "Apache 2.0" } -requires-python = ">=3.9" +requires-python = ">=3.10" dependencies = [ "cuda-python", - "numpy>=1.23,<2.0a0", + "numpy>=1.23,<3.0a0", + "nvidia-cublas", + "nvidia-curand", + "nvidia-cusolver", + "nvidia-cusparse", "pylibraft==24.10.*,>=0.0.0a0", ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. classifiers = [ "Intended Audience :: Developers", "Programming Language :: Python", - "Programming Language :: Python :: 3.9", "Programming Language :: Python :: 3.10", + "Programming Language :: Python :: 3.11", + "Programming Language :: Python :: 3.12", ] [project.optional-dependencies] @@ -128,7 +133,7 @@ requires = [ ] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. build-backend = "scikit_build_core.build" dependencies-file = "../../dependencies.yaml" -matrix-entry = "cuda_suffixed=true" +matrix-entry = "cuda_suffixed=true;use_cuda_wheels=true" [tool.pytest.ini_options] filterwarnings = [ diff --git a/python/cuvs_bench/LICENSE b/python/cuvs_bench/LICENSE new file mode 120000 index 000000000..30cff7403 --- /dev/null +++ b/python/cuvs_bench/LICENSE @@ -0,0 +1 @@ +../../LICENSE \ No newline at end of file diff --git a/python/cuvs_bench/cuvs_bench/VERSION b/python/cuvs_bench/cuvs_bench/VERSION new file mode 120000 index 000000000..d62dc733e --- /dev/null +++ b/python/cuvs_bench/cuvs_bench/VERSION @@ -0,0 +1 @@ +../../../VERSION \ No newline at end of file diff --git a/python/cuvs_bench/pyproject.toml b/python/cuvs_bench/pyproject.toml new file mode 100644 index 000000000..7bb9e2f8d --- /dev/null +++ b/python/cuvs_bench/pyproject.toml @@ -0,0 +1,70 @@ +# Copyright (c) 2024, NVIDIA CORPORATION. + +[build-system] +build-backend = "rapids_build_backend.build" +requires = [ + "rapids-build-backend>=0.3.0,<0.4.0.dev0", + "setuptools", + "wheel", +] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. + +[project] +name = "cuvs-bench" +dynamic = ["version"] +description = "cuVS benchmarks" +authors = [ + { name = "NVIDIA Corporation" }, +] +license = { text = "Apache 2.0" } +requires-python = ">=3.10" +dependencies = [ +] # This list was generated by `rapids-dependency-file-generator`. To make changes, edit ../../dependencies.yaml and run `rapids-dependency-file-generator`. +classifiers = [ + "Intended Audience :: Developers", + "Topic :: Database", + "Topic :: Scientific/Engineering", + "License :: OSI Approved :: Apache Software License", + "Programming Language :: Python", + "Programming Language :: Python :: 3.10", + "Programming Language :: Python :: 3.11", + "Programming Language :: Python :: 3.12", +] + +[project.urls] +Homepage = "https://github.com/rapidsai/raft" + +[tool.setuptools.packages.find] +where = ["src"] + +[tool.setuptools.package-data] +"*" = ["*.*", "VERSION"] + +[tool.isort] +line_length = 79 +multi_line_output = 3 +include_trailing_comma = true +force_grid_wrap = 0 +combine_as_imports = true +order_by_type = true +skip = [ + "thirdparty", + ".eggs", + ".git", + ".hg", + ".mypy_cache", + ".tox", + ".venv", + "_build", + "buck-out", + "build", + "dist", +] + +[tool.setuptools.dynamic] +version = { file = "cuvs_bench/VERSION" } + +[tool.rapids-build-backend] +build-backend = "scikit_build_core.build" +requires = [] +dependencies-file = "../../dependencies.yaml" +matrix-entry = "cuda_suffixed=true"