diff --git a/.devcontainer/cuda11.8-conda/devcontainer.json b/.devcontainer/cuda11.8-conda/devcontainer.json index 281859a2d..5e203a384 100644 --- a/.devcontainer/cuda11.8-conda/devcontainer.json +++ b/.devcontainer/cuda11.8-conda/devcontainer.json @@ -5,12 +5,17 @@ "args": { "CUDA": "11.8", "PYTHON_PACKAGE_MANAGER": "conda", - "BASE": "rapidsai/devcontainers:24.04-cpp-llvm16-cuda11.8-mambaforge-ubuntu22.04" + "BASE": "rapidsai/devcontainers:24.06-cpp-cuda11.8-mambaforge-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda11.8-conda" + ], "hostRequirements": {"gpu": "optional"}, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.4": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} }, "overrideFeatureInstallOrder": [ "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" diff --git a/.devcontainer/cuda11.8-pip/devcontainer.json b/.devcontainer/cuda11.8-pip/devcontainer.json index 5f042ffd4..682b2d155 100644 --- a/.devcontainer/cuda11.8-pip/devcontainer.json +++ b/.devcontainer/cuda11.8-pip/devcontainer.json @@ -5,12 +5,17 @@ "args": { "CUDA": "11.8", "PYTHON_PACKAGE_MANAGER": "pip", - "BASE": "rapidsai/devcontainers:24.04-cpp-cuda11.8-ubuntu22.04" + "BASE": "rapidsai/devcontainers:24.06-cpp-cuda11.8-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda11.8-pip" + ], "hostRequirements": {"gpu": "optional"}, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.4": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} }, "overrideFeatureInstallOrder": [ "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" diff --git a/.devcontainer/cuda12.2-conda/devcontainer.json b/.devcontainer/cuda12.2-conda/devcontainer.json index ab04745b7..d3bd5b769 100644 --- a/.devcontainer/cuda12.2-conda/devcontainer.json +++ b/.devcontainer/cuda12.2-conda/devcontainer.json @@ -5,12 +5,17 @@ "args": { "CUDA": "12.2", "PYTHON_PACKAGE_MANAGER": "conda", - "BASE": "rapidsai/devcontainers:24.04-cpp-mambaforge-ubuntu22.04" + "BASE": "rapidsai/devcontainers:24.06-cpp-mambaforge-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda12.2-conda" + ], "hostRequirements": {"gpu": "optional"}, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.4": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} }, "overrideFeatureInstallOrder": [ "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" diff --git a/.devcontainer/cuda12.2-pip/devcontainer.json b/.devcontainer/cuda12.2-pip/devcontainer.json index d4f567086..a667a3fd0 100644 --- a/.devcontainer/cuda12.2-pip/devcontainer.json +++ b/.devcontainer/cuda12.2-pip/devcontainer.json @@ -5,12 +5,17 @@ "args": { "CUDA": "12.2", "PYTHON_PACKAGE_MANAGER": "pip", - "BASE": "rapidsai/devcontainers:24.04-cpp-cuda12.2-ubuntu22.04" + "BASE": "rapidsai/devcontainers:24.06-cpp-cuda12.2-ubuntu22.04" } }, + "runArgs": [ + "--rm", + "--name", + "${localEnv:USER}-rapids-${localWorkspaceFolderBasename}-24.06-cuda12.2-pip" + ], "hostRequirements": {"gpu": "optional"}, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.4": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:24.6": {} }, "overrideFeatureInstallOrder": [ "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index ba2bf3bc8..06cb83f7d 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.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -37,7 +37,7 @@ jobs: python-build: needs: [cpp-build] secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -46,7 +46,7 @@ jobs: upload-conda: needs: [cpp-build, python-build] secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-upload-packages.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-upload-packages.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -56,7 +56,7 @@ jobs: if: github.ref_type == 'branch' needs: python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.06 with: arch: "amd64" branch: ${{ inputs.branch }} @@ -68,7 +68,7 @@ jobs: sha: ${{ inputs.sha }} wheel-build-cuspatial: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -78,7 +78,7 @@ jobs: wheel-publish-cuspatial: needs: wheel-build-cuspatial secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -87,7 +87,7 @@ jobs: package-name: cuspatial wheel-build-cuproj: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} @@ -97,7 +97,7 @@ jobs: wheel-publish-cuproj: needs: wheel-build-cuproj secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-publish.yaml@branch-24.06 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 201691502..be33cfd04 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -25,40 +25,40 @@ jobs: - wheel-tests-cuproj - devcontainer secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/pr-builder.yaml@branch-24.06 checks: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/checks.yaml@branch-24.06 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.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-build.yaml@branch-24.06 with: build_type: pull-request conda-cpp-tests: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.06 with: build_type: pull-request conda-python-build: needs: conda-cpp-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-build.yaml@branch-24.06 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.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@branch-24.06 with: build_type: pull-request conda-notebook-tests: needs: conda-python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.06 with: build_type: pull-request node_type: "gpu-v100-latest-1" @@ -68,7 +68,7 @@ jobs: docs-build: needs: conda-python-build secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/custom-job.yaml@branch-24.06 with: build_type: pull-request node_type: "gpu-v100-latest-1" @@ -78,34 +78,34 @@ jobs: wheel-build-cuspatial: needs: checks secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.06 with: build_type: pull-request script: ci/build_wheel_cuspatial.sh wheel-tests-cuspatial: needs: wheel-build-cuspatial secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.06 with: build_type: pull-request script: ci/test_wheel_cuspatial.sh wheel-build-cuproj: needs: checks secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-build.yaml@branch-24.06 with: build_type: pull-request script: ci/build_wheel_cuproj.sh wheel-tests-cuproj: needs: [wheel-build-cuspatial, wheel-build-cuproj] secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.06 with: build_type: pull-request script: ci/test_wheel_cuproj.sh devcontainer: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/build-in-devcontainer.yaml@branch-24.06 with: arch: '["amd64"]' cuda: '["12.2"]' diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 2adba7b3b..00312c6a0 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -16,7 +16,7 @@ on: jobs: conda-cpp-tests: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-cpp-tests.yaml@branch-24.06 with: build_type: nightly branch: ${{ inputs.branch }} @@ -24,7 +24,7 @@ jobs: sha: ${{ inputs.sha }} conda-python-tests: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/conda-python-tests.yaml@branch-24.06 with: build_type: nightly branch: ${{ inputs.branch }} @@ -32,7 +32,7 @@ jobs: sha: ${{ inputs.sha }} wheel-tests-cuspatial: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.06 with: build_type: nightly branch: ${{ inputs.branch }} @@ -41,7 +41,7 @@ jobs: script: ci/test_wheel_cuspatial.sh wheel-tests-cuproj: secrets: inherit - uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.04 + uses: rapidsai/shared-workflows/.github/workflows/wheels-test.yaml@branch-24.06 with: build_type: nightly branch: ${{ inputs.branch }} diff --git a/CHANGELOG.md b/CHANGELOG.md index fa7e8b254..264cef755 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -1,3 +1,27 @@ +# cuspatial 24.06.00 (5 Jun 2024) + +## 🚨 Breaking Changes + +- Replace rmm::mr::device_memory_resource* with rmm::device_async_resource_ref ([#1373](https://github.com/rapidsai/cuspatial/pull/1373)) [@harrism](https://github.com/harrism) + +## 🐛 Bug Fixes + +- create conda ci test env in one step ([#1387](https://github.com/rapidsai/cuspatial/pull/1387)) [@msarahan](https://github.com/msarahan) + +## 🛠️ Improvements + +- Fix up imports for cudf changes ([#1383](https://github.com/rapidsai/cuspatial/pull/1383)) [@vyasr](https://github.com/vyasr) +- Fix building cuspatial with CCCL main ([#1382](https://github.com/rapidsai/cuspatial/pull/1382)) [@trxcllnt](https://github.com/trxcllnt) +- Fix quadtree spatial join OOMs on large numbers of input polygons ([#1381](https://github.com/rapidsai/cuspatial/pull/1381)) [@trxcllnt](https://github.com/trxcllnt) +- Enable warnings-as-errors for cuproj tests ([#1379](https://github.com/rapidsai/cuspatial/pull/1379)) [@mroeschke](https://github.com/mroeschke) +- Always use a static gtest and gbench ([#1377](https://github.com/rapidsai/cuspatial/pull/1377)) [@trxcllnt](https://github.com/trxcllnt) +- Migrate to `{{ stdlib("c") }}` ([#1376](https://github.com/rapidsai/cuspatial/pull/1376)) [@hcho3](https://github.com/hcho3) +- add --rm and --name to devcontainer run args ([#1375](https://github.com/rapidsai/cuspatial/pull/1375)) [@trxcllnt](https://github.com/trxcllnt) +- Replace rmm::mr::device_memory_resource* with rmm::device_async_resource_ref ([#1373](https://github.com/rapidsai/cuspatial/pull/1373)) [@harrism](https://github.com/harrism) +- Enable all tests for `arm` jobs ([#1365](https://github.com/rapidsai/cuspatial/pull/1365)) [@galipremsagar](https://github.com/galipremsagar) +- Enable pytest failures on warnings on FutureWarnings (Replace deprecated `geopandas.dataset` module) ([#1360](https://github.com/rapidsai/cuspatial/pull/1360)) [@mroeschke](https://github.com/mroeschke) +- Fix `JOIN_POINT_IN_POLYGON_LARGE_TEST_EXP` test ([#1346](https://github.com/rapidsai/cuspatial/pull/1346)) [@trxcllnt](https://github.com/trxcllnt) + # cuSpatial 24.04.00 (10 Apr 2024) ## 🐛 Bug Fixes diff --git a/README.md b/README.md index 5ceb75c38..97ba2859a 100644 --- a/README.md +++ b/README.md @@ -113,7 +113,7 @@ An example command from the Release Selector: docker run --gpus all --pull always --rm -it \ --shm-size=1g --ulimit memlock=-1 --ulimit stack=67108864 \ -p 8888:8888 -p 8787:8787 -p 8786:8786 \ - nvcr.io/nvidia/rapidsai/notebooks:24.04-cuda11.8-py3.10 + nvcr.io/nvidia/rapidsai/notebooks:24.06-cuda11.8-py3.10 ``` ### Install with Conda @@ -125,7 +125,7 @@ cuSpatial can be installed with conda (miniconda, or the full Anaconda distribut ```shell conda install -c rapidsai -c conda-forge -c nvidia \ - cuspatial=24.04 python=3.11 cudatoolkit=11.8 + cuspatial=24.06 python=3.11 cudatoolkit=11.8 ``` We also provide nightly Conda packages built from the HEAD of our latest development branch. @@ -184,7 +184,7 @@ To build and install cuSpatial from source please see the [build documentation]( If you find cuSpatial useful in your published work, please consider citing the repository. ```bibtex -@misc{cuspatial:24.04, +@misc{cuspatial:24.06, author = {{NVIDIA Corporation}}, title = {cuSpatial: GPU-Accelerated Geospatial and Spatiotemporal Algorithms}, year = {2023}, diff --git a/VERSION b/VERSION index 4a2fe8aa5..0bff6981a 100644 --- a/VERSION +++ b/VERSION @@ -1 +1 @@ -24.04.00 +24.06.00 diff --git a/ci/build_docs.sh b/ci/build_docs.sh index 2f0d97622..0d4e97488 100755 --- a/ci/build_docs.sh +++ b/ci/build_docs.sh @@ -6,30 +6,24 @@ set -euo pipefail rapids-logger "Create test conda environment" . /opt/conda/etc/profile.d/conda.sh +rapids-logger "Downloading artifacts from previous jobs" +CPP_CHANNEL="$(rapids-download-conda-from-s3 cpp)" +PYTHON_CHANNEL="$(rapids-download-conda-from-s3 python)" + rapids-dependency-file-generator \ --output conda \ --file_key docs \ - --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" | tee env.yaml + --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" \ + --prepend-channel "${CPP_CHANNEL}" --prepend-channel "${PYTHON_CHANNEL}" | tee env.yaml rapids-mamba-retry env create --yes -f env.yaml -n docs conda activate docs rapids-print-env -rapids-logger "Downloading artifacts from previous jobs" -CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) -PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python) - -rapids-mamba-retry install \ - --channel "${CPP_CHANNEL}" \ - --channel "${PYTHON_CHANNEL}" \ - libcuspatial \ - cuspatial \ - cuproj - export RAPIDS_VERSION="$(rapids-version)" export RAPIDS_VERSION_MAJOR_MINOR="$(rapids-version-major-minor)" -export RAPIDS_VERSION_NUMBER="24.04" +export RAPIDS_VERSION_NUMBER="24.06" export RAPIDS_DOCS_DIR="$(mktemp -d)" rapids-logger "Build cuSpatial CPP docs" diff --git a/ci/release/update-version.sh b/ci/release/update-version.sh index 8bf47d77f..ae13d7629 100755 --- a/ci/release/update-version.sh +++ b/ci/release/update-version.sh @@ -45,7 +45,10 @@ NEXT_SHORT_TAG_PEP440=$(python -c "from setuptools.extern import packaging; prin DEPENDENCIES=( cudf cuml + cuspatial libcudf + libcuspatial + libcuspatial-tests librmm rmm cuspatial @@ -81,4 +84,5 @@ sed_runner "s/notebooks:[0-9]\+\.[0-9]\+/notebooks:${NEXT_SHORT_TAG}/g" README.m find .devcontainer/ -type f -name devcontainer.json -print0 | while IFS= read -r -d '' filename; do sed_runner "s@rapidsai/devcontainers:[0-9.]*@rapidsai/devcontainers:${NEXT_SHORT_TAG}@g" "${filename}" sed_runner "s@rapidsai/devcontainers/features/rapids-build-utils:[0-9.]*@rapidsai/devcontainers/features/rapids-build-utils:${NEXT_SHORT_TAG_PEP440}@" "${filename}" + sed_runner "s@rapids-\${localWorkspaceFolderBasename}-[0-9.]*@rapids-\${localWorkspaceFolderBasename}-${NEXT_SHORT_TAG}@g" "${filename}" done diff --git a/ci/test_cpp.sh b/ci/test_cpp.sh index 35a981e93..0f0a09ef1 100755 --- a/ci/test_cpp.sh +++ b/ci/test_cpp.sh @@ -5,11 +5,16 @@ set -euo pipefail . /opt/conda/etc/profile.d/conda.sh + +rapids-logger "Downloading artifacts from previous jobs" +CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) + rapids-logger "Generate C++ testing dependencies" rapids-dependency-file-generator \ --output conda \ - --file_key test_cpp \ - --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch)" | tee env.yaml + --file-key test_cpp \ + --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch)" \ + --prepend-channel "${CPP_CHANNEL}" | tee env.yaml rapids-mamba-retry env create --yes -f env.yaml -n test @@ -18,7 +23,6 @@ set +u conda activate test set -u -CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"}/ mkdir -p "${RAPIDS_TESTS_DIR}" @@ -27,10 +31,6 @@ export CUSPATIAL_HOME="${PWD}" rapids-print-env -rapids-mamba-retry install \ - --channel "${CPP_CHANNEL}" \ - libcuspatial libcuspatial-tests - rapids-logger "Check GPU usage" nvidia-smi diff --git a/ci/test_notebooks.sh b/ci/test_notebooks.sh index a4ae9c21c..61ae4f194 100755 --- a/ci/test_notebooks.sh +++ b/ci/test_notebooks.sh @@ -5,11 +5,16 @@ set -euo pipefail . /opt/conda/etc/profile.d/conda.sh +rapids-logger "Downloading artifacts from previous jobs" +CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) +PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python) + rapids-logger "Generate notebook testing dependencies" rapids-dependency-file-generator \ --output conda \ - --file_key test_notebooks \ - --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" | tee env.yaml + --file-key test_notebooks \ + --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" \ + --prepend-channel "${CPP_CHANNEL}" --prepend-channel "${PYTHON_CHANNEL}" | tee env.yaml rapids-mamba-retry env create --yes -f env.yaml -n test @@ -20,15 +25,6 @@ set -u rapids-print-env -rapids-logger "Downloading artifacts from previous jobs" -CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) -PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python) - -rapids-mamba-retry install \ - --channel "${CPP_CHANNEL}" \ - --channel "${PYTHON_CHANNEL}" \ - cuspatial libcuspatial cuproj - NBTEST="$(realpath "$(dirname "$0")/utils/nbtest.sh")" pushd notebooks diff --git a/ci/test_python.sh b/ci/test_python.sh index c6b80ed45..9a7cada37 100755 --- a/ci/test_python.sh +++ b/ci/test_python.sh @@ -5,11 +5,16 @@ set -euo pipefail . /opt/conda/etc/profile.d/conda.sh +rapids-logger "Downloading artifacts from previous jobs" +CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) +PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python) + rapids-logger "Generate Python testing dependencies" rapids-dependency-file-generator \ --output conda \ - --file_key test_python \ - --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" | tee env.yaml + --file-key test_python \ + --matrix "cuda=${RAPIDS_CUDA_VERSION%.*};arch=$(arch);py=${RAPIDS_PY_VERSION}" \ + --prepend-channel "${CPP_CHANNEL}" --prepend-channel "${PYTHON_CHANNEL}" | tee env.yaml rapids-mamba-retry env create --yes -f env.yaml -n test @@ -18,10 +23,6 @@ set +u conda activate test set -u -rapids-logger "Downloading artifacts from previous jobs" -CPP_CHANNEL=$(rapids-download-conda-from-s3 cpp) -PYTHON_CHANNEL=$(rapids-download-conda-from-s3 python) - RAPIDS_TESTS_DIR=${RAPIDS_TESTS_DIR:-"${PWD}/test-results"} RAPIDS_COVERAGE_DIR=${RAPIDS_COVERAGE_DIR:-"${PWD}/coverage-results"} mkdir -p "${RAPIDS_TESTS_DIR}" "${RAPIDS_COVERAGE_DIR}" @@ -31,11 +32,6 @@ export CUSPATIAL_HOME="${PWD}" rapids-print-env -rapids-mamba-retry install \ - --channel "${CPP_CHANNEL}" \ - --channel "${PYTHON_CHANNEL}" \ - libcuspatial cuspatial cuproj - rapids-logger "Check GPU usage" nvidia-smi diff --git a/ci/test_wheel_cuproj.sh b/ci/test_wheel_cuproj.sh index e6888134d..9451190e0 100755 --- a/ci/test_wheel_cuproj.sh +++ b/ci/test_wheel_cuproj.sh @@ -19,16 +19,12 @@ python -m pip install --no-deps ./local-cuspatial-dep/cuspatial*.whl # echo to expand wildcard before adding `[extra]` requires for pip python -m pip install $(echo ./dist/cuproj*.whl)[test] -if [[ "$(arch)" == "aarch64" && ${RAPIDS_BUILD_TYPE} == "pull-request" ]]; then - python ./ci/wheel_smoke_test_cuproj.py -else - rapids-logger "pytest cuproj" - pushd python/cuproj/cuproj - python -m pytest \ - --cache-clear \ - --junitxml="${RAPIDS_TESTS_DIR}/junit-cuproj.xml" \ - --numprocesses=8 \ - --dist=worksteal \ - tests - popd -fi +rapids-logger "pytest cuproj" +pushd python/cuproj/cuproj +python -m pytest \ + --cache-clear \ + --junitxml="${RAPIDS_TESTS_DIR}/junit-cuproj.xml" \ + --numprocesses=8 \ + --dist=worksteal \ + tests +popd diff --git a/ci/test_wheel_cuspatial.sh b/ci/test_wheel_cuspatial.sh index ff4a2405f..aa36681bc 100755 --- a/ci/test_wheel_cuspatial.sh +++ b/ci/test_wheel_cuspatial.sh @@ -15,16 +15,12 @@ python -m pip install --no-binary fiona 'fiona>=1.8.19,<1.9' # echo to expand wildcard before adding `[extra]` requires for pip python -m pip install $(echo ./dist/cuspatial*.whl)[test] -if [[ "$(arch)" == "aarch64" && ${RAPIDS_BUILD_TYPE} == "pull-request" ]]; then - python ./ci/wheel_smoke_test_cuspatial.py -else - rapids-logger "pytest cuspatial" - pushd python/cuspatial/cuspatial - python -m pytest \ - --cache-clear \ - --junitxml="${RAPIDS_TESTS_DIR}/junit-cuspatial.xml" \ - --numprocesses=8 \ - --dist=worksteal \ - tests - popd -fi +rapids-logger "pytest cuspatial" +pushd python/cuspatial/cuspatial +python -m pytest \ + --cache-clear \ + --junitxml="${RAPIDS_TESTS_DIR}/junit-cuspatial.xml" \ + --numprocesses=8 \ + --dist=worksteal \ + tests +popd diff --git a/ci/wheel_smoke_test_cuproj.py b/ci/wheel_smoke_test_cuproj.py deleted file mode 100644 index d5b9b494a..000000000 --- a/ci/wheel_smoke_test_cuproj.py +++ /dev/null @@ -1,17 +0,0 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. - -from cuproj import Transformer as cuTransformer - -from cupy.testing import assert_allclose - -if __name__ == '__main__': - # Sydney opera house latitude and longitude - lat = -33.8587 - lon = 151.2140 - - # Transform to UTM using cuproj - cu_transformer = cuTransformer.from_crs("epsg:4326", "EPSG:32756") - cuproj_x, cuproj_y = cu_transformer.transform(lat, lon) - - assert_allclose(cuproj_x, 334783.9544807102) - assert_allclose(cuproj_y, 6252075.961741454) diff --git a/ci/wheel_smoke_test_cuspatial.py b/ci/wheel_smoke_test_cuspatial.py deleted file mode 100644 index df2af9abd..000000000 --- a/ci/wheel_smoke_test_cuspatial.py +++ /dev/null @@ -1,28 +0,0 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. - -import numpy as np -import cudf -import cuspatial -import pyarrow as pa -from shapely.geometry import Point - -if __name__ == '__main__': - order, quadtree = cuspatial.quadtree_on_points( - cuspatial.GeoSeries([Point(0.5, 0.5), Point(1.5, 1.5)]), - *(0, 2, 0, 2), # bbox - 1, # scale - 1, # max_depth - 1, # min_size - ) - cudf.testing.assert_frame_equal( - quadtree, - cudf.DataFrame( - { - "key": cudf.Series(pa.array([0, 3], type=pa.uint32())), - "level": cudf.Series(pa.array([0, 0], type=pa.uint8())), - "is_internal_node": cudf.Series(pa.array([False, False], type=pa.bool_())), - "length": cudf.Series(pa.array([1, 1], type=pa.uint32())), - "offset": cudf.Series(pa.array([0, 1], type=pa.uint32())), - } - ), - ) diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 506060cf9..ae933a67c 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -11,8 +11,8 @@ dependencies: - cmake>=3.26.4 - cuda-version=11.8 - cudatoolkit -- cudf==24.4.* -- cuml==24.4.* +- cudf==24.6.* +- cuml==24.6.* - cupy>=12.0.0 - curl - cxx-compiler @@ -20,12 +20,12 @@ dependencies: - doxygen - gcc_linux-64=11.* - geopandas>=0.11.0 -- gmock>=1.13.0 -- gtest>=1.13.0 - ipython - ipywidgets -- libcudf==24.4.* -- librmm==24.4.* +- libcudf==24.6.* +- libcuspatial-tests==24.6.* +- libcuspatial==24.6.* +- librmm==24.6.* - myst-parser - nbsphinx - ninja @@ -41,7 +41,7 @@ dependencies: - pytest-cov - pytest-xdist - python>=3.9,<3.12 -- rmm==24.4.* +- rmm==24.6.* - scikit-build-core>=0.7.0 - scikit-image - shapely diff --git a/conda/environments/all_cuda-122_arch-x86_64.yaml b/conda/environments/all_cuda-122_arch-x86_64.yaml index 566aa2327..4f48db253 100644 --- a/conda/environments/all_cuda-122_arch-x86_64.yaml +++ b/conda/environments/all_cuda-122_arch-x86_64.yaml @@ -14,8 +14,8 @@ dependencies: - cuda-nvcc - cuda-nvrtc-dev - cuda-version=12.2 -- cudf==24.4.* -- cuml==24.4.* +- cudf==24.6.* +- cuml==24.6.* - cupy>=12.0.0 - curl - cxx-compiler @@ -23,12 +23,12 @@ dependencies: - doxygen - gcc_linux-64=11.* - geopandas>=0.11.0 -- gmock>=1.13.0 -- gtest>=1.13.0 - ipython - ipywidgets -- libcudf==24.4.* -- librmm==24.4.* +- libcudf==24.6.* +- libcuspatial-tests==24.6.* +- libcuspatial==24.6.* +- librmm==24.6.* - myst-parser - nbsphinx - ninja @@ -43,7 +43,7 @@ dependencies: - pytest-cov - pytest-xdist - python>=3.9,<3.12 -- rmm==24.4.* +- rmm==24.6.* - scikit-build-core>=0.7.0 - scikit-image - shapely diff --git a/conda/recipes/cuproj/conda_build_config.yaml b/conda/recipes/cuproj/conda_build_config.yaml index 92f617ce1..67f968412 100644 --- a/conda/recipes/cuproj/conda_build_config.yaml +++ b/conda/recipes/cuproj/conda_build_config.yaml @@ -10,7 +10,10 @@ cuda_compiler: cuda11_compiler: - nvcc -sysroot_version: +c_stdlib: + - sysroot + +c_stdlib_version: - "2.17" cmake_version: diff --git a/conda/recipes/cuproj/meta.yaml b/conda/recipes/cuproj/meta.yaml index 22309116e..4721daefb 100644 --- a/conda/recipes/cuproj/meta.yaml +++ b/conda/recipes/cuproj/meta.yaml @@ -1,4 +1,4 @@ -# Copyright (c) 2018-2023, NVIDIA CORPORATION. +# Copyright (c) 2018-2024, NVIDIA CORPORATION. {% set version = environ['RAPIDS_PACKAGE_VERSION'].lstrip('v') %} {% set minor_version = version.split('.')[0] + '.' + version.split('.')[1] %} @@ -52,7 +52,7 @@ requirements: {% endif %} - cuda-version ={{ cuda_version }} - ninja - - sysroot_{{ target_platform }} {{ sysroot_version }} + - {{ stdlib("c") }} host: {% if cuda_major != "11" %} - cuda-cudart-dev diff --git a/conda/recipes/cuspatial/conda_build_config.yaml b/conda/recipes/cuspatial/conda_build_config.yaml index e28b98da7..e3ca633eb 100644 --- a/conda/recipes/cuspatial/conda_build_config.yaml +++ b/conda/recipes/cuspatial/conda_build_config.yaml @@ -10,7 +10,10 @@ cuda_compiler: cuda11_compiler: - nvcc -sysroot_version: +c_stdlib: + - sysroot + +c_stdlib_version: - "2.17" cmake_version: diff --git a/conda/recipes/cuspatial/meta.yaml b/conda/recipes/cuspatial/meta.yaml index c2c60133d..1bc68677f 100644 --- a/conda/recipes/cuspatial/meta.yaml +++ b/conda/recipes/cuspatial/meta.yaml @@ -52,7 +52,7 @@ requirements: {% endif %} - cuda-version ={{ cuda_version }} - ninja - - sysroot_{{ target_platform }} {{ sysroot_version }} + - {{ stdlib("c") }} host: {% if cuda_major != "11" %} - cuda-cudart-dev diff --git a/conda/recipes/libcuspatial/conda_build_config.yaml b/conda/recipes/libcuspatial/conda_build_config.yaml index 37d54ccc2..51cfd482f 100644 --- a/conda/recipes/libcuspatial/conda_build_config.yaml +++ b/conda/recipes/libcuspatial/conda_build_config.yaml @@ -16,5 +16,8 @@ cmake_version: gtest_version: - ">=1.13.0" -sysroot_version: +c_stdlib: + - sysroot + +c_stdlib_version: - "2.17" diff --git a/conda/recipes/libcuspatial/meta.yaml b/conda/recipes/libcuspatial/meta.yaml index 659aa511c..3e212690c 100644 --- a/conda/recipes/libcuspatial/meta.yaml +++ b/conda/recipes/libcuspatial/meta.yaml @@ -1,4 +1,4 @@ -# Copyright (c) 2018-2023, NVIDIA CORPORATION. +# Copyright (c) 2018-2024, NVIDIA CORPORATION. {% set version = environ['RAPIDS_PACKAGE_VERSION'].lstrip('v') %} {% set minor_version = version.split('.')[0] + '.' + version.split('.')[1] %} @@ -42,12 +42,10 @@ requirements: - cuda-version ={{ cuda_version }} - cmake {{ cmake_version }} - ninja - - sysroot_{{ target_platform }} {{ sysroot_version }} + - {{ stdlib("c") }} host: - cuda-version ={{ cuda_version }} - doxygen - - gmock {{ gtest_version }} - - gtest {{ gtest_version }} - libcudf ={{ minor_version }} - librmm ={{ minor_version }} - sqlite @@ -129,5 +127,3 @@ outputs: - cuda-cudart {% endif %} - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} - - gmock {{ gtest_version }} - - gtest {{ gtest_version }} diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 1de36d515..8908617b1 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -107,7 +107,8 @@ rapids_cpm_init() include(cmake/thirdparty/get_cudf.cmake) # find or install GoogleTest if (CUSPATIAL_BUILD_TESTS) - include(cmake/thirdparty/get_gtest.cmake) + include(${rapids-cmake-dir}/cpm/gtest.cmake) + rapids_cpm_gtest(BUILD_STATIC) endif() # find or add ranger include (cmake/thirdparty/get_ranger.cmake) @@ -222,7 +223,7 @@ endif() if(CUSPATIAL_BUILD_BENCHMARKS) # Find or install GoogleBench include(${rapids-cmake-dir}/cpm/gbench.cmake) - rapids_cpm_gbench() + rapids_cpm_gbench(BUILD_STATIC) # Find or install NVBench Temporarily force downloading of fmt because current versions of nvbench # do not support the latest version of fmt, which is automatically pulled into our conda diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 99780a677..10e626f3f 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -1,5 +1,5 @@ #============================================================================= -# Copyright (c) 2019-2021, NVIDIA CORPORATION. +# Copyright (c) 2019-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. @@ -23,6 +23,17 @@ add_library(cuspatial_benchmark_common OBJECT target_compile_features(cuspatial_benchmark_common PUBLIC cxx_std_17 cuda_std_17) +set_target_properties(cuspatial_benchmark_common + PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" + INSTALL_RPATH "\$ORIGIN/../../../lib" + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED ON + POSITION_INDEPENDENT_CODE ON + INTERFACE_POSITION_INDEPENDENT_CODE ON +) + target_link_libraries(cuspatial_benchmark_common PUBLIC benchmark::benchmark cudf::cudftestutil @@ -43,6 +54,10 @@ function(ConfigureBench CMAKE_BENCH_NAME) set_target_properties(${CMAKE_BENCH_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" INSTALL_RPATH "\$ORIGIN/../../../lib" + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED ON ) target_link_libraries(${CMAKE_BENCH_NAME} PRIVATE benchmark::benchmark_main cuspatial_benchmark_common) install( @@ -61,7 +76,11 @@ function(ConfigureNVBench CMAKE_BENCH_NAME) ${CMAKE_BENCH_NAME} PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" INSTALL_RPATH "\$ORIGIN/../../../lib" - ) + CXX_STANDARD 17 + CXX_STANDARD_REQUIRED ON + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED ON +) target_link_libraries( ${CMAKE_BENCH_NAME} PRIVATE cuspatial_benchmark_common nvbench::main ) diff --git a/cpp/cmake/thirdparty/get_gtest.cmake b/cpp/cmake/thirdparty/get_gtest.cmake deleted file mode 100644 index 45ef26f0a..000000000 --- a/cpp/cmake/thirdparty/get_gtest.cmake +++ /dev/null @@ -1,39 +0,0 @@ -# ============================================================================= -# Copyright (c) 2023, 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. -# ============================================================================= - -# This function finds gtest and sets any additional necessary environment variables. -function(find_and_configure_gtest) - include(${rapids-cmake-dir}/cpm/gtest.cmake) - - # Find or install GoogleTest - rapids_cpm_gtest(BUILD_EXPORT_SET cuspatial-testing-exports INSTALL_EXPORT_SET cuspatial-testing-exports) - - if(GTest_ADDED) - rapids_export( - BUILD GTest - VERSION ${GTest_VERSION} - EXPORT_SET GTestTargets - GLOBAL_TARGETS gtest gmock gtest_main gmock_main - NAMESPACE GTest:: - ) - - include("${rapids-cmake-dir}/export/find_package_root.cmake") - rapids_export_find_package_root( - BUILD GTest [=[${CMAKE_CURRENT_LIST_DIR}]=] cuspatial-testing-exports - ) - endif() - -endfunction() - -find_and_configure_gtest() diff --git a/cpp/cuproj/CMakeLists.txt b/cpp/cuproj/CMakeLists.txt index 35ba0277e..061a627ca 100644 --- a/cpp/cuproj/CMakeLists.txt +++ b/cpp/cuproj/CMakeLists.txt @@ -95,8 +95,9 @@ include(cmake/thirdparty/get_rmm.cmake) # find or install GoogleTest and Proj if (CUPROJ_BUILD_TESTS) - include(cmake/thirdparty/get_gtest.cmake) - include(cmake/thirdparty/get_proj.cmake) + include(${rapids-cmake-dir}/cpm/gtest.cmake) + rapids_cpm_gtest(BUILD_STATIC) + include(cmake/thirdparty/get_proj.cmake) endif() ################################################################################################### @@ -162,7 +163,7 @@ endif() if(CUPROJ_BUILD_BENCHMARKS) # Find or install GoogleBench include(${rapids-cmake-dir}/cpm/gbench.cmake) - rapids_cpm_gbench() + rapids_cpm_gbench(BUILD_STATIC) # Find or install NVBench Temporarily force downloading of fmt because current versions of nvbench # do not support the latest version of fmt, which is automatically pulled into our conda diff --git a/cpp/cuproj/cmake/thirdparty/get_gtest.cmake b/cpp/cuproj/cmake/thirdparty/get_gtest.cmake deleted file mode 100644 index 0eaa54d25..000000000 --- a/cpp/cuproj/cmake/thirdparty/get_gtest.cmake +++ /dev/null @@ -1,39 +0,0 @@ -# ============================================================================= -# Copyright (c) 2023, 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. -# ============================================================================= - -# This function finds gtest and sets any additional necessary environment variables. -function(find_and_configure_gtest) - include(${rapids-cmake-dir}/cpm/gtest.cmake) - - # Find or install GoogleTest - rapids_cpm_gtest(BUILD_EXPORT_SET cuproj-testing-exports INSTALL_EXPORT_SET cuproj-testing-exports) - - if(GTest_ADDED) - rapids_export( - BUILD GTest - VERSION ${GTest_VERSION} - EXPORT_SET GTestTargets - GLOBAL_TARGETS gtest gmock gtest_main gmock_main - NAMESPACE GTest:: - ) - - include("${rapids-cmake-dir}/export/find_package_root.cmake") - rapids_export_find_package_root( - BUILD GTest [=[${CMAKE_CURRENT_LIST_DIR}]=] cuproj-testing-exports - ) - endif() - -endfunction() - -find_and_configure_gtest() diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md index 74db9350f..325c9ce79 100644 --- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md +++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md @@ -84,7 +84,7 @@ Examples: ```c++ template -void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_memory_resource* mr) +void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_async_resource_ref mr) { ... } @@ -233,9 +233,10 @@ std::unique_ptr points_in_spatial_window( cudf::column_view const& y); ``` -## RMM Memory Resources (`rmm::device_memory_resource`) +## Memory Resources (`rmm::device_memory_resource`) -libcuspatial allocates all device memory via RMM memory resources (MR). See the +libcuspatial allocates all device memory via RMM memory resources (MR) or CUDA MRs. Either type +can be passed to libcuspatial functions via `rmm::device_async_resource_ref` parameters. See the [RMM documentation](https://github.com/rapidsai/rmm/blob/main/README.md) for details. ### Current Device Memory Resource @@ -245,6 +246,27 @@ RMM provides a "default" memory resource for each device that can be accessed an respectively. All memory resource parameters should be defaulted to use the return value of `rmm::mr::get_current_device_resource()`. +### Resource Refs + +Memory resources are passed via resource ref parameters. A resource ref is memory resource wrapper +that enables consumers to specify properties of resources that they expect. These are defined +in the `cuda::mr` namespace of libcu++, but RMM provides some convenience wrappers in +`rmm/resource_ref.hpp`: + - `rmm::device_resource_ref` accepts a memory resource that provides synchronous allocation + of device-accessible memory. + - `rmm::device_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of device-accessible memory. + - `rmm::host_resource_ref` accepts a memory resource that provides synchronous allocation of host- + accessible memory. + - `rmm::host_async_resource_ref` accepts a memory resource that provides stream-ordered allocation + of host-accessible memory. + - `rmm::host_device_resource_ref` accepts a memory resource that provides synchronous allocation of + host- and device-accessible memory. + - `rmm::host_device_async_resource_ref` accepts a memory resource that provides stream-ordered + allocation of host- and device-accessible memory. + +See the libcu++ [docs on `resource_ref`](https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_resource/resource_ref.html) for more information. + # libcuspatial API and Implementation This section provides specifics about the structure and implementation of cuSpatial API functions. @@ -439,8 +461,8 @@ There are a few key points to notice. cuSpatial APIs will not need to use this returned iterator. 9. All APIs that run CUDA device code (including Thrust algorithms) or allocate memory take a CUDA stream on which to execute the device code and allocate memory. - 10. Any API that allocate and return device data (not shown here) should also take an - `rmm::device_memory_resource` to use for output memory allocation. + 10. Any API that allocates and returns device data (not shown here) should also take an + `rmm::device_async_resource_ref` to use for output memory allocation. ### (Multiple) Return Values @@ -542,22 +564,28 @@ control how device memory is allocated. ### Output Memory -Any libcuspatial API that allocates memory that is *returned* to a user must accept a pointer to a -`device_memory_resource` as the last parameter. Inside the API, this memory resource must be used -to allocate any memory for returned objects. It should therefore be passed into functions whose -outputs will be returned. Example: +Any libcuspatial API that allocates memory that is *returned* to a user must accept a +`rmm::device_async_resource_ref` as the last parameter. Inside the API, this memory resource must +be used to allocate any memory for returned objects. It should therefore be passed into functions +whose outputs will be returned. Example: ```c++ // Returned `column` contains newly allocated memory, // therefore the API must accept a memory resource pointer std::unique_ptr returns_output_memory( - ..., rmm::device_memory_resource * mr = rmm::mr::get_current_device_resource()); + ..., rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); // This API does not allocate any new *output* memory, therefore // a memory resource is unnecessary void does_not_allocate_output_memory(...); ``` +This rule automatically applies to all detail APIs that allocate memory. Any detail API may be +called by any public API, and therefore could be allocating memory that is returned to the user. +To support such uses cases, all detail APIs allocating memory resources should accept an `mr` +parameter. Callers are responsible for either passing through a provided `mr` or +`rmm::mr::get_current_device_resource()` as needed. + ### Temporary Memory Not all memory allocated within a libcuspatial API is returned to the caller. Often algorithms must @@ -566,7 +594,7 @@ obtained from `rmm::mr::get_current_device_resource()` for temporary memory allo ```c++ rmm::device_buffer some_function( - ..., rmm::mr::device_memory_resource mr * = rmm::mr::get_current_device_resource()) { + ..., rmm::mr::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) { rmm::device_buffer returned_buffer(..., mr); // Returned buffer uses the passed in MR ... rmm::device_buffer temporary_buffer(...); // Temporary buffer uses default MR @@ -578,12 +606,12 @@ rmm::device_buffer some_function( ### Memory Management libcuspatial code eschews raw pointers and direct memory allocation. Use RMM classes built to -use [`device_memory_resource`](https://github.com/rapidsai/rmm/#device_memory_resource) for device +use [memory resources](https://github.com/rapidsai/rmm/#device_memory_resource) for device memory allocation with automated lifetime management. #### rmm::device_buffer Allocates a specified number of bytes of untyped, uninitialized device memory using a -`device_memory_resource`. If no resource is explicitly provided, uses +memory resource. If no `rmm::device_async_resource_ref` is explicitly provided, uses `rmm::mr::get_current_device_resource()`. `rmm::device_buffer` is movable and copyable on a stream. A copy performs a deep copy of the diff --git a/cpp/include/cuspatial/bounding_boxes.hpp b/cpp/include/cuspatial/bounding_boxes.hpp index 9229eb7c2..2134bc1d1 100644 --- a/cpp/include/cuspatial/bounding_boxes.hpp +++ b/cpp/include/cuspatial/bounding_boxes.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -53,7 +54,7 @@ std::unique_ptr linestring_bounding_boxes( cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute minimum bounding box for each polygon in a list. @@ -80,8 +81,8 @@ std::unique_ptr polygon_bounding_boxes( cudf::column_view const& ring_offsets, cudf::column_view const& x, cudf::column_view const& y, - double expansion_radius = 0.0, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + double expansion_radius = 0.0, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh index 3d06c93b1..a444ae64a 100644 --- a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh +++ b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -65,7 +66,7 @@ compute_point_keys_and_sorted_indices(PointIt points_first, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); rmm::device_uvector keys(num_points, stream); @@ -259,7 +260,7 @@ inline auto make_full_levels(PointIt points_first, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); // Compute point keys and sort into bottom-level quadrants diff --git a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh index ffbb79c0d..df28ecca1 100644 --- a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh +++ b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -21,6 +21,7 @@ #include #include #include +#include #include #include @@ -84,9 +85,9 @@ inline rmm::device_uvector flatten_point_keys( keys_and_levels + num_valid_nodes, flattened_keys.begin(), [last_level = max_depth - 1] __device__(auto const& val) { - bool is_parent{false}; - uint32_t key{}, level{}; - thrust::tie(key, level, is_parent) = val; + auto& key = thrust::get<0>(val); + auto& level = thrust::get<1>(val); + auto& is_parent = thrust::get<2>(val); // if this is a parent node, return max_key. otherwise // compute the key for one level up the tree. Leaf nodes // whose keys are zero will be removed in a subsequent @@ -309,7 +310,7 @@ inline rmm::device_uvector construct_non_leaf_indicator( int32_t num_parent_nodes, int32_t num_valid_nodes, int32_t max_size, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { // diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index 1d88e5e36..818c6560b 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -34,6 +34,7 @@ #include #include #include +#include #include #include @@ -142,7 +143,7 @@ std::unique_ptr> compute_types_buffer( OffsetRangeB points_offset, OffsetRangeB segments_offset, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto types_buffer = std::make_unique>(union_column_size, stream, mr); thrust::tabulate(rmm::exec_policy(stream), @@ -162,7 +163,7 @@ std::unique_ptr> compute_types_buffer( template std::unique_ptr> compute_offset_buffer( rmm::device_uvector const& types_buffer, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto N = types_buffer.size(); @@ -202,7 +203,7 @@ template pairwise_linestring_intersection( MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { using types_t = typename linestring_intersection_result::types_t; diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh index e513cf651..f1a365555 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -29,6 +29,7 @@ #include #include #include +#include #include #include @@ -228,7 +229,7 @@ struct linestring_intersection_intermediates { /** @brief Construct a zero-pair, zero-geometry intermediate object */ linestring_intersection_intermediates(rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) : offsets(std::make_unique>(1, stream)), geoms(std::make_unique>(0, stream, mr)), lhs_linestring_ids(std::make_unique>(0, stream)), @@ -244,7 +245,7 @@ struct linestring_intersection_intermediates { std::size_t num_geoms, rmm::device_uvector const& num_geoms_per_pair, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) : offsets(std::make_unique>(num_pairs + 1, stream)), geoms(std::make_unique>(num_geoms, stream, mr)), lhs_linestring_ids(std::make_unique>(num_geoms, stream)), @@ -472,7 +473,7 @@ std::pair, index_t>, linestring_intersection_intermediates, index_t>> pairwise_linestring_intersection_with_duplicates(MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { static_assert(std::is_integral_v, "Index type must be integral."); diff --git a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh index 43ff7b0f6..472bb29b0 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh @@ -22,6 +22,8 @@ #include #include +#include + #include #include @@ -39,7 +41,7 @@ join_quadtree_and_bounding_boxes(point_quadtree_ref quadtree, T scale, int8_t max_depth, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { static_assert(is_same>(), "Iterator value_type mismatch"); diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh index 1ae0ff1bf..68a2e9596 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -24,13 +24,16 @@ #include #include +#include #include #include +#include #include +#include #include -#include +#include namespace cuspatial { namespace detail { @@ -56,7 +59,7 @@ struct compute_poly_and_point_indices { using IndexType = iterator_value_type; inline thrust::tuple __device__ - operator()(IndexType const global_index) const + operator()(std::uint64_t const global_index) const { auto const [quad_poly_index, local_point_index] = get_quad_and_local_point_indices(global_index, point_offsets_begin, point_offsets_end); @@ -108,7 +111,7 @@ std::pair, rmm::device_uvector> quadtr PointIterator points_first, MultiPolygonRange polygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using T = iterator_vec_base_type; @@ -117,16 +120,26 @@ std::pair, rmm::device_uvector> quadtr auto num_poly_quad_pairs = std::distance(poly_indices_first, poly_indices_last); - auto quad_lengths_iter = - thrust::make_permutation_iterator(quadtree.length_begin(), quad_indices_first); + // The quadtree length is an iterator of uint32_t, but we have to transform into uint64_t values + // so the thrust::inclusive_scan accumulates into uint64_t outputs. Changing the output iterator + // to uint64_t isn't sufficient to achieve this behavior. + auto quad_lengths_iter = thrust::make_transform_iterator( + thrust::make_permutation_iterator(quadtree.length_begin(), quad_indices_first), + cuda::proclaim_return_type([] __device__(IndexType const& i) -> std::uint64_t { + return static_cast(i); + })); auto quad_offsets_iter = thrust::make_permutation_iterator(quadtree.offset_begin(), quad_indices_first); - // Compute a "local" set of zero-based point offsets from number of points in each quadrant + // Compute a "local" set of zero-based point offsets from the number of points in each quadrant. + // // Use `num_poly_quad_pairs + 1` as the length so that the last element produced by // `inclusive_scan` is the total number of points to be tested against any polygon. - rmm::device_uvector local_point_offsets(num_poly_quad_pairs + 1, stream); + // + // Accumulate into uint64_t, because the prefix sums can overflow the size of uint32_t + // when testing a large number of polygons against a large quadtree. + rmm::device_uvector local_point_offsets(num_poly_quad_pairs + 1, stream); // inclusive scan of quad_lengths_iter thrust::inclusive_scan(rmm::exec_policy(stream), @@ -135,21 +148,27 @@ std::pair, rmm::device_uvector> quadtr local_point_offsets.begin() + 1); // Ensure local point offsets starts at 0 - IndexType init{0}; + std::uint64_t init{0}; local_point_offsets.set_element_async(0, init, stream); // The last element is the total number of points to test against any polygon. auto num_total_points = local_point_offsets.back_element(stream); - // Allocate the output polygon and point index pair vectors - rmm::device_uvector poly_indices(num_total_points, stream); - rmm::device_uvector point_indices(num_total_points, stream); - - auto poly_and_point_indices = - thrust::make_zip_iterator(poly_indices.begin(), point_indices.begin()); - - // Enumerate the point X/Ys using the sorted `point_indices` (from quadtree construction) - auto point_xys_iter = thrust::make_permutation_iterator(points_first, point_indices_first); + // The largest supported input size for thrust::count_if/copy_if is INT32_MAX. + // This functor iterates over the input space and processes up to INT32_MAX elements at a time. + std::uint64_t max_points_to_test = std::numeric_limits::max(); + auto count_in_chunks = [&](auto const& func) { + std::uint64_t memo{}; + for (std::uint64_t offset{0}; offset < num_total_points; offset += max_points_to_test) { + memo += func(memo, offset, std::min(max_points_to_test, num_total_points - offset)); + } + return memo; + }; + + detail::test_poly_point_intersection test_poly_point_pair{ + // Enumerate the point X/Ys using the sorted `point_indices` (from quadtree construction) + thrust::make_permutation_iterator(points_first, point_indices_first), + polygons}; // Compute the combination of polygon and point index pairs. For each polygon/quadrant pair, // enumerate pairs of (poly_index, point_index) for each point in each quadrant. @@ -162,28 +181,57 @@ std::pair, rmm::device_uvector> quadtr // pp_pairs.append((polygon, point)) // ``` // - auto global_to_poly_and_point_indices = detail::make_counting_transform_iterator( - 0, - detail::compute_poly_and_point_indices{quad_offsets_iter, - local_point_offsets.begin(), - local_point_offsets.end(), - poly_indices_first}); - - // Compute the number of intersections by removing (poly, point) pairs that don't intersect - auto num_intersections = thrust::distance( - poly_and_point_indices, - thrust::copy_if(rmm::exec_policy(stream), - global_to_poly_and_point_indices, - global_to_poly_and_point_indices + num_total_points, - poly_and_point_indices, - detail::test_poly_point_intersection{point_xys_iter, polygons})); - - poly_indices.resize(num_intersections, stream); - poly_indices.shrink_to_fit(stream); - point_indices.resize(num_intersections, stream); - point_indices.shrink_to_fit(stream); - - return std::pair{std::move(poly_indices), std::move(point_indices)}; + auto global_to_poly_and_point_indices = [&](auto offset = 0) { + return detail::make_counting_transform_iterator( + offset, + detail::compute_poly_and_point_indices{quad_offsets_iter, + local_point_offsets.begin(), + local_point_offsets.end(), + poly_indices_first}); + }; + + auto run_quadtree_point_in_polygon = [&](auto output_size) { + // Allocate the output polygon and point index pair vectors + rmm::device_uvector poly_indices(output_size, stream); + rmm::device_uvector point_indices(output_size, stream); + + auto num_intersections = count_in_chunks([&](auto memo, auto offset, auto size) { + auto poly_and_point_indices = + thrust::make_zip_iterator(poly_indices.begin(), point_indices.begin()) + memo; + // Remove (poly, point) pairs that don't intersect + return thrust::distance(poly_and_point_indices, + thrust::copy_if(rmm::exec_policy(stream), + global_to_poly_and_point_indices(offset), + global_to_poly_and_point_indices(offset) + size, + poly_and_point_indices, + test_poly_point_pair)); + }); + + if (num_intersections < output_size) { + poly_indices.resize(num_intersections, stream); + point_indices.resize(num_intersections, stream); + poly_indices.shrink_to_fit(stream); + point_indices.shrink_to_fit(stream); + } + + return std::pair{std::move(poly_indices), std::move(point_indices)}; + }; + + try { + // First attempt to run the hit test assuming allocating space for all possible intersections + // fits into the available memory. + return run_quadtree_point_in_polygon(num_total_points); + } catch (rmm::out_of_memory const&) { + // If we OOM the first time, pre-compute the number of hits and allocate only that amount of + // space for the output buffers. This halves performance, but it should at least return valid + // results. + return run_quadtree_point_in_polygon(count_in_chunks([&](auto memo, auto offset, auto size) { + return thrust::count_if(rmm::exec_policy(stream), + global_to_poly_and_point_indices(offset), + global_to_poly_and_point_indices(offset) + size, + test_poly_point_pair); + })); + } } } // namespace cuspatial diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh index 6335f133f..5b61882b6 100644 --- a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh +++ b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -156,7 +157,7 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_ PointIterator points_first, MultiLinestringRange linestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestrings.num_multilinestrings() == linestrings.num_linestrings(), "Only one linestring per multilinestring currently supported."); diff --git a/cpp/include/cuspatial/detail/point_quadtree.cuh b/cpp/include/cuspatial/detail/point_quadtree.cuh index 7e0d70dac..19922af35 100644 --- a/cpp/include/cuspatial/detail/point_quadtree.cuh +++ b/cpp/include/cuspatial/detail/point_quadtree.cuh @@ -24,6 +24,7 @@ #include #include #include +#include #include #include @@ -48,7 +49,7 @@ inline point_quadtree make_quad_tree(rmm::device_uvector& keys, int32_t max_size, int32_t level_1_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // count the number of child nodes auto num_child_nodes = thrust::reduce(rmm::exec_policy(stream), @@ -155,7 +156,7 @@ inline point_quadtree make_leaf_tree(rmm::device_uvector& keys, rmm::device_uvector& lengths, int32_t num_top_quads, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { rmm::device_uvector levels(num_top_quads, stream, mr); rmm::device_uvector is_internal_node(num_top_quads, stream, mr); @@ -195,7 +196,7 @@ inline std::pair, point_quadtree> construct_quadtr T scale, int8_t max_depth, int32_t max_size, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { // Construct the full set of non-empty subquadrants starting from the lowest level. @@ -243,7 +244,7 @@ std::pair, point_quadtree> quadtree_on_points( int8_t max_depth, int32_t max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = thrust::distance(points_first, points_last); if (num_points <= 0) { diff --git a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh index b9b53bfc0..03ad0fe27 100644 --- a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -74,6 +74,7 @@ template +CUSPATIAL_HOST_DEVICE multilinestring_range::multilinestring_range( GeometryIterator geometry_begin, GeometryIterator geometry_end, diff --git a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh index 6da914181..11b8a594d 100644 --- a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh +++ b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -63,7 +64,7 @@ void order_trajectories(IdInputIt ids_first, PointOutputIt points_out_first, TimestampOutputIt timestamps_out_first, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using id_type = iterator_value_type; using timestamp_type = iterator_value_type; @@ -116,7 +117,7 @@ std::unique_ptr> derive_trajectories( PointOutputIt points_out_first, TimestampOutputIt timestamps_out_first, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { detail::order_trajectories(ids_first, ids_last, diff --git a/cpp/include/cuspatial/distance.hpp b/cpp/include/cuspatial/distance.hpp index 11cec4fae..76c920b11 100644 --- a/cpp/include/cuspatial/distance.hpp +++ b/cpp/include/cuspatial/distance.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -25,6 +25,7 @@ #include #include +#include #include @@ -53,8 +54,8 @@ std::unique_ptr haversine_distance( cudf::column_view const& a_lat, cudf::column_view const& b_lon, cudf::column_view const& b_lat, - double const radius = EARTH_RADIUS_KM, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + double const radius = EARTH_RADIUS_KM, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief computes Hausdorff distances for all pairs in a collection of spaces @@ -124,7 +125,7 @@ std::pair, cudf::table_view> directed_hausdorff_di cudf::column_view const& xs, cudf::column_view const& ys, cudf::column_view const& space_offsets, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)point-to-(multi)point Cartesian distance @@ -144,7 +145,7 @@ std::pair, cudf::table_view> directed_hausdorff_di std::unique_ptr pairwise_point_distance( geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)points-to-(multi)linestrings Cartesian distance @@ -166,7 +167,7 @@ std::unique_ptr pairwise_point_distance( std::unique_ptr pairwise_point_linestring_distance( geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)point-to-(multi)polygon Cartesian distance @@ -190,7 +191,7 @@ std::unique_ptr pairwise_point_linestring_distance( std::unique_ptr pairwise_point_polygon_distance( geometry_column_view const& multipoints, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)linestring-to-(multi)linestring Cartesian distance @@ -212,7 +213,7 @@ std::unique_ptr pairwise_point_polygon_distance( std::unique_ptr pairwise_linestring_distance( geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)linestring-to-(multi)polygon Cartesian distance @@ -237,7 +238,7 @@ std::unique_ptr pairwise_linestring_distance( std::unique_ptr pairwise_linestring_polygon_distance( geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute pairwise (multi)polygon-to-(multi)polygon Cartesian distance @@ -256,7 +257,7 @@ std::unique_ptr pairwise_linestring_polygon_distance( std::unique_ptr pairwise_polygon_distance( geometry_column_view const& multipolygons1, geometry_column_view const& multipolygons2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/geometry/box.hpp b/cpp/include/cuspatial/geometry/box.hpp index 1041c4de2..4a9f97639 100644 --- a/cpp/include/cuspatial/geometry/box.hpp +++ b/cpp/include/cuspatial/geometry/box.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -40,9 +40,9 @@ class alignas(sizeof(Vertex)) box { private: /** - * @brief Output stream operator for `vec_2d` for human-readable formatting + * @brief Output stream operator for `box` for human-readable formatting */ - friend std::ostream& operator<<(std::ostream& os, cuspatial::box const& b) + friend std::ostream& operator<<(std::ostream& os, cuspatial::box const& b) { return os << "{" << b.v1 << ", " << b.v2 << "}"; } diff --git a/cpp/include/cuspatial/intersection.cuh b/cpp/include/cuspatial/intersection.cuh index 331e6eb0a..c086d84cc 100644 --- a/cpp/include/cuspatial/intersection.cuh +++ b/cpp/include/cuspatial/intersection.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -22,6 +22,7 @@ #include #include #include +#include #include @@ -87,8 +88,8 @@ template pairwise_linestring_intersection( MultiLinestringRange1 multilinestrings1, MultiLinestringRange2 multilinestrings2, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), - rmm::cuda_stream_view stream = rmm::cuda_stream_default); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource(), + rmm::cuda_stream_view stream = rmm::cuda_stream_default); } // namespace cuspatial diff --git a/cpp/include/cuspatial/intersection.hpp b/cpp/include/cuspatial/intersection.hpp index aedaeb521..1869768b1 100644 --- a/cpp/include/cuspatial/intersection.hpp +++ b/cpp/include/cuspatial/intersection.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -21,6 +21,7 @@ #include #include +#include namespace cuspatial { /** @@ -54,6 +55,6 @@ struct linestring_intersection_column_result { linestring_intersection_column_result pairwise_linestring_intersection( geometry_column_view const& multilinestrings_lhs, geometry_column_view const& multilinestrings_rhs, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/nearest_points.hpp b/cpp/include/cuspatial/nearest_points.hpp index 640a1c9b6..46f0c14a2 100644 --- a/cpp/include/cuspatial/nearest_points.hpp +++ b/cpp/include/cuspatial/nearest_points.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,6 +19,8 @@ #include #include +#include + #include namespace cuspatial { @@ -165,7 +167,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points( std::optional> multilinestring_geometry_offsets, cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp index 9bfb13700..0df60c269 100644 --- a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp +++ b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp @@ -21,6 +21,7 @@ #include #include +#include #include @@ -69,6 +70,6 @@ namespace cuspatial { std::unique_ptr pairwise_multipoint_equals_count( geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/point_in_polygon.hpp b/cpp/include/cuspatial/point_in_polygon.hpp index 11e7381c3..c66e2e260 100644 --- a/cpp/include/cuspatial/point_in_polygon.hpp +++ b/cpp/include/cuspatial/point_in_polygon.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -21,6 +21,7 @@ #include #include +#include #include @@ -78,7 +79,7 @@ std::unique_ptr point_in_polygon( cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Given (point, polygon pairs), tests whether the point of each pair is inside the polygon @@ -127,7 +128,7 @@ std::unique_ptr pairwise_point_in_polygon( cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/point_quadtree.cuh b/cpp/include/cuspatial/point_quadtree.cuh index a84b07558..f64cee6f0 100644 --- a/cpp/include/cuspatial/point_quadtree.cuh +++ b/cpp/include/cuspatial/point_quadtree.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -22,6 +22,7 @@ #include #include #include +#include #include #include @@ -175,8 +176,8 @@ std::pair, point_quadtree> quadtree_on_points( T scale, int8_t max_depth, int32_t max_size, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/point_quadtree.hpp b/cpp/include/cuspatial/point_quadtree.hpp index 17e156ea6..72306730f 100644 --- a/cpp/include/cuspatial/point_quadtree.hpp +++ b/cpp/include/cuspatial/point_quadtree.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -74,7 +75,7 @@ std::pair, std::unique_ptr> quadtree_ double scale, int8_t max_depth, cudf::size_type max_size, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/points_in_range.hpp b/cpp/include/cuspatial/points_in_range.hpp index 7bd23c2c0..8637e2b24 100644 --- a/cpp/include/cuspatial/points_in_range.hpp +++ b/cpp/include/cuspatial/points_in_range.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -21,6 +21,7 @@ #include #include +#include #include @@ -60,7 +61,7 @@ std::unique_ptr points_in_range( double range_max_y, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/projection.hpp b/cpp/include/cuspatial/projection.hpp index 1158d8c7c..e71164148 100644 --- a/cpp/include/cuspatial/projection.hpp +++ b/cpp/include/cuspatial/projection.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -50,7 +51,7 @@ std::pair, std::unique_ptr> sinusoid double origin_lat, cudf::column_view const& input_lon, cudf::column_view const& input_lat, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/spatial_join.cuh b/cpp/include/cuspatial/spatial_join.cuh index 9815d5f4f..096af5fb2 100644 --- a/cpp/include/cuspatial/spatial_join.cuh +++ b/cpp/include/cuspatial/spatial_join.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -23,6 +23,7 @@ #include #include +#include #include #include @@ -69,8 +70,8 @@ join_quadtree_and_bounding_boxes( vec_2d const& v_min, T scale, int8_t max_depth, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Test whether the specified points are inside any of the specified polygons. @@ -127,8 +128,8 @@ std::pair, rmm::device_uvector> quadtr PointIndexIterator point_indices_last, PointIterator points_first, MultiPolygonRange polygons, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances @@ -184,8 +185,8 @@ quadtree_point_to_nearest_linestring( PointIndexIterator point_indices_last, PointIterator points_first, MultiLinestringRange linestrings, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); } // namespace cuspatial diff --git a/cpp/include/cuspatial/spatial_join.hpp b/cpp/include/cuspatial/spatial_join.hpp index 579a3c4f1..e3bec92d8 100644 --- a/cpp/include/cuspatial/spatial_join.hpp +++ b/cpp/include/cuspatial/spatial_join.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -19,6 +19,7 @@ #include #include +#include #include @@ -71,7 +72,7 @@ std::unique_ptr join_quadtree_and_bounding_boxes( double y_max, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Test whether the specified points are inside any of the specified polygons. @@ -122,7 +123,7 @@ std::unique_ptr quadtree_point_in_polygon( cudf::column_view const& ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances @@ -171,7 +172,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_offsets, cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial/trajectory.cuh b/cpp/include/cuspatial/trajectory.cuh index f6c5d76f4..dc5ec9708 100644 --- a/cpp/include/cuspatial/trajectory.cuh +++ b/cpp/include/cuspatial/trajectory.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -85,8 +86,8 @@ std::unique_ptr> derive_trajectories( IdOutputIt ids_output_first, PointOutputIt points_output_first, TimestampOutputIt timestamps_output_first, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the total distance (in meters) and average speed (in m/s) of objects in diff --git a/cpp/include/cuspatial/trajectory.hpp b/cpp/include/cuspatial/trajectory.hpp index 5eb5eaf0d..308063569 100644 --- a/cpp/include/cuspatial/trajectory.hpp +++ b/cpp/include/cuspatial/trajectory.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -22,6 +22,7 @@ #include #include +#include #include @@ -62,7 +63,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the distance and speed of objects in a trajectory. Groups the @@ -95,7 +96,7 @@ std::unique_ptr trajectory_distances_and_speeds( cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @brief Compute the spatial bounding boxes of trajectories. Groups the x, y, @@ -127,7 +128,7 @@ std::unique_ptr trajectory_bounding_boxes( cudf::column_view const& object_id, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()); /** * @} // end of doxygen group diff --git a/cpp/include/cuspatial_test/base_fixture.hpp b/cpp/include/cuspatial_test/base_fixture.hpp index c9927b489..820e54e77 100644 --- a/cpp/include/cuspatial_test/base_fixture.hpp +++ b/cpp/include/cuspatial_test/base_fixture.hpp @@ -20,6 +20,7 @@ #include #include +#include namespace cuspatial { namespace test { @@ -28,7 +29,7 @@ namespace test { * @brief Mixin to supply rmm resources for fixtures */ class RMMResourceMixin { - rmm::mr::device_memory_resource* _mr{rmm::mr::get_current_device_resource()}; + rmm::device_async_resource_ref _mr{rmm::mr::get_current_device_resource()}; rmm::cuda_stream_view _stream{rmm::cuda_stream_default}; public: @@ -37,7 +38,7 @@ class RMMResourceMixin { * all tests inheriting from this fixture * @return pointer to memory resource */ - rmm::mr::device_memory_resource* mr() { return _mr; } + rmm::device_async_resource_ref mr() { return _mr; } /** * @brief Returns `cuda_stream_view` that should be used for computation in diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index a3efa5812..8c17ff291 100644 --- a/cpp/include/cuspatial_test/geometry_generator.cuh +++ b/cpp/include/cuspatial_test/geometry_generator.cuh @@ -33,6 +33,8 @@ #include +#include + namespace cuspatial { namespace test { @@ -399,8 +401,9 @@ auto generate_multipoint_array(multipoint_generator_parameter params, std::size_t{0}, params.num_points_per_multipoints); - auto engine_x = deterministic_engine(params.num_points()); - auto engine_y = deterministic_engine(2 * params.num_points()); + auto golden_ratio = (1 + std::sqrt(T{5})) / 2; + auto engine_x = deterministic_engine(golden_ratio * params.num_points()); + auto engine_y = deterministic_engine((1 / golden_ratio) * params.num_points()); auto x_dist = make_uniform_dist(params.lower_left.x, params.upper_right.x); auto y_dist = make_uniform_dist(params.lower_left.y, params.upper_right.y); diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 139d411c5..9d3e8a31a 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -28,6 +28,7 @@ #include #include #include +#include #include #include @@ -56,7 +57,7 @@ auto make_device_vector(std::initializer_list inl) template auto make_device_uvector(std::initializer_list inl, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { std::vector hvec(inl.begin(), inl.end()); auto res = rmm::device_uvector(inl.size(), stream, mr); @@ -362,6 +363,17 @@ class multipoint_array { _geometry_offsets.begin(), _geometry_offsets.end(), _coordinates.begin(), _coordinates.end()}; } + /** + * @brief Copy the offset arrays to host. + */ + auto to_host() const + { + auto geometry_offsets = cuspatial::test::to_host(_geometry_offsets); + auto coordinate_offsets = cuspatial::test::to_host(_coordinates); + + return std::tuple{geometry_offsets, coordinate_offsets}; + } + /// Release ownership auto release() { return std::pair{std::move(_geometry_offsets), std::move(_coordinates)}; } diff --git a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu index 86b6aeb75..7716d231c 100644 --- a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu +++ b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -26,6 +26,7 @@ #include #include +#include #include @@ -43,7 +44,7 @@ std::unique_ptr compute_linestring_bounding_boxes( cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0; @@ -93,7 +94,7 @@ struct dispatch_compute_linestring_bounding_boxes { cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return compute_linestring_bounding_boxes( linestring_offsets, x, y, static_cast(expansion_radius), stream, mr); @@ -109,7 +110,7 @@ std::unique_ptr linestring_bounding_boxes(cudf::column_view const& cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_compute_linestring_bounding_boxes{}, @@ -127,7 +128,7 @@ std::unique_ptr linestring_bounding_boxes(cudf::column_view const& cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0; diff --git a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu index e38e12dfd..c4c092e55 100644 --- a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu +++ b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu @@ -27,6 +27,7 @@ #include #include +#include #include @@ -43,7 +44,7 @@ std::unique_ptr compute_polygon_bounding_boxes(cudf::column_view co cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_polygons = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0; @@ -96,7 +97,7 @@ struct dispatch_compute_polygon_bounding_boxes { cudf::column_view const& y, T expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return compute_polygon_bounding_boxes( poly_offsets, ring_offsets, x, y, expansion_radius, stream, mr); @@ -113,7 +114,7 @@ std::unique_ptr polygon_bounding_boxes(cudf::column_view const& pol cudf::column_view const& y, double expansion_radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_compute_polygon_bounding_boxes{}, @@ -133,7 +134,7 @@ std::unique_ptr polygon_bounding_boxes(cudf::column_view const& pol cudf::column_view const& x, cudf::column_view const& y, double expansion_radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_polys = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0; auto num_rings = ring_offsets.size() > 0 ? ring_offsets.size() - 1 : 0; diff --git a/cpp/src/distance/hausdorff.cu b/cpp/src/distance/hausdorff.cu index 9706b5bf6..b4f2a62d1 100644 --- a/cpp/src/distance/hausdorff.cu +++ b/cpp/src/distance/hausdorff.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -32,6 +32,7 @@ #include #include +#include #include #include @@ -81,7 +82,7 @@ struct hausdorff_functor { cudf::column_view const& ys, cudf::column_view const& space_offsets, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto const num_points = static_cast(xs.size()); auto const num_spaces = static_cast(space_offsets.size()); @@ -120,7 +121,7 @@ std::pair, cudf::table_view> directed_hausdorff_di cudf::column_view const& xs, cudf::column_view const& ys, cudf::column_view const& space_offsets, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(xs.type() == ys.type(), "Inputs `xs` and `ys` must have same type."); CUSPATIAL_EXPECTS(xs.size() == ys.size(), "Inputs `xs` and `ys` must have same length."); diff --git a/cpp/src/distance/haversine.cu b/cpp/src/distance/haversine.cu index 3c4e631da..18a55cd2d 100644 --- a/cpp/src/distance/haversine.cu +++ b/cpp/src/distance/haversine.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -49,7 +50,7 @@ struct haversine_functor { cudf::column_view const& b_lat, T radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { if (a_lon.is_empty()) { return cudf::empty_like(a_lon); } @@ -81,7 +82,7 @@ std::unique_ptr haversine_distance(cudf::column_view const& a_lon, cudf::column_view const& b_lat, double radius, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(radius > 0, "radius must be positive."); @@ -108,7 +109,7 @@ std::unique_ptr haversine_distance(cudf::column_view const& a_lon, cudf::column_view const& b_lon, cudf::column_view const& b_lat, double radius, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::haversine_distance( a_lon, a_lat, b_lon, b_lat, radius, rmm::cuda_stream_default, mr); diff --git a/cpp/src/distance/linestring_distance.cu b/cpp/src/distance/linestring_distance.cu index fe7b99fda..840d82b43 100644 --- a/cpp/src/distance/linestring_distance.cu +++ b/cpp/src/distance/linestring_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -42,7 +43,7 @@ struct pairwise_linestring_distance_launch { std::unique_ptr operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multilinestrings1.size(); @@ -71,7 +72,7 @@ struct pairwise_linestring_distance_functor { std::unique_ptr operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multilinestrings1.geometry_type() == geometry_type_id::LINESTRING && multilinestrings2.geometry_type() == geometry_type_id::LINESTRING, @@ -96,7 +97,7 @@ struct pairwise_linestring_distance_functor { std::unique_ptr pairwise_linestring_distance( geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multilinestrings1.collection_type(), diff --git a/cpp/src/distance/linestring_polygon_distance.cu b/cpp/src/distance/linestring_polygon_distance.cu index b6d370b90..552ab792d 100644 --- a/cpp/src/distance/linestring_polygon_distance.cu +++ b/cpp/src/distance/linestring_polygon_distance.cu @@ -34,6 +34,7 @@ #include #include +#include #include @@ -54,7 +55,7 @@ struct pairwise_linestring_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto multilinestrings_range = make_multilinestring_range(multilinestrings); @@ -87,7 +88,7 @@ struct pairwise_linestring_polygon_distance { std::unique_ptr operator()(geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multilinestrings.geometry_type() == geometry_type_id::LINESTRING && multipolygons.geometry_type() == geometry_type_id::POLYGON, @@ -114,7 +115,7 @@ struct pairwise_linestring_polygon_distance { std::unique_ptr pairwise_linestring_polygon_distance( geometry_column_view const& multilinestrings, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multilinestrings.collection_type(), diff --git a/cpp/src/distance/point_distance.cu b/cpp/src/distance/point_distance.cu index d2349b7a2..960c500d1 100644 --- a/cpp/src/distance/point_distance.cu +++ b/cpp/src/distance/point_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -27,6 +27,7 @@ #include #include +#include #include #include @@ -40,7 +41,7 @@ struct pairwise_point_distance_impl { geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multipoints1.size(); @@ -68,7 +69,7 @@ struct pairwise_point_distance_functor { std::unique_ptr operator()(geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints1.geometry_type() == geometry_type_id::POINT && multipoints2.geometry_type() == geometry_type_id::POINT, @@ -94,7 +95,7 @@ struct pairwise_point_distance_functor { std::unique_ptr pairwise_point_distance(geometry_column_view const& multipoints1, geometry_column_view const& multipoints2, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multipoints1.collection_type(), diff --git a/cpp/src/distance/point_linestring_distance.cu b/cpp/src/distance/point_linestring_distance.cu index 59cb1b460..42d45b152 100644 --- a/cpp/src/distance/point_linestring_distance.cu +++ b/cpp/src/distance/point_linestring_distance.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -48,7 +49,7 @@ struct pairwise_point_linestring_distance_impl { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = multipoints.size(); @@ -79,7 +80,7 @@ struct pairwise_point_linestring_distance_functor { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT && multilinestrings.geometry_type() == geometry_type_id::LINESTRING, @@ -106,7 +107,7 @@ struct pairwise_point_linestring_distance_functor { std::unique_ptr pairwise_point_linestring_distance( geometry_column_view const& multipoints, geometry_column_view const& multilinestrings, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( multipoints.collection_type(), diff --git a/cpp/src/distance/point_polygon_distance.cu b/cpp/src/distance/point_polygon_distance.cu index dd7364fbd..b75b328bc 100644 --- a/cpp/src/distance/point_polygon_distance.cu +++ b/cpp/src/distance/point_polygon_distance.cu @@ -34,6 +34,7 @@ #include #include +#include #include @@ -54,7 +55,7 @@ struct pairwise_point_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto multipoints_range = make_multipoint_range(multipoints); auto multipolygons_range = @@ -83,7 +84,7 @@ struct pairwise_point_polygon_distance { std::unique_ptr operator()(geometry_column_view const& multipoints, geometry_column_view const& multipolygons, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( multipoints.coordinate_type(), @@ -100,7 +101,7 @@ struct pairwise_point_polygon_distance { std::unique_ptr pairwise_point_polygon_distance( geometry_column_view const& multipoints, geometry_column_view const& multipolygons, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT && multipolygons.geometry_type() == geometry_type_id::POLYGON, diff --git a/cpp/src/distance/polygon_distance.cu b/cpp/src/distance/polygon_distance.cu index 678aa3f56..c5a6b0ed1 100644 --- a/cpp/src/distance/polygon_distance.cu +++ b/cpp/src/distance/polygon_distance.cu @@ -32,6 +32,7 @@ #include #include +#include #include @@ -52,7 +53,7 @@ struct pairwise_polygon_distance_impl { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto lhs_range = make_multipolygon_range(lhs); auto rhs_range = make_multipolygon_range(rhs); @@ -80,7 +81,7 @@ struct pairwise_polygon_distance { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POLYGON && rhs.geometry_type() == geometry_type_id::POLYGON, @@ -106,7 +107,7 @@ struct pairwise_polygon_distance { std::unique_ptr pairwise_polygon_distance(geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return multi_geometry_double_dispatch( lhs.collection_type(), rhs.collection_type(), lhs, rhs, rmm::cuda_stream_default, mr); diff --git a/cpp/src/equality/pairwise_multipoint_equals_count.cu b/cpp/src/equality/pairwise_multipoint_equals_count.cu index 71e2b6d2b..abfde9e9c 100644 --- a/cpp/src/equality/pairwise_multipoint_equals_count.cu +++ b/cpp/src/equality/pairwise_multipoint_equals_count.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -26,6 +26,7 @@ #include #include +#include #include #include @@ -47,7 +48,7 @@ struct pairwise_multipoint_equals_count_impl { std::unique_ptr operator()(geometry_column_view const& lhs, geometry_column_view const& rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = lhs.size(); // lhs is a buffer of xy coords auto type = cudf::data_type(cudf::type_to_id()); @@ -78,7 +79,7 @@ struct pairwise_multipoint_equals_count { std::unique_ptr operator()(geometry_column_view lhs, geometry_column_view rhs, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( lhs.coordinate_type(), @@ -94,7 +95,7 @@ struct pairwise_multipoint_equals_count { std::unique_ptr pairwise_multipoint_equals_count(geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POINT && rhs.geometry_type() == geometry_type_id::POINT, diff --git a/cpp/src/indexing/point_quadtree.cu b/cpp/src/indexing/point_quadtree.cu index 4e2f5987f..a8e53f1e1 100644 --- a/cpp/src/indexing/point_quadtree.cu +++ b/cpp/src/indexing/point_quadtree.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -26,6 +26,7 @@ #include #include #include +#include #include @@ -72,7 +73,7 @@ struct dispatch_construct_quadtree { int8_t max_depth, cudf::size_type max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto points = cuspatial::make_vec_2d_iterator(x.begin(), y.begin()); auto [point_indices, tree] = @@ -131,7 +132,7 @@ std::pair, std::unique_ptr> quadtree_ int8_t max_depth, cudf::size_type max_size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(x.type(), dispatch_construct_quadtree{}, @@ -160,7 +161,7 @@ std::pair, std::unique_ptr> quadtree_ double scale, int8_t max_depth, cudf::size_type max_size, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(x.size() == y.size(), "x and y columns must have the same length"); if (x.is_empty() || y.is_empty()) { diff --git a/cpp/src/intersection/linestring_intersection.cu b/cpp/src/intersection/linestring_intersection.cu index 30e60e798..46734f69d 100644 --- a/cpp/src/intersection/linestring_intersection.cu +++ b/cpp/src/intersection/linestring_intersection.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -34,6 +34,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ namespace detail { std::unique_ptr even_sequence(cudf::size_type size, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto res = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id()}, size, @@ -67,7 +68,7 @@ struct pairwise_linestring_intersection_launch { operator()(geometry_column_view const& multilinestrings1, geometry_column_view const& multilinestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using index_t = cudf::size_type; @@ -155,7 +156,7 @@ struct pairwise_linestring_intersection { linestring_intersection_column_result operator()(geometry_column_view const& linestrings1, geometry_column_view const& linestrings2, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestrings1.coordinate_type() == linestrings2.coordinate_type(), "Input linestring coordinates must be the same type."); @@ -176,7 +177,7 @@ struct pairwise_linestring_intersection { linestring_intersection_column_result pairwise_linestring_intersection( geometry_column_view const& lhs, geometry_column_view const& rhs, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::LINESTRING && rhs.geometry_type() == geometry_type_id::LINESTRING, diff --git a/cpp/src/join/quadtree_bbox_filtering.cu b/cpp/src/join/quadtree_bbox_filtering.cu index 59bd369cd..1757938f3 100644 --- a/cpp/src/join/quadtree_bbox_filtering.cu +++ b/cpp/src/join/quadtree_bbox_filtering.cu @@ -27,6 +27,7 @@ #include #include #include +#include #include @@ -42,7 +43,7 @@ struct dispatch_quadtree_bounding_box_join { double y_min, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto bbox_min = cuspatial::make_vec_2d_iterator(bbox.column(0).template begin(), @@ -94,7 +95,7 @@ std::unique_ptr join_quadtree_and_bounding_boxes(cudf::table_view c double y_max, double scale, int8_t max_depth, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(quadtree.num_columns() == 5, "quadtree table must have 5 columns"); CUSPATIAL_EXPECTS(bbox.num_columns() == 4, "bbox table must have 4 columns"); diff --git a/cpp/src/join/quadtree_point_in_polygon.cu b/cpp/src/join/quadtree_point_in_polygon.cu index 75b723aa8..48f9055ba 100644 --- a/cpp/src/join/quadtree_point_in_polygon.cu +++ b/cpp/src/join/quadtree_point_in_polygon.cu @@ -28,6 +28,7 @@ #include #include +#include namespace cuspatial { namespace detail { @@ -52,7 +53,7 @@ struct compute_quadtree_point_in_polygon { cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto poly_indices = poly_quad_pairs.column(0); auto quad_indices = poly_quad_pairs.column(1); @@ -120,7 +121,7 @@ std::unique_ptr quadtree_point_in_polygon(cudf::table_view const& p cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(point_x.type(), compute_quadtree_point_in_polygon{}, @@ -148,7 +149,7 @@ std::unique_ptr quadtree_point_in_polygon(cudf::table_view const& p cudf::column_view const& ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(poly_quad_pairs.num_columns() == 2, "a quadrant-polygon table must have 2 columns"); diff --git a/cpp/src/join/quadtree_point_to_nearest_linestring.cu b/cpp/src/join/quadtree_point_to_nearest_linestring.cu index c423c1cb4..8288ecbc3 100644 --- a/cpp/src/join/quadtree_point_to_nearest_linestring.cu +++ b/cpp/src/join/quadtree_point_to_nearest_linestring.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -28,6 +28,7 @@ #include #include #include +#include #include #include @@ -54,7 +55,7 @@ struct compute_quadtree_point_to_nearest_linestring { cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto linestring_indices = linestring_quad_pairs.column(0); auto quad_indices = linestring_quad_pairs.column(1); @@ -120,7 +121,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(point_x.type(), compute_quadtree_point_to_nearest_linestring{}, @@ -147,7 +148,7 @@ std::unique_ptr quadtree_point_to_nearest_linestring( cudf::column_view const& linestring_offsets, cudf::column_view const& linestring_points_x, cudf::column_view const& linestring_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(linestring_quad_pairs.num_columns() == 2, "a quadrant-linestring table must have 2 columns"); diff --git a/cpp/src/nearest_points/point_linestring_nearest_points.cu b/cpp/src/nearest_points/point_linestring_nearest_points.cu index 9e15cc225..00a4d365f 100644 --- a/cpp/src/nearest_points/point_linestring_nearest_points.cu +++ b/cpp/src/nearest_points/point_linestring_nearest_points.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -56,7 +57,7 @@ struct pairwise_point_linestring_nearest_points_impl { cudf::device_span linestring_offsets, cudf::column_view linestring_points_xy, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto num_points = static_cast(points_xy.size() / 2); auto num_linestring_points = static_cast(linestring_points_xy.size() / 2); @@ -219,7 +220,7 @@ struct pairwise_point_linestring_nearest_points_functor { cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(points_xy.size() % 2 == 0 && linestring_points_xy.size() % 2 == 0, "Points array must contain even number of coordinates."); @@ -258,7 +259,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points( std::optional> multilinestring_geometry_offsets, cudf::device_span linestring_part_offsets, cudf::column_view linestring_points_xy, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return double_boolean_dispatch( multipoint_geometry_offsets.has_value(), diff --git a/cpp/src/point_in_polygon/point_in_polygon.cu b/cpp/src/point_in_polygon/point_in_polygon.cu index e6eb56d08..c22a8da94 100644 --- a/cpp/src/point_in_polygon/point_in_polygon.cu +++ b/cpp/src/point_in_polygon/point_in_polygon.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -28,6 +28,7 @@ #include #include +#include #include #include @@ -56,7 +57,7 @@ struct point_in_polygon_functor { cudf::column_view const& poly_points_y, bool pairwise, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = test_points_x.size(); auto tid = pairwise ? cudf::type_to_id() : cudf::type_to_id(); @@ -112,7 +113,7 @@ std::unique_ptr point_in_polygon(cudf::column_view const& test_poi cudf::column_view const& poly_points_y, bool pairwise, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( test_points_x.size() == test_points_y.size() and poly_points_x.size() == poly_points_y.size(), @@ -155,7 +156,7 @@ std::unique_ptr point_in_polygon(cudf::column_view const& test_poi cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::point_in_polygon(test_points_x, test_points_y, @@ -174,7 +175,7 @@ std::unique_ptr pairwise_point_in_polygon(cudf::column_view const& cudf::column_view const& poly_ring_offsets, cudf::column_view const& poly_points_x, cudf::column_view const& poly_points_y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cuspatial::detail::point_in_polygon(test_points_x, test_points_y, diff --git a/cpp/src/points_in_range/points_in_range.cu b/cpp/src/points_in_range/points_in_range.cu index df6f60c6a..edb843a93 100644 --- a/cpp/src/points_in_range/points_in_range.cu +++ b/cpp/src/points_in_range/points_in_range.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2023, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -42,7 +43,7 @@ struct points_in_range_dispatch { cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto points_begin = cuspatial::make_vec_2d_iterator(x.begin(), y.begin()); @@ -98,7 +99,7 @@ std::unique_ptr points_in_range(double range_min_x, cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(x.type() == y.type(), "Type mismatch between x and y arrays"); CUSPATIAL_EXPECTS(x.size() == y.size(), "Size mismatch between x and y arrays"); @@ -129,7 +130,7 @@ std::unique_ptr points_in_range(double range_min_x, double range_max_y, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return detail::points_in_range( range_min_x, range_max_x, range_min_y, range_max_y, x, y, rmm::cuda_stream_default, mr); diff --git a/cpp/src/projection/sinusoidal_projection.cu b/cpp/src/projection/sinusoidal_projection.cu index 1be4ade40..2552601a4 100644 --- a/cpp/src/projection/sinusoidal_projection.cu +++ b/cpp/src/projection/sinusoidal_projection.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2023, NVIDIA CORPORATION. + * Copyright (c) 2019-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. @@ -26,6 +26,7 @@ #include #include +#include #include #include @@ -53,7 +54,7 @@ struct dispatch_sinusoidal_projection { cudf::column_view const& input_lon, cudf::column_view const& input_lat, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto size = input_lon.size(); auto type = cudf::data_type{cudf::type_to_id()}; @@ -87,7 +88,7 @@ pair_of_columns sinusoidal_projection(double origin_lon, cudf::column_view const& input_lon, cudf::column_view const& input_lat, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( origin_lon >= -180 && origin_lon <= 180 && origin_lat >= -90 && origin_lat <= 90, @@ -116,7 +117,7 @@ pair_of_columns sinusoidal_projection(double origin_lon, double origin_lat, cudf::column_view const& input_lon, cudf::column_view const& input_lat, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return detail::sinusoidal_projection( origin_lon, origin_lat, input_lon, input_lat, rmm::cuda_stream_default, mr); diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu index 8356369ef..e4e6f3745 100644 --- a/cpp/src/trajectory/derive_trajectories.cu +++ b/cpp/src/trajectory/derive_trajectories.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ struct derive_trajectories_dispatch { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto cols = std::vector>{}; cols.reserve(4); @@ -95,7 +96,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::double_type_dispatcher(x.type(), timestamp.type(), @@ -114,7 +115,7 @@ std::pair, std::unique_ptr> derive_tr cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(), diff --git a/cpp/src/trajectory/trajectory_bounding_boxes.cu b/cpp/src/trajectory/trajectory_bounding_boxes.cu index 8d441b9c4..baeea23e4 100644 --- a/cpp/src/trajectory/trajectory_bounding_boxes.cu +++ b/cpp/src/trajectory/trajectory_bounding_boxes.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020-2022, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -27,6 +27,7 @@ #include #include +#include #include @@ -45,7 +46,7 @@ struct dispatch_element { cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // Construct output columns auto type = cudf::data_type{cudf::type_to_id()}; @@ -101,7 +102,7 @@ std::unique_ptr trajectory_bounding_boxes(cudf::size_type num_traje cudf::column_view const& x, cudf::column_view const& y, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( x.type(), dispatch_element{}, num_trajectories, object_id, x, y, stream, mr); @@ -112,7 +113,7 @@ std::unique_ptr trajectory_bounding_boxes(cudf::size_type num_traje cudf::column_view const& object_id, cudf::column_view const& x, cudf::column_view const& y, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS(object_id.size() == x.size() && x.size() == y.size(), "Data size mismatch"); CUSPATIAL_EXPECTS(x.type().id() == y.type().id(), "Data type mismatch"); diff --git a/cpp/src/trajectory/trajectory_distances_and_speeds.cu b/cpp/src/trajectory/trajectory_distances_and_speeds.cu index a27f8b7f8..89474e925 100644 --- a/cpp/src/trajectory/trajectory_distances_and_speeds.cu +++ b/cpp/src/trajectory/trajectory_distances_and_speeds.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-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. @@ -29,6 +29,7 @@ #include #include +#include #include #include @@ -54,7 +55,7 @@ struct dispatch_timestamp { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { // Construct output columns auto type = cudf::data_type{cudf::type_to_id()}; @@ -102,7 +103,7 @@ struct dispatch_element { cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher(timestamp.type(), dispatch_timestamp{}, @@ -132,7 +133,7 @@ std::unique_ptr trajectory_distances_and_speeds(cudf::size_type num cudf::column_view const& y, cudf::column_view const& timestamp, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { return cudf::type_dispatcher( x.type(), dispatch_element{}, num_trajectories, object_id, x, y, timestamp, stream, mr); @@ -144,7 +145,7 @@ std::unique_ptr trajectory_distances_and_speeds(cudf::size_type num cudf::column_view const& x, cudf::column_view const& y, cudf::column_view const& timestamp, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { CUSPATIAL_EXPECTS( x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(), diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 7f7cb112a..f6752c860 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -1,5 +1,5 @@ #============================================================================= -# Copyright (c) 2019-2023, NVIDIA CORPORATION. +# Copyright (c) 2019-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. @@ -158,6 +158,9 @@ ConfigureTest(JOIN_POINT_IN_POLYGON_SMALL_TEST_EXP ConfigureTest(JOIN_POINT_IN_POLYGON_LARGE_TEST_EXP join/quadtree_point_in_polygon_test_large.cu) +ConfigureTest(JOIN_POINT_IN_POLYGON_OOM_TEST_EXP + join/quadtree_point_in_polygon_test_oom.cu) + ConfigureTest(JOIN_POINT_TO_LINESTRING_SMALL_TEST_EXP join/quadtree_point_to_nearest_linestring_test_small.cu) diff --git a/cpp/tests/bounding_boxes/linestring_bounding_boxes_test.cpp b/cpp/tests/bounding_boxes/linestring_bounding_boxes_test.cpp index fb4fab6db..9e5a61ef8 100644 --- a/cpp/tests/bounding_boxes/linestring_bounding_boxes_test.cpp +++ b/cpp/tests/bounding_boxes/linestring_bounding_boxes_test.cpp @@ -21,6 +21,8 @@ #include #include +#include + struct LinestringBoundingBoxErrorTest : public ::testing::Test {}; using T = float; diff --git a/cpp/tests/bounding_boxes/polygon_bounding_boxes_test.cpp b/cpp/tests/bounding_boxes/polygon_bounding_boxes_test.cpp index 83be1f845..f80c2734d 100644 --- a/cpp/tests/bounding_boxes/polygon_bounding_boxes_test.cpp +++ b/cpp/tests/bounding_boxes/polygon_bounding_boxes_test.cpp @@ -23,6 +23,8 @@ #include #include +#include + struct PolygonBoundingBoxErrorTest : public ::testing::Test {}; using T = float; diff --git a/cpp/tests/distance/linestring_polygon_distance_test.cu b/cpp/tests/distance/linestring_polygon_distance_test.cu index 7cc3b6b55..9a09fd84b 100644 --- a/cpp/tests/distance/linestring_polygon_distance_test.cu +++ b/cpp/tests/distance/linestring_polygon_distance_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -26,6 +26,7 @@ #include #include #include +#include #include @@ -35,7 +36,7 @@ using namespace cuspatial::test; template struct PairwiseLinestringPolygonDistanceTest : public BaseFixture { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } void run_single(std::initializer_list multilinestring_geometry_offsets, std::initializer_list multilinestring_part_offsets, diff --git a/cpp/tests/distance/point_polygon_distance_test.cu b/cpp/tests/distance/point_polygon_distance_test.cu index c16d324e0..244ac5ee6 100644 --- a/cpp/tests/distance/point_polygon_distance_test.cu +++ b/cpp/tests/distance/point_polygon_distance_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2023, NVIDIA CORPORATION. + * 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. @@ -25,6 +25,7 @@ #include #include #include +#include #include @@ -36,7 +37,7 @@ double constexpr PI = 3.14159265358979323846; template struct PairwisePointPolygonDistanceTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } void run_single(std::initializer_list>> multipoints, std::initializer_list multipolygon_geometry_offsets, diff --git a/cpp/tests/intersection/intersection_test_utils.cuh b/cpp/tests/intersection/intersection_test_utils.cuh index ab0971648..3a2112d07 100644 --- a/cpp/tests/intersection/intersection_test_utils.cuh +++ b/cpp/tests/intersection/intersection_test_utils.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -22,6 +22,8 @@ #include #include +#include + #include #include #include @@ -89,7 +91,7 @@ struct order_key_value_pairs { template linestring_intersection_result segment_sort_intersection_result( linestring_intersection_result& result, - rmm::mr::device_memory_resource* mr, + rmm::device_async_resource_ref mr, rmm::cuda_stream_view stream) { auto const num_points = result.points_coords->size(); @@ -189,7 +191,7 @@ auto make_linestring_intersection_result( std::initializer_list rhs_linestring_ids, std::initializer_list rhs_segment_ids, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto d_geometry_collection_offset = make_device_uvector(geometry_collection_offset, stream, mr); diff --git a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu index eb36d5e35..760ab72a4 100644 --- a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu +++ b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022-2023, NVIDIA CORPORATION. + * Copyright (c) 2022-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. @@ -27,6 +27,7 @@ #include #include #include +#include #include @@ -44,7 +45,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list std::initializer_list rhs_linestring_ids, std::initializer_list rhs_segment_ids, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { auto d_offset = make_device_uvector(offsets, stream, mr); auto d_geoms = make_device_uvector(geoms, stream, mr); @@ -66,7 +67,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list template struct LinestringIntersectionIntermediatesRemoveIfTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single(IntermediateType& intermediates, diff --git a/cpp/tests/intersection/linestring_intersection_test.cpp b/cpp/tests/intersection/linestring_intersection_test.cpp index 9b9318104..d91d99ddb 100644 --- a/cpp/tests/intersection/linestring_intersection_test.cpp +++ b/cpp/tests/intersection/linestring_intersection_test.cpp @@ -37,6 +37,7 @@ #include #include +#include #include #include @@ -116,7 +117,7 @@ std::pair> make_linestring_col struct LinestringIntersectionTestBase : public BaseFixture { rmm::cuda_stream_view stream{rmm::cuda_stream_default}; - rmm::mr::device_memory_resource* mr{rmm::mr::get_current_device_resource()}; + rmm::device_async_resource_ref mr{rmm::mr::get_current_device_resource()}; }; template diff --git a/cpp/tests/intersection/linestring_intersection_test.cu b/cpp/tests/intersection/linestring_intersection_test.cu index 7f758821c..d6ec4affd 100644 --- a/cpp/tests/intersection/linestring_intersection_test.cu +++ b/cpp/tests/intersection/linestring_intersection_test.cu @@ -30,6 +30,7 @@ #include #include +#include #include #include @@ -43,7 +44,7 @@ using namespace cuspatial::test; template struct LinestringIntersectionTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single_test(MultiLinestringRange lhs, diff --git a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu index 359b45700..720eb2efa 100644 --- a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu +++ b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu @@ -31,6 +31,7 @@ #include #include #include +#include #include #include @@ -72,7 +73,7 @@ using namespace cuspatial::test; template Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediate, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) + rmm::device_async_resource_ref mr) { using GeomType = typename Intermediate::geometry_t; using IndexType = typename Intermediate::index_t; @@ -126,7 +127,7 @@ Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediat template struct LinestringIntersectionDuplicatesTest : public ::testing::Test { rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; } - rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); } + rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); } template void run_single(MultilinestringRange1 lhs, diff --git a/cpp/tests/join/quadtree_point_in_polygon_test.cpp b/cpp/tests/join/quadtree_point_in_polygon_test.cpp index 592e7a52e..51884756f 100644 --- a/cpp/tests/join/quadtree_point_in_polygon_test.cpp +++ b/cpp/tests/join/quadtree_point_in_polygon_test.cpp @@ -28,6 +28,8 @@ #include #include +#include + #include using T = float; diff --git a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu index d80709486..554d197cc 100644 --- a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu +++ b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu @@ -31,6 +31,7 @@ #include #include #include +#include #include #include @@ -52,18 +53,18 @@ template inline auto generate_points( std::vector> const& quads, uint32_t points_per_quad, + cuspatial::vec_2d v_min, + cuspatial::vec_2d v_max, std::size_t seed, rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) { - auto engine = cuspatial::test::deterministic_engine(0); - auto uniform = cuspatial::test::make_normal_dist(0.0, 1.0); - auto pgen = cuspatial::test::point_generator(cuspatial::vec_2d{0.0, 0.0}, - cuspatial::vec_2d{1.0, 1.0}, - engine, - engine, - uniform, - uniform); + auto engine_x = cuspatial::test::deterministic_engine(42); + auto engine_y = cuspatial::test::deterministic_engine(137); + auto uniform_x = cuspatial::test::make_normal_dist(v_min.x, v_max.x); + auto uniform_y = cuspatial::test::make_normal_dist(v_min.y, v_max.y); + auto pgen = + cuspatial::test::point_generator(v_min, v_max, engine_x, engine_y, uniform_x, uniform_y); auto num_points = quads.size() * points_per_quad; rmm::device_uvector> points(num_points, stream, mr); @@ -96,12 +97,12 @@ TYPED_TEST(PIPRefineTestLarge, TestLarge) {7, 8, 3, 4}, {0, 4, 4, 8}}; - auto points_in = generate_points(quads, min_size, 0, this->stream()); + auto points_in = generate_points(quads, min_size, v_min, v_max, 0, this->stream()); - auto [point_indices, quadtree] = quadtree_on_points( + auto [point_indices, quadtree] = cuspatial::quadtree_on_points( points_in.begin(), points_in.end(), v_min, v_max, scale, max_depth, min_size, this->stream()); - auto points = rmm::device_uvector>(quads.size() * min_size, this->stream()); + auto points = rmm::device_uvector>(point_indices.size(), this->stream()); thrust::gather(rmm::exec_policy(this->stream()), point_indices.begin(), point_indices.end(), @@ -159,10 +160,13 @@ TYPED_TEST(PIPRefineTestLarge, TestLarge) quadtree, point_indices.begin(), point_indices.end(), - points.begin(), + points_in.begin(), multipolygons, this->stream()); + EXPECT_GT(actual_point_indices.size(), 0); + EXPECT_GT(actual_poly_indices.size(), 0); + thrust::stable_sort_by_key(rmm::exec_policy(this->stream()), actual_point_indices.begin(), actual_point_indices.end(), @@ -171,17 +175,17 @@ TYPED_TEST(PIPRefineTestLarge, TestLarge) { // verify rmm::device_uvector hits(points.size(), this->stream()); - auto points_range = make_multipoint_range( + auto points_range = cuspatial::make_multipoint_range( points.size(), thrust::make_counting_iterator(0), points.size(), points.begin()); - auto polygons_range = make_multipolygon_range(multipolygons.size(), - thrust::make_counting_iterator(0), - multipolygons.size(), - multipolygons.part_offset_begin(), - multipolygons.num_rings(), - multipolygons.ring_offset_begin(), - multipolygons.num_points(), - multipolygons.point_begin()); + auto polygons_range = cuspatial::make_multipolygon_range(multipolygons.size(), + thrust::make_counting_iterator(0), + multipolygons.size(), + multipolygons.part_offset_begin(), + multipolygons.num_rings(), + multipolygons.ring_offset_begin(), + multipolygons.num_points(), + multipolygons.point_begin()); auto hits_end = cuspatial::point_in_polygon(points_range, polygons_range, hits.begin(), this->stream()); @@ -208,6 +212,9 @@ TYPED_TEST(PIPRefineTestLarge, TestLarge) auto d_expected_poly_indices = rmm::device_vector(expected_poly_indices); auto d_expected_point_indices = rmm::device_vector(expected_point_indices); + EXPECT_GT(d_expected_poly_indices.size(), 0); + EXPECT_GT(d_expected_point_indices.size(), 0); + thrust::stable_sort_by_key(rmm::exec_policy(this->stream()), d_expected_point_indices.begin(), d_expected_point_indices.end(), diff --git a/cpp/tests/join/quadtree_point_in_polygon_test_oom.cu b/cpp/tests/join/quadtree_point_in_polygon_test_oom.cu new file mode 100644 index 000000000..2f66918e0 --- /dev/null +++ b/cpp/tests/join/quadtree_point_in_polygon_test_oom.cu @@ -0,0 +1,140 @@ +/* + * 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 +#include +#include + +#include +#include +#include +#include +#include +#include + +#include + +/* + * The test uses the same quadtree structure as in pip_refine_test_small. However, the number of + * randomly generated points under all quadrants (min_size) are increased to be more than the + * number of threads per-block. + */ + +template +struct PIPRefineTestLarge : public cuspatial::test::BaseFixture {}; + +template +struct test_point_in_poly_functor { + cuspatial::multipolygon_range*> + polys; + __device__ inline T operator()(cuspatial::vec_2d point) + { + auto it = cuspatial::make_counting_transform_iterator(0, [&](auto i) { return polys[i][0]; }); + return thrust::count_if(thrust::seq, it, it + polys.num_multipolygons(), [&](auto poly) { + return cuspatial::is_point_in_polygon(point, poly); + }); + } +}; + +using TestTypes = ::testing::Types; + +TYPED_TEST_CASE(PIPRefineTestLarge, TestTypes); + +TYPED_TEST(PIPRefineTestLarge, TestOOM) +{ + using T = TypeParam; + using cuspatial::vec_2d; + using cuspatial::test::make_device_vector; + + using point_t = typename cuspatial::test::multipoint_array, + rmm::device_uvector>>; + + using polys_t = typename cuspatial::test::multipolygon_array, + rmm::device_uvector, + rmm::device_uvector, + rmm::device_uvector>>; + + vec_2d v_min{0.0, 0.0}; + vec_2d v_max{1'000.0, 1'000.0}; + T const scale{1.0}; + std::uint8_t const max_depth{15}; + std::uint32_t const min_size{10'000}; + + std::size_t const num_points{1'000'000}; + std::size_t const num_polys{24'000}; + + rmm::device_uvector> points = [&]() { + point_t points = cuspatial::test::generate_multipoint_array( + cuspatial::test::multipoint_generator_parameter{1, num_points, v_min, v_max}, + this->stream()); + return points.release().second; + }(); + + polys_t multipoly_array = cuspatial::test::generate_multipolygon_array( + cuspatial::test::multipolygon_generator_parameter{ + num_polys, + 1, + 0, + 4, + vec_2d{(v_max - v_min).x / 2, (v_max - v_min).y / 2}, + (v_max - v_min).x / 8}, + this->stream()); + + auto multipolygons = multipoly_array.range(); + + auto expected_size = thrust::count_if(rmm::exec_policy(this->stream()), + points.begin(), + points.end(), + test_point_in_poly_functor{multipolygons}); + + auto [point_indices, quadtree] = cuspatial::quadtree_on_points( + points.begin(), points.end(), v_min, v_max, scale, max_depth, min_size, this->stream()); + + auto bboxes = + rmm::device_uvector>(multipolygons.num_polygons(), this->stream()); + + cuspatial::polygon_bounding_boxes(multipolygons.part_offset_begin(), + multipolygons.part_offset_end(), + multipolygons.ring_offset_begin(), + multipolygons.ring_offset_end(), + multipolygons.point_begin(), + multipolygons.point_end(), + bboxes.begin(), + T{0}, + this->stream()); + + EXPECT_GT(bboxes.size(), 0); + + auto [poly_indices, quad_indices] = cuspatial::join_quadtree_and_bounding_boxes( + quadtree, bboxes.begin(), bboxes.end(), v_min, scale, max_depth, this->stream()); + + auto [actual_poly_indices, actual_point_indices] = + cuspatial::quadtree_point_in_polygon(poly_indices.begin(), + poly_indices.end(), + quad_indices.begin(), + quadtree, + point_indices.begin(), + point_indices.end(), + points.begin(), + multipolygons, + this->stream()); + + EXPECT_GT(actual_point_indices.size(), 0); + EXPECT_GT(actual_poly_indices.size(), 0); + + EXPECT_EQ(actual_point_indices.size(), expected_size); + EXPECT_EQ(actual_poly_indices.size(), expected_size); +} diff --git a/cpp/tests/join/quadtree_point_to_nearest_linestring_test.cpp b/cpp/tests/join/quadtree_point_to_nearest_linestring_test.cpp index bfca5b024..d04848010 100644 --- a/cpp/tests/join/quadtree_point_to_nearest_linestring_test.cpp +++ b/cpp/tests/join/quadtree_point_to_nearest_linestring_test.cpp @@ -28,6 +28,8 @@ #include #include +#include + #include using T = float; diff --git a/dependencies.yaml b/dependencies.yaml index fabb4149e..bd5d25e11 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -18,6 +18,7 @@ files: - depends_on_cuml - depends_on_cupy - run_python_cuspatial + - test_libcuspatial - test_python_cuspatial - test_python_cuproj - notebooks @@ -25,6 +26,7 @@ files: output: none includes: - cuda_version + - test_libcuspatial test_python: output: none includes: @@ -32,6 +34,7 @@ files: - py_version - test_python_cuspatial - test_python_cuproj + - test_cuspatial test_notebooks: output: none includes: @@ -39,6 +42,7 @@ files: - depends_on_cuml - notebooks - py_version + - test_cuspatial checks: output: none includes: @@ -50,6 +54,7 @@ files: - cuda_version - docs - py_version + - test_cuspatial py_build_cuspatial: output: [pyproject] pyproject_dir: python/cuspatial @@ -105,6 +110,7 @@ files: includes: - test_python_cuproj - depends_on_cuspatial + - test_cuspatial channels: - rapidsai @@ -122,10 +128,8 @@ dependencies: packages: - c-compiler - cxx-compiler - - gmock>=1.13.0 - - gtest>=1.13.0 - - libcudf==24.4.* - - librmm==24.4.* + - libcudf==24.6.* + - librmm==24.6.* - proj - sqlite specific: @@ -167,9 +171,7 @@ dependencies: packages: - c-compiler - cxx-compiler - - gmock>=1.13.0 - - gtest>=1.13.0 - - librmm==24.4.* + - librmm==24.6.* - proj - sqlite specific: @@ -360,7 +362,7 @@ dependencies: common: - output_types: conda packages: - - &rmm_conda rmm==24.4.* + - &rmm_conda rmm==24.6.* - output_types: requirements packages: # pip recognizes the index as a global option for the requirements.txt file @@ -371,17 +373,17 @@ dependencies: matrices: - matrix: {cuda: "12.*"} packages: - - rmm-cu12==24.4.* + - rmm-cu12==24.6.* - matrix: {cuda: "11.*"} packages: - - rmm-cu11==24.4.* + - rmm-cu11==24.6.* - {matrix: null, packages: [*rmm_conda]} depends_on_cudf: common: - output_types: conda packages: - - &cudf_conda cudf==24.4.* + - &cudf_conda cudf==24.6.* - output_types: requirements packages: # pip recognizes the index as a global option for the requirements.txt file @@ -392,17 +394,17 @@ dependencies: matrices: - matrix: {cuda: "12.*"} packages: - - cudf-cu12==24.4.* + - cudf-cu12==24.6.* - matrix: {cuda: "11.*"} packages: - - cudf-cu11==24.4.* + - cudf-cu11==24.6.* - {matrix: null, packages: [*cudf_conda]} depends_on_cuml: common: - output_types: conda packages: - - &cuml_conda cuml==24.4.* + - &cuml_conda cuml==24.6.* - output_types: requirements packages: # pip recognizes the index as a global option for the requirements.txt file @@ -413,17 +415,17 @@ dependencies: matrices: - matrix: {cuda: "12.*"} packages: - - cuml-cu12==24.4.* + - cuml-cu12==24.6.* - matrix: {cuda: "11.*"} packages: - - cuml-cu11==24.4.* + - cuml-cu11==24.6.* - {matrix: null, packages: [*cuml_conda]} depends_on_cuspatial: common: - output_types: conda packages: - - &cuspatial_conda cuspatial==24.4.* + - &cuspatial_conda cuspatial==24.6.* - output_types: requirements packages: # pip recognizes the index as a global option for the requirements.txt file @@ -434,10 +436,10 @@ dependencies: matrices: - matrix: {cuda: "12.*"} packages: - - cuspatial-cu12==24.4.* + - cuspatial-cu12==24.6.* - matrix: {cuda: "11.*"} packages: - - cuspatial-cu11==24.4.* + - cuspatial-cu11==24.6.* - {matrix: null, packages: [*cuspatial_conda]} depends_on_cupy: @@ -455,3 +457,16 @@ dependencies: packages: - cupy-cuda11x>=12.0.0 - {matrix: null, packages: [cupy-cuda11x>=12.0.0]} + test_libcuspatial: + common: + - output_types: conda + packages: + - libcuspatial==24.6.* + - libcuspatial-tests==24.6.* + test_cuspatial: + common: + - output_types: conda + packages: + - libcuspatial==24.6.* + - cuspatial==24.6.* + - cuproj==24.6.* diff --git a/docs/cuproj/source/user_guide/cuproj_api_examples.ipynb b/docs/cuproj/source/user_guide/cuproj_api_examples.ipynb index a319082b4..8dd25cb28 100644 --- a/docs/cuproj/source/user_guide/cuproj_api_examples.ipynb +++ b/docs/cuproj/source/user_guide/cuproj_api_examples.ipynb @@ -45,7 +45,7 @@ "metadata": {}, "outputs": [], "source": [ - "# !conda create -n rapids-24.04 --solver=libmamba -c rapidsai -c conda-forge -c nvidia \\ \n", + "# !conda create -n rapids-24.06 --solver=libmamba -c rapidsai -c conda-forge -c nvidia \\ \n", "# cuproj-23.12 python=3.10 cuda-version=12.0" ] }, diff --git a/docs/source/user_guide/cuspatial_api_examples.ipynb b/docs/source/user_guide/cuspatial_api_examples.ipynb index 4ec66b4b8..a20b8cf1e 100644 --- a/docs/source/user_guide/cuspatial_api_examples.ipynb +++ b/docs/source/user_guide/cuspatial_api_examples.ipynb @@ -57,8 +57,8 @@ "metadata": {}, "outputs": [], "source": [ - "# !conda create -n rapids-24.04 -c rapidsai -c conda-forge -c nvidia \\ \n", - "# cuspatial=24.04 python=3.9 cudatoolkit=11.5 " + "# !conda create -n rapids-24.06 -c rapidsai -c conda-forge -c nvidia \\ \n", + "# cuspatial=24.06 python=3.9 cudatoolkit=11.5 " ] }, { diff --git a/python/cuproj/pyproject.toml b/python/cuproj/pyproject.toml index 698482319..54c4a3709 100644 --- a/python/cuproj/pyproject.toml +++ b/python/cuproj/pyproject.toml @@ -18,7 +18,7 @@ requires = [ "cmake>=3.26.4", "cython>=3.0.0", "ninja", - "rmm==24.4.*", + "rmm==24.6.*", "scikit-build-core[pyproject]>=0.7.0", "setuptools", "wheel", @@ -34,7 +34,7 @@ license = { text = "Apache 2.0" } requires-python = ">=3.9" dependencies = [ "cupy-cuda11x>=12.0.0", - "rmm==24.4.*", + "rmm==24.6.*", ] # 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", @@ -49,7 +49,7 @@ classifiers = [ [project.optional-dependencies] test = [ - "cuspatial==24.4.*", + "cuspatial==24.6.*", "geopandas>=0.11.0", "numpy>=1.23,<2.0a0", "pyproj>=3.6.0,<3.7a0", @@ -111,3 +111,11 @@ wheel.packages = ["cuproj"] provider = "scikit_build_core.metadata.regex" input = "cuproj/VERSION" regex = "(?P.*)" + +[tool.pytest.ini_options] +xfail_strict = true +filterwarnings = [ + "error", + # https://github.com/pytest-dev/pytest-cov/issues/557 + "ignore:The --rsyncdir command line argument:DeprecationWarning", +] diff --git a/python/cuspatial/benchmarks/api/bench_api.py b/python/cuspatial/benchmarks/api/bench_api.py index d93e45ebb..d8e24cf33 100644 --- a/python/cuspatial/benchmarks/api/bench_api.py +++ b/python/cuspatial/benchmarks/api/bench_api.py @@ -1,4 +1,6 @@ -# Copyright (c) 2022-2023, NVIDIA CORPORATION. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. +import pathlib + import cupy import geopandas import pytest @@ -223,12 +225,16 @@ def bench_quadtree_point_to_nearest_linestring(benchmark): SCALE = 3 MAX_DEPTH = 7 MIN_SIZE = 125 - host_countries = geopandas.read_file( - geopandas.datasets.get_path("naturalearth_lowres") - ) - host_cities = geopandas.read_file( - geopandas.datasets.get_path("naturalearth_cities") + data_dir = ( + pathlib.Path(__file__).parent.parent.parent + / "cuspatial" + / "tests" + / "data" ) + naturalearth_lowres = data_dir / "naturalearth_lowres.shp" + naturalearth_cities = data_dir / "naturalearth_cities.shp" + host_countries = geopandas.read_file(naturalearth_lowres) + host_cities = geopandas.read_file(naturalearth_cities) gpu_countries = cuspatial.from_geopandas( host_countries[host_countries["geometry"].type == "Polygon"] ) diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance.pxd index 1af0f97e0..98925a516 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/distance.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/distance.pxd @@ -1,11 +1,11 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. +# Copyright (c) 2023-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.utility cimport pair -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table_view cimport table_view +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table_view cimport table_view from cuspatial._lib.cpp.column.geometry_column_view cimport ( geometry_column_view, diff --git a/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd b/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd index 62f3a318c..3a92adae3 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/distance/polygon_distance.pxd @@ -1,8 +1,8 @@ -# Copyright (c) 2023, NVIDIA CORPORATION. +# Copyright (c) 2023-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr -from cudf._lib.cpp.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column cimport column from cuspatial._lib.cpp.column.geometry_column_view cimport ( geometry_column_view, diff --git a/python/cuspatial/cuspatial/_lib/cpp/linestring_bounding_boxes.pxd b/python/cuspatial/cuspatial/_lib/cpp/linestring_bounding_boxes.pxd index fa458a6cf..8622a417c 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/linestring_bounding_boxes.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/linestring_bounding_boxes.pxd @@ -1,9 +1,9 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table cdef extern from "cuspatial/bounding_boxes.hpp" \ diff --git a/python/cuspatial/cuspatial/_lib/cpp/points_in_range.pxd b/python/cuspatial/cuspatial/_lib/cpp/points_in_range.pxd index e6534dc7b..d03b3a133 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/points_in_range.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/points_in_range.pxd @@ -1,9 +1,9 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from cudf._lib.column cimport column, column_view -from cudf._lib.cpp.table.table cimport table, table_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table, table_view cdef extern from "cuspatial/points_in_range.hpp" namespace "cuspatial" nogil: diff --git a/python/cuspatial/cuspatial/_lib/cpp/polygon_bounding_boxes.pxd b/python/cuspatial/cuspatial/_lib/cpp/polygon_bounding_boxes.pxd index a6d28f965..45910626b 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/polygon_bounding_boxes.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/polygon_bounding_boxes.pxd @@ -1,9 +1,9 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table cdef extern from "cuspatial/bounding_boxes.hpp" \ diff --git a/python/cuspatial/cuspatial/_lib/cpp/projection.pxd b/python/cuspatial/cuspatial/_lib/cpp/projection.pxd index 3c85461e7..724269486 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/projection.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/projection.pxd @@ -1,10 +1,10 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view cdef extern from "cuspatial/projection.hpp" namespace "cuspatial" \ diff --git a/python/cuspatial/cuspatial/_lib/cpp/quadtree.pxd b/python/cuspatial/cuspatial/_lib/cpp/quadtree.pxd index be34d98a9..db2dabaca 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/quadtree.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/quadtree.pxd @@ -1,13 +1,13 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libc.stdint cimport int8_t from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table -from cudf._lib.cpp.types cimport size_type +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table +from cudf._lib.pylibcudf.libcudf.types cimport size_type cdef extern from "cuspatial/point_quadtree.hpp" namespace "cuspatial" nogil: diff --git a/python/cuspatial/cuspatial/_lib/cpp/spatial_join.pxd b/python/cuspatial/cuspatial/_lib/cpp/spatial_join.pxd index 948aa77c9..ab96cf53c 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/spatial_join.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/spatial_join.pxd @@ -1,10 +1,10 @@ -# Copyright (c) 2020-2022, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libc.stdint cimport int8_t from libcpp.memory cimport unique_ptr from cudf._lib.column cimport column_view -from cudf._lib.cpp.table.table cimport table, table_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table, table_view cdef extern from "cuspatial/spatial_join.hpp" namespace "cuspatial" nogil: diff --git a/python/cuspatial/cuspatial/_lib/cpp/trajectory.pxd b/python/cuspatial/cuspatial/_lib/cpp/trajectory.pxd index 9031d17eb..5940349cc 100644 --- a/python/cuspatial/cuspatial/_lib/cpp/trajectory.pxd +++ b/python/cuspatial/cuspatial/_lib/cpp/trajectory.pxd @@ -1,12 +1,12 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table -from cudf._lib.cpp.types cimport size_type +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table +from cudf._lib.pylibcudf.libcudf.types cimport size_type cdef extern from "cuspatial/trajectory.hpp" namespace "cuspatial" nogil: diff --git a/python/cuspatial/cuspatial/_lib/distance.pyx b/python/cuspatial/cuspatial/_lib/distance.pyx index cf61773e0..90e7b5ac0 100644 --- a/python/cuspatial/cuspatial/_lib/distance.pyx +++ b/python/cuspatial/cuspatial/_lib/distance.pyx @@ -1,12 +1,12 @@ -# Copyright (c) 2022-2023, NVIDIA CORPORATION. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. from libcpp.memory cimport make_shared, shared_ptr, unique_ptr from libcpp.utility cimport move, pair from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table_view cimport table_view +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table_view cimport table_view from cudf._lib.utils cimport columns_from_table_view from cuspatial._lib.cpp.column.geometry_column_view cimport ( diff --git a/python/cuspatial/cuspatial/_lib/linestring_bounding_boxes.pyx b/python/cuspatial/cuspatial/_lib/linestring_bounding_boxes.pyx index 28815cc57..4023f4965 100644 --- a/python/cuspatial/cuspatial/_lib/linestring_bounding_boxes.pyx +++ b/python/cuspatial/cuspatial/_lib/linestring_bounding_boxes.pyx @@ -1,11 +1,11 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table from cudf._lib.utils cimport columns_from_unique_ptr from cuspatial._lib.cpp.linestring_bounding_boxes cimport ( diff --git a/python/cuspatial/cuspatial/_lib/nearest_points.pyx b/python/cuspatial/cuspatial/_lib/nearest_points.pyx index d8ecebc58..da24357d3 100644 --- a/python/cuspatial/cuspatial/_lib/nearest_points.pyx +++ b/python/cuspatial/cuspatial/_lib/nearest_points.pyx @@ -1,9 +1,9 @@ -# Copyright (c) 2022-2023, NVIDIA CORPORATION. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view from cuspatial._lib.cpp.nearest_points cimport ( pairwise_point_linestring_nearest_points as c_func, diff --git a/python/cuspatial/cuspatial/_lib/points_in_range.pyx b/python/cuspatial/cuspatial/_lib/points_in_range.pyx index e54b60a6b..4c25c2dbf 100644 --- a/python/cuspatial/cuspatial/_lib/points_in_range.pyx +++ b/python/cuspatial/cuspatial/_lib/points_in_range.pyx @@ -1,10 +1,10 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.utility cimport move from cudf._lib.column cimport Column, column_view -from cudf._lib.cpp.table.table cimport table +from cudf._lib.pylibcudf.libcudf.table.table cimport table from cudf._lib.utils cimport data_from_unique_ptr from cuspatial._lib.cpp.points_in_range cimport ( diff --git a/python/cuspatial/cuspatial/_lib/polygon_bounding_boxes.pyx b/python/cuspatial/cuspatial/_lib/polygon_bounding_boxes.pyx index d9419fe7b..0f963fd5a 100644 --- a/python/cuspatial/cuspatial/_lib/polygon_bounding_boxes.pyx +++ b/python/cuspatial/cuspatial/_lib/polygon_bounding_boxes.pyx @@ -1,11 +1,11 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table from cudf._lib.utils cimport columns_from_unique_ptr from cuspatial._lib.cpp.polygon_bounding_boxes cimport ( diff --git a/python/cuspatial/cuspatial/_lib/quadtree.pyx b/python/cuspatial/cuspatial/_lib/quadtree.pyx index de8246bc3..eb548a882 100644 --- a/python/cuspatial/cuspatial/_lib/quadtree.pyx +++ b/python/cuspatial/cuspatial/_lib/quadtree.pyx @@ -1,4 +1,4 @@ -# Copyright (c) 2020, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libc.stdint cimport int8_t from libcpp.memory cimport unique_ptr @@ -6,10 +6,10 @@ from libcpp.pair cimport pair from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table -from cudf._lib.cpp.types cimport size_type +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table +from cudf._lib.pylibcudf.libcudf.types cimport size_type from cudf._lib.utils cimport data_from_unique_ptr from cuspatial._lib.cpp.quadtree cimport ( diff --git a/python/cuspatial/cuspatial/_lib/spatial.pyx b/python/cuspatial/cuspatial/_lib/spatial.pyx index 916d66cec..6698da904 100644 --- a/python/cuspatial/cuspatial/_lib/spatial.pyx +++ b/python/cuspatial/cuspatial/_lib/spatial.pyx @@ -1,12 +1,12 @@ -# Copyright (c) 2019, NVIDIA CORPORATION. +# Copyright (c) 2019-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view from cuspatial._lib.cpp.projection cimport ( sinusoidal_projection as cpp_sinusoidal_projection, diff --git a/python/cuspatial/cuspatial/_lib/spatial_join.pyx b/python/cuspatial/cuspatial/_lib/spatial_join.pyx index aab04c938..045014aba 100644 --- a/python/cuspatial/cuspatial/_lib/spatial_join.pyx +++ b/python/cuspatial/cuspatial/_lib/spatial_join.pyx @@ -1,11 +1,11 @@ -# Copyright (c) 2020-2022, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. from libc.stdint cimport int8_t from libcpp.memory cimport unique_ptr from libcpp.utility cimport move from cudf._lib.column cimport Column, column_view -from cudf._lib.cpp.table.table cimport table, table_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table, table_view from cudf._lib.utils cimport data_from_unique_ptr, table_view_from_table from cuspatial._lib.cpp.spatial_join cimport ( diff --git a/python/cuspatial/cuspatial/_lib/trajectory.pyx b/python/cuspatial/cuspatial/_lib/trajectory.pyx index 012439904..df29dc05a 100644 --- a/python/cuspatial/cuspatial/_lib/trajectory.pyx +++ b/python/cuspatial/cuspatial/_lib/trajectory.pyx @@ -1,14 +1,14 @@ -# Copyright (c) 2019-2020, NVIDIA CORPORATION. +# Copyright (c) 2019-2024, NVIDIA CORPORATION. from libcpp.memory cimport unique_ptr from libcpp.pair cimport pair from libcpp.utility cimport move from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column cimport column -from cudf._lib.cpp.column.column_view cimport column_view -from cudf._lib.cpp.table.table cimport table -from cudf._lib.cpp.types cimport size_type +from cudf._lib.pylibcudf.libcudf.column.column cimport column +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.table.table cimport table +from cudf._lib.pylibcudf.libcudf.types cimport size_type from cudf._lib.utils cimport data_from_unique_ptr from cuspatial._lib.cpp.trajectory cimport ( diff --git a/python/cuspatial/cuspatial/_lib/utils.pxd b/python/cuspatial/cuspatial/_lib/utils.pxd index 0567434d1..ad1714016 100644 --- a/python/cuspatial/cuspatial/_lib/utils.pxd +++ b/python/cuspatial/cuspatial/_lib/utils.pxd @@ -1,5 +1,5 @@ -# Copyright (c) 2022, NVIDIA CORPORATION. -from cudf._lib.cpp.column.column_view cimport column_view +# Copyright (c) 2022-2024, NVIDIA CORPORATION. +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view from cuspatial._lib.cpp.optional cimport nullopt, optional diff --git a/python/cuspatial/cuspatial/_lib/utils.pyx b/python/cuspatial/cuspatial/_lib/utils.pyx index cabad333c..4aea62863 100644 --- a/python/cuspatial/cuspatial/_lib/utils.pyx +++ b/python/cuspatial/cuspatial/_lib/utils.pyx @@ -1,6 +1,6 @@ -# Copyright (c) 2022, NVIDIA CORPORATION. +# Copyright (c) 2022-2024, NVIDIA CORPORATION. from cudf._lib.column cimport Column -from cudf._lib.cpp.column.column_view cimport column_view +from cudf._lib.pylibcudf.libcudf.column.column_view cimport column_view from cuspatial._lib.cpp.optional cimport nullopt, optional diff --git a/python/cuspatial/cuspatial/tests/conftest.py b/python/cuspatial/cuspatial/tests/conftest.py index 7b6ae2d1b..caa53dcea 100644 --- a/python/cuspatial/cuspatial/tests/conftest.py +++ b/python/cuspatial/cuspatial/tests/conftest.py @@ -1,4 +1,5 @@ -# Copyright (c) 2020-2023, NVIDIA CORPORATION. +# Copyright (c) 2020-2024, NVIDIA CORPORATION. +import pathlib import geopandas as gpd import numpy as np @@ -308,13 +309,18 @@ def factory(length): @pytest.fixture -def naturalearth_cities(): - return gpd.read_file(gpd.datasets.get_path("naturalearth_cities")) +def data_dir(): + return pathlib.Path(__file__).parent / "data" @pytest.fixture -def naturalearth_lowres(): - return gpd.read_file(gpd.datasets.get_path("naturalearth_lowres")) +def naturalearth_cities(data_dir): + return gpd.read_file(data_dir / "naturalearth_cities.shp") + + +@pytest.fixture +def naturalearth_lowres(data_dir): + return gpd.read_file(data_dir / "naturalearth_lowres.shp") @pytest.fixture(scope="session") diff --git a/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shp b/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shp new file mode 100644 index 000000000..a7bcefa0a Binary files /dev/null and b/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shp differ diff --git a/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shx b/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shx new file mode 100644 index 000000000..ce83da735 Binary files /dev/null and b/python/cuspatial/cuspatial/tests/data/naturalearth_cities.shx differ diff --git a/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shp b/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shp new file mode 100644 index 000000000..5a3440b25 Binary files /dev/null and b/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shp differ diff --git a/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shx b/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shx new file mode 100644 index 000000000..275512f80 Binary files /dev/null and b/python/cuspatial/cuspatial/tests/data/naturalearth_lowres.shx differ diff --git a/python/cuspatial/cuspatial/tests/test_geocolumn_accessor.py b/python/cuspatial/cuspatial/tests/test_geocolumn_accessor.py index db074ae88..c4ec694b6 100644 --- a/python/cuspatial/cuspatial/tests/test_geocolumn_accessor.py +++ b/python/cuspatial/cuspatial/tests/test_geocolumn_accessor.py @@ -1,4 +1,4 @@ -# Copyright 2022 NVIDIA Corporation +# Copyright (c) 2022-2024, NVIDIA CORPORATION import cupy as cp import geopandas as gpd @@ -18,10 +18,10 @@ "range, expected", [[slice(0, 3), [0, 3, 4, 5]], [slice(3, 6), [0, 30, 40, 41]]], ) -def test_GeoColumnAccessor_polygon_offset(range, expected): - gpdf = cuspatial.from_geopandas( - gpd.read_file(gpd.datasets.get_path("naturalearth_lowres")) - ) +def test_GeoColumnAccessor_polygon_offset( + range, expected, naturalearth_lowres +): + gpdf = cuspatial.from_geopandas(naturalearth_lowres) shorter = gpdf[range]["geometry"] expected = cp.array(expected) got = shorter.polygons.geometry_offset diff --git a/python/cuspatial/cuspatial/tests/test_geodataframe.py b/python/cuspatial/cuspatial/tests/test_geodataframe.py index 876e30dc5..718046c17 100644 --- a/python/cuspatial/cuspatial/tests/test_geodataframe.py +++ b/python/cuspatial/cuspatial/tests/test_geodataframe.py @@ -189,10 +189,9 @@ def test_interleaved_polygons(gpdf, polys): ) -def test_to_geopandas_with_geopandas_dataset(): - df = gpd.read_file(gpd.datasets.get_path("naturalearth_lowres")) - gdf = cuspatial.from_geopandas(df) - assert_eq_geo_df(df, gdf.to_geopandas()) +def test_to_geopandas_with_geopandas_dataset(naturalearth_lowres): + gdf = cuspatial.from_geopandas(naturalearth_lowres) + assert_eq_geo_df(naturalearth_lowres, gdf.to_geopandas()) def test_to_shapely_random(): @@ -327,14 +326,12 @@ def test_boolmask(gpdf, df_boolmask): reason="Size discrepancies between Python versions. See " "https://github.com/rapidsai/cuspatial/issues/1352", ) -def test_memory_usage(gs): +def test_memory_usage(gs, data_dir): assert gs.memory_usage() == 224 - host_dataframe = gpd.read_file( - gpd.datasets.get_path("naturalearth_lowres") - ) + host_dataframe = gpd.read_file(data_dir / "naturalearth_lowres.shp") gpu_dataframe = cuspatial.from_geopandas(host_dataframe) # The df size is 8kb of cudf rows and 217kb of the geometry column - assert gpu_dataframe.memory_usage().sum() == 224945 + assert gpu_dataframe.memory_usage().sum() == 216793 def test_from_dict(): diff --git a/python/cuspatial/cuspatial/tests/test_geoseries.py b/python/cuspatial/cuspatial/tests/test_geoseries.py index 395b58512..6e97d037b 100644 --- a/python/cuspatial/cuspatial/tests/test_geoseries.py +++ b/python/cuspatial/cuspatial/tests/test_geoseries.py @@ -551,11 +551,8 @@ def test_memory_usage_simple(gs): assert cugs.memory_usage() == 1616 -def test_memory_usage_large(): - host_dataframe = gpd.read_file( - gpd.datasets.get_path("naturalearth_lowres") - ) - geometry = cuspatial.from_geopandas(host_dataframe)["geometry"] +def test_memory_usage_large(naturalearth_lowres): + geometry = cuspatial.from_geopandas(naturalearth_lowres)["geometry"] # the geometry column from naturalearth_lowres is 217kb of coordinates assert geometry.memory_usage() == 216793 diff --git a/python/cuspatial/pyproject.toml b/python/cuspatial/pyproject.toml index e58b6bd54..b4ae45e9e 100644 --- a/python/cuspatial/pyproject.toml +++ b/python/cuspatial/pyproject.toml @@ -16,10 +16,10 @@ build-backend = "scikit_build_core.build" requires = [ "cmake>=3.26.4", - "cudf==24.4.*", + "cudf==24.6.*", "cython>=3.0.0", "ninja", - "rmm==24.4.*", + "rmm==24.6.*", "scikit-build-core[pyproject]>=0.7.0", "setuptools", "wheel", @@ -36,10 +36,10 @@ authors = [ license = { text = "Apache 2.0" } requires-python = ">=3.9" dependencies = [ - "cudf==24.4.*", + "cudf==24.6.*", "geopandas>=0.11.0", "numpy>=1.23,<2.0a0", - "rmm==24.4.*", + "rmm==24.6.*", ] # 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", @@ -124,5 +124,9 @@ regex = "(?P.*)" [tool.pytest.ini_options] xfail_strict = true filterwarnings = [ - "error:::cudf" + "error:::cudf", + "error::FutureWarning", + "error::DeprecationWarning", + # https://github.com/pytest-dev/pytest-cov/issues/557 + "ignore:The --rsyncdir command line argument:DeprecationWarning", ]