diff --git a/.devcontainer/Dockerfile b/.devcontainer/Dockerfile index 688975418..68194d0f9 100644 --- a/.devcontainer/Dockerfile +++ b/.devcontainer/Dockerfile @@ -1,8 +1,8 @@ # syntax=docker/dockerfile:1.5 -ARG CUDA=11.8 +ARG CUDA=12.0 ARG LLVM=16 -ARG RAPIDS=23.06 +ARG RAPIDS=23.08 ARG DISTRO=ubuntu22.04 ARG REPO=rapidsai/devcontainers @@ -10,19 +10,19 @@ ARG PYTHON_PACKAGE_MANAGER=conda FROM ${REPO}:${RAPIDS}-cpp-llvm${LLVM}-cuda${CUDA}-${DISTRO} as pip-base -FROM ${REPO}:${RAPIDS}-cpp-llvm${LLVM}-cuda${CUDA}-mambaforge-${DISTRO} as conda-base +FROM ${REPO}:${RAPIDS}-cpp-mambaforge-${DISTRO} as conda-base + +COPY --from=pip-base /etc/skel/.config/clangd/config.yaml /etc/skel/.config/clangd/config.yaml FROM ${PYTHON_PACKAGE_MANAGER}-base +ARG CUDA +ENV CUDAARCHS="RAPIDS" +ENV CUDA_VERSION="${CUDA_VERSION:-${CUDA}}" + ARG PYTHON_PACKAGE_MANAGER ENV PYTHON_PACKAGE_MANAGER="${PYTHON_PACKAGE_MANAGER}" -USER coder - -RUN /bin/bash -c 'mkdir -m 0755 -p ~/.{aws,cache,conda,config/pip,local}' - -WORKDIR /home/coder/ - ENV PYTHONSAFEPATH="1" ENV PYTHONUNBUFFERED="1" ENV PYTHONDONTWRITEBYTECODE="1" @@ -30,5 +30,4 @@ ENV PYTHONDONTWRITEBYTECODE="1" ENV SCCACHE_REGION="us-east-2" ENV SCCACHE_BUCKET="rapids-sccache-devs" ENV VAULT_HOST="https://vault.ops.k8s.rapids.ai" - ENV HISTFILE="/home/coder/.cache/._bash_history" diff --git a/.devcontainer/README.md b/.devcontainer/README.md index 40d330b1c..c2f5fe9bc 100644 --- a/.devcontainer/README.md +++ b/.devcontainer/README.md @@ -87,11 +87,18 @@ $ .devcontainer/launch.sh unified pip #### Isolated mode -`.devcontainer/launch.sh isolated` launches the devcontainer without the deps/repo bind mounts, and instead contains a unique copy of the `cuspatial` source in the container's file system. +`.devcontainer/launch.sh isolated` launches the devcontainer without the deps/repo bind mounts, and instead contains a unique copy of the `cuspatial` source in a Docker [volume](https://docs.docker.com/storage/volumes/). Use this mode to launch multiple isolated development containers that can be checked out to separate branches of `cuspatial`. -**Be sure to push any commits you want to persist. Once this container is removed, any unpushed changes will be lost!** +The Docker volume persists after the devcontainer is removed, ensuring you don't pending lose work by accidentally removing the devcontainer. + +However, you will need to manually remove the volume once you've committed and pushed your changes: + +* Use the [`docker volume ls`](https://docs.docker.com/engine/reference/commandline/volume_ls/) command to list all volumes +* Use [`docker volume rm`](https://docs.docker.com/engine/reference/commandline/volume_rm/) or [`docker volume prune`](https://docs.docker.com/engine/reference/commandline/volume_prune/) to clean up unused volumes + +Alternatively, use the "Dev Volumes" tab of the VSCode Dev Containers extension to view and remove unused volumes. Examples: ```bash diff --git a/.devcontainer/conda/devcontainer.json b/.devcontainer/conda/devcontainer.json deleted file mode 100644 index 902276b71..000000000 --- a/.devcontainer/conda/devcontainer.json +++ /dev/null @@ -1,63 +0,0 @@ -{ - "shutdownAction": "stopContainer", - - "build": { - "context": "${localWorkspaceFolder}/.devcontainer", - "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", - "args": { - "CUDA": "11.8", - "LLVM": "16", - "PYTHON_PACKAGE_MANAGER": "conda" - } - }, - "hostRequirements": { - "gpu": true - }, - - "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} - }, - - "overrideFeatureInstallOrder": [ - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" - ], - - "initializeCommand": [ - "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}/single}" - ], - "updateContentCommand": ["rapids-make-vscode-workspace", "--update"], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-conda-env"], - - "containerEnv": { - "DEFAULT_CONDA_ENV": "rapids" - }, - - "workspaceFolder": "/home/coder", - "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", - "mounts": [ - "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent" - ], - - "customizations": { - "vscode": { - "extensions": [ - "mutantdino.resourcemonitor", - "tamasfe.even-better-toml" - ], - "settings": { - "files.trimFinalNewlines": true, - "files.insertFinalNewline": true, - "files.trimTrailingWhitespace": true, - "files.watcherExclude": { - "**/target/**": true - }, - "python.linting.flake8Enabled": true - } - } - } -} diff --git a/.devcontainer/conda/isolated/.devcontainer/devcontainer.json b/.devcontainer/conda/isolated/.devcontainer/devcontainer.json index aeebb0e7c..e55b955e4 100644 --- a/.devcontainer/conda/isolated/.devcontainer/devcontainer.json +++ b/.devcontainer/conda/isolated/.devcontainer/devcontainer.json @@ -5,7 +5,7 @@ "context": "${localWorkspaceFolder}/.devcontainer", "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", "args": { - "CUDA": "11.8", + "CUDA": "12.0", "LLVM": "16", "PYTHON_PACKAGE_MANAGER": "conda" } @@ -15,7 +15,7 @@ }, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} }, "overrideFeatureInstallOrder": [ @@ -27,10 +27,8 @@ ], "updateContentCommand": [ - "/bin/bash", "-c", "cp -ar /workspaces/${localWorkspaceFolderBasename} /home/coder/${localWorkspaceFolderBasename} && rapids-make-vscode-workspace --update" + "/bin/bash", "-c", "sudo chown -R $(id -u):$(id -g) /home/coder/${localWorkspaceFolderBasename} && mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml && cp -ar /workspaces/${localWorkspaceFolderBasename} /home/coder/ && rapids-make-vscode-workspace --update" ], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-conda-env"], "containerEnv": { "DEFAULT_CONDA_ENV": "rapids" @@ -42,23 +40,53 @@ "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", + "source=${localWorkspaceFolderBasename}-conda-isolated-${devcontainerId},target=/home/coder/${localWorkspaceFolderBasename},type=volume" ], "customizations": { "vscode": { "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", "tamasfe.even-better-toml" ], "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", "files.trimFinalNewlines": true, "files.insertFinalNewline": true, "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, "**/target/**": true }, - "python.linting.flake8Enabled": true + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } } } } diff --git a/.devcontainer/conda/single/.devcontainer/devcontainer.json b/.devcontainer/conda/single/.devcontainer/devcontainer.json deleted file mode 120000 index cec0647b2..000000000 --- a/.devcontainer/conda/single/.devcontainer/devcontainer.json +++ /dev/null @@ -1 +0,0 @@ -../../devcontainer.json \ No newline at end of file diff --git a/.devcontainer/conda/single/.devcontainer/devcontainer.json b/.devcontainer/conda/single/.devcontainer/devcontainer.json new file mode 100644 index 000000000..7b0865cf1 --- /dev/null +++ b/.devcontainer/conda/single/.devcontainer/devcontainer.json @@ -0,0 +1,95 @@ +{ + "shutdownAction": "stopContainer", + + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "12.0", + "LLVM": "16", + "PYTHON_PACKAGE_MANAGER": "conda" + } + }, + "hostRequirements": { + "gpu": true + }, + + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} + }, + + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + + "initializeCommand": [ + "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}/single}" + ], + + "updateContentCommand": [ + "/bin/bash", + "-c", + "mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml" + ], + + "containerEnv": { + "DEFAULT_CONDA_ENV": "rapids" + }, + + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/single,target=/home/coder/.conda/envs,type=bind,consistency=consistent" + ], + + "customizations": { + "vscode": { + "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", + "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", + "tamasfe.even-better-toml" + ], + "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", + "files.trimFinalNewlines": true, + "files.insertFinalNewline": true, + "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, + "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, + "**/target/**": true + }, + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } + } + } + } +} diff --git a/.devcontainer/conda/unified/.devcontainer/devcontainer.json b/.devcontainer/conda/unified/.devcontainer/devcontainer.json index 68ca35f5f..da741fc06 100644 --- a/.devcontainer/conda/unified/.devcontainer/devcontainer.json +++ b/.devcontainer/conda/unified/.devcontainer/devcontainer.json @@ -5,7 +5,7 @@ "context": "${localWorkspaceFolder}/.devcontainer", "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", "args": { - "CUDA": "11.8", + "CUDA": "12.0", "LLVM": "16", "PYTHON_PACKAGE_MANAGER": "conda" } @@ -15,7 +15,7 @@ }, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} }, "overrideFeatureInstallOrder": [ @@ -23,11 +23,14 @@ ], "initializeCommand": [ - "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}/unified}" + "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/unified} ${localWorkspaceFolder}/../{rmm,kvikio,cudf,raft,cumlprims_mg,cuml,cugraph-ops,cugraph,cuspatial}" + ], + + "updateContentCommand": [ + "/bin/bash", + "-c", + "mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml" ], - "updateContentCommand": ["rapids-make-vscode-workspace", "--update"], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-conda-env"], "containerEnv": { "DEFAULT_CONDA_ENV": "rapids" @@ -37,28 +40,63 @@ "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", "mounts": [ "source=${localWorkspaceFolder}/../rmm,target=/home/coder/rmm,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../kvikio,target=/home/coder/kvikio,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../cudf,target=/home/coder/cudf,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../raft,target=/home/coder/raft,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cumlprims_mg,target=/home/coder/cumlprims_mg,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cuml,target=/home/coder/cuml,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cugraph-ops,target=/home/coder/cugraph-ops,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cugraph,target=/home/coder/cugraph,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.conda/pkgs,target=/home/coder/.conda/pkgs,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.conda/${localWorkspaceFolderBasename}/unified,target=/home/coder/.conda/envs,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.conda/unified,target=/home/coder/.conda/envs,type=bind,consistency=consistent" ], "customizations": { "vscode": { "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", "tamasfe.even-better-toml" ], "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", "files.trimFinalNewlines": true, "files.insertFinalNewline": true, "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, "**/target/**": true }, - "python.linting.flake8Enabled": true + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } } } } diff --git a/.devcontainer/devcontainer.json b/.devcontainer/devcontainer.json index 902276b71..841df8838 100644 --- a/.devcontainer/devcontainer.json +++ b/.devcontainer/devcontainer.json @@ -5,7 +5,7 @@ "context": "${localWorkspaceFolder}/.devcontainer", "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", "args": { - "CUDA": "11.8", + "CUDA": "12.0", "LLVM": "16", "PYTHON_PACKAGE_MANAGER": "conda" } @@ -15,7 +15,7 @@ }, "features": { - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} }, "overrideFeatureInstallOrder": [ @@ -25,9 +25,6 @@ "initializeCommand": [ "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config,conda/pkgs,conda/${localWorkspaceFolderBasename}/single}" ], - "updateContentCommand": ["rapids-make-vscode-workspace", "--update"], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-conda-env"], "containerEnv": { "DEFAULT_CONDA_ENV": "rapids" @@ -46,17 +43,46 @@ "customizations": { "vscode": { "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", "tamasfe.even-better-toml" ], "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", "files.trimFinalNewlines": true, "files.insertFinalNewline": true, "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, "**/target/**": true }, - "python.linting.flake8Enabled": true + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } } } } diff --git a/.devcontainer/pip/devcontainer.json b/.devcontainer/pip/devcontainer.json deleted file mode 100644 index 9c51daab1..000000000 --- a/.devcontainer/pip/devcontainer.json +++ /dev/null @@ -1,66 +0,0 @@ -{ - "shutdownAction": "stopContainer", - - "build": { - "context": "${localWorkspaceFolder}/.devcontainer", - "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", - "args": { - "CUDA": "11.8", - "LLVM": "16", - "PYTHON_PACKAGE_MANAGER": "pip" - } - }, - "hostRequirements": { - "gpu": true - }, - - "features": { - "ghcr.io/devcontainers/features/python:1": {}, - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} - }, - - "overrideFeatureInstallOrder": [ - "ghcr.io/devcontainers/features/python", - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" - ], - - "initializeCommand": [ - "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/${localWorkspaceFolderBasename}/single}" - ], - "updateContentCommand": ["rapids-make-vscode-workspace", "--update"], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-pip-env"], - - "containerEnv": { - "PYTHONSAFEPATH": "true", - "PYTHONUNBUFFERED": "true", - "DEFAULT_VIRTUAL_ENV": "rapids" - }, - - "workspaceFolder": "/home/coder", - "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", - "mounts": [ - "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.local/${localWorkspaceFolderBasename}/single,target=/home/coder/.local,type=bind,consistency=consistent" - ], - - "customizations": { - "vscode": { - "extensions": [ - "mutantdino.resourcemonitor", - "tamasfe.even-better-toml" - ], - "settings": { - "files.trimFinalNewlines": true, - "files.insertFinalNewline": true, - "files.trimTrailingWhitespace": true, - "files.watcherExclude": { - "**/target/**": true - }, - "python.linting.flake8Enabled": true - } - } - } -} diff --git a/.devcontainer/pip/isolated/.devcontainer/devcontainer.json b/.devcontainer/pip/isolated/.devcontainer/devcontainer.json index c398cab0e..f1b0e9d6a 100644 --- a/.devcontainer/pip/isolated/.devcontainer/devcontainer.json +++ b/.devcontainer/pip/isolated/.devcontainer/devcontainer.json @@ -5,7 +5,7 @@ "context": "${localWorkspaceFolder}/.devcontainer", "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", "args": { - "CUDA": "11.8", + "CUDA": "12.0", "LLVM": "16", "PYTHON_PACKAGE_MANAGER": "pip" } @@ -15,12 +15,10 @@ }, "features": { - "ghcr.io/devcontainers/features/python:1": {}, - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} }, "overrideFeatureInstallOrder": [ - "ghcr.io/devcontainers/features/python", "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" ], @@ -29,10 +27,8 @@ ], "updateContentCommand": [ - "/bin/bash", "-c", "cp -ar /workspaces/${localWorkspaceFolderBasename} /home/coder/${localWorkspaceFolderBasename} && rapids-make-vscode-workspace --update" + "/bin/bash", "-c", "sudo chown -R $(id -u):$(id -g) /home/coder/${localWorkspaceFolderBasename} && mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml && cp -ar /workspaces/${localWorkspaceFolderBasename} /home/coder/ && rapids-make-vscode-workspace --update" ], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-pip-env"], "containerEnv": { "DEFAULT_VIRTUAL_ENV": "rapids" @@ -43,23 +39,53 @@ "mounts": [ "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolderBasename}-pip-isolated-${devcontainerId},target=/home/coder/${localWorkspaceFolderBasename},type=volume" ], "customizations": { "vscode": { "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", "tamasfe.even-better-toml" ], "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", "files.trimFinalNewlines": true, "files.insertFinalNewline": true, "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, "**/target/**": true }, - "python.linting.flake8Enabled": true + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } } } } diff --git a/.devcontainer/pip/single/.devcontainer/devcontainer.json b/.devcontainer/pip/single/.devcontainer/devcontainer.json deleted file mode 120000 index cec0647b2..000000000 --- a/.devcontainer/pip/single/.devcontainer/devcontainer.json +++ /dev/null @@ -1 +0,0 @@ -../../devcontainer.json \ No newline at end of file diff --git a/.devcontainer/pip/single/.devcontainer/devcontainer.json b/.devcontainer/pip/single/.devcontainer/devcontainer.json new file mode 100644 index 000000000..a93626950 --- /dev/null +++ b/.devcontainer/pip/single/.devcontainer/devcontainer.json @@ -0,0 +1,96 @@ +{ + "shutdownAction": "stopContainer", + + "build": { + "context": "${localWorkspaceFolder}/.devcontainer", + "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", + "args": { + "CUDA": "12.0", + "LLVM": "16", + "PYTHON_PACKAGE_MANAGER": "pip" + } + }, + "hostRequirements": { + "gpu": true + }, + + "features": { + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} + }, + + "overrideFeatureInstallOrder": [ + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" + ], + + "initializeCommand": [ + "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/${localWorkspaceFolderBasename}/single}" + ], + + "updateContentCommand": [ + "/bin/bash", + "-c", + "mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml" + ], + + "containerEnv": { + "PYTHONSAFEPATH": "true", + "PYTHONUNBUFFERED": "true", + "DEFAULT_VIRTUAL_ENV": "rapids" + }, + + "workspaceFolder": "/home/coder", + "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", + "mounts": [ + "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../.local/${localWorkspaceFolderBasename}/single,target=/home/coder/.local,type=bind,consistency=consistent" + ], + + "customizations": { + "vscode": { + "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", + "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", + "tamasfe.even-better-toml" + ], + "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", + "files.trimFinalNewlines": true, + "files.insertFinalNewline": true, + "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, + "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, + "**/target/**": true + }, + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } + } + } + } +} diff --git a/.devcontainer/pip/unified/.devcontainer/devcontainer.json b/.devcontainer/pip/unified/.devcontainer/devcontainer.json index 6a7acb47d..dff5e2d7c 100644 --- a/.devcontainer/pip/unified/.devcontainer/devcontainer.json +++ b/.devcontainer/pip/unified/.devcontainer/devcontainer.json @@ -5,7 +5,7 @@ "context": "${localWorkspaceFolder}/.devcontainer", "dockerfile": "${localWorkspaceFolder}/.devcontainer/Dockerfile", "args": { - "CUDA": "11.8", + "CUDA": "12.0", "LLVM": "16", "PYTHON_PACKAGE_MANAGER": "pip" } @@ -15,21 +15,22 @@ }, "features": { - "ghcr.io/devcontainers/features/python:1": {}, - "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.6": {} + "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils:23.8": {} }, "overrideFeatureInstallOrder": [ - "ghcr.io/devcontainers/features/python", "ghcr.io/rapidsai/devcontainers/features/rapids-build-utils" ], "initializeCommand": [ - "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/${localWorkspaceFolderBasename}/unified}" + "/bin/bash", "-c", "mkdir -m 0755 -p ${localWorkspaceFolder}/../.{aws,cache,config/pip,local/unified} ${localWorkspaceFolder}/../{rmm,kvikio,cudf,raft,cumlprims_mg,cuml,cugraph-ops,cugraph,cuspatial}" + ], + + "updateContentCommand": [ + "/bin/bash", + "-c", + "mkdir -m 0755 -p ~/.config/clangd && cp -n /etc/skel/.config/clangd/config.yaml ~/.config/clangd/config.yaml" ], - "updateContentCommand": ["rapids-make-vscode-workspace", "--update"], - "postCreateCommand": ["rapids-make-vscode-workspace", "--update"], - "postAttachCommand": ["rapids-make-pip-env"], "containerEnv": { "DEFAULT_VIRTUAL_ENV": "rapids" @@ -39,27 +40,62 @@ "workspaceMount": "source=${localWorkspaceFolder},target=/home/coder/${localWorkspaceFolderBasename},type=bind,consistency=consistent", "mounts": [ "source=${localWorkspaceFolder}/../rmm,target=/home/coder/rmm,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../kvikio,target=/home/coder/kvikio,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../cudf,target=/home/coder/cudf,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../raft,target=/home/coder/raft,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cumlprims_mg,target=/home/coder/cumlprims_mg,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cuml,target=/home/coder/cuml,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cugraph-ops,target=/home/coder/cugraph-ops,type=bind,consistency=consistent", + "source=${localWorkspaceFolder}/../cugraph,target=/home/coder/cugraph,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.aws,target=/home/coder/.aws,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.cache,target=/home/coder/.cache,type=bind,consistency=consistent", "source=${localWorkspaceFolder}/../.config,target=/home/coder/.config,type=bind,consistency=consistent", - "source=${localWorkspaceFolder}/../.local/${localWorkspaceFolderBasename}/unified,target=/home/coder/.local,type=bind,consistency=consistent" + "source=${localWorkspaceFolder}/../.local/unified,target=/home/coder/.local,type=bind,consistency=consistent" ], "customizations": { "vscode": { "extensions": [ + "llvm-vs-code-extensions.vscode-clangd", "mutantdino.resourcemonitor", + "ms-vscode.cpptools", + "nvidia.nsight-vscode-edition", + "seaube.clangformat", "tamasfe.even-better-toml" ], "settings": { + "C_Cpp.vcpkg.enabled": false, + "C_Cpp.formatting": "disabled", + "C_Cpp.autocomplete": "disabled", + "C_Cpp.errorSquiggles": "disabled", + "C_Cpp.intelliSenseEngine": "disabled", + "C_Cpp.configurationWarnings": "disabled", + "C_Cpp.autoAddFileAssociations": false, + "clang-format.fallbackStyle": "none", "files.trimFinalNewlines": true, "files.insertFinalNewline": true, "files.trimTrailingWhitespace": true, + "files.associations": { + "*.cu": "cuda-cpp", + "*.cuh": "cuda-cpp", + "**/libcudacxx/include/**/*": "cpp", + "**/libcudacxx-src/include/**/*": "cpp" + }, "files.watcherExclude": { + "**/build/**": true, + "**/_skbuild/**": true, "**/target/**": true }, - "python.linting.flake8Enabled": true + "python.linting.flake8Enabled": true, + "[c]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + }, + "[cuda-cpp]": { + "editor.defaultFormatter": "seaube.clangformat" + } } } } diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index a872e289f..ebe3619b1 100644 --- a/.github/workflows/build.yaml +++ b/.github/workflows/build.yaml @@ -68,19 +68,17 @@ jobs: sha: ${{ inputs.sha }} wheel-build: secrets: inherit - uses: rapidsai/shared-action-workflows/.github/workflows/wheels-manylinux-build.yml@branch-23.08 + uses: rapidsai/shared-action-workflows/.github/workflows/wheels-build.yaml@branch-23.08 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} sha: ${{ inputs.sha }} date: ${{ inputs.date }} - package-name: cuspatial - package-dir: python/cuspatial - skbuild-configure-options: "-DCUSPATIAL_BUILD_WHEELS=ON" + script: ci/build_wheel.sh wheel-publish: needs: wheel-build secrets: inherit - uses: rapidsai/shared-action-workflows/.github/workflows/wheels-manylinux-publish.yml@branch-23.08 + uses: rapidsai/shared-action-workflows/.github/workflows/wheels-publish.yaml@branch-23.08 with: build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 53cc04bd2..c5300ccd5 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -75,20 +75,14 @@ jobs: wheel-build: needs: checks secrets: inherit - uses: rapidsai/shared-action-workflows/.github/workflows/wheels-manylinux-build.yml@branch-23.08 + uses: rapidsai/shared-action-workflows/.github/workflows/wheels-build.yaml@branch-23.08 with: build_type: pull-request - package-dir: python/cuspatial - package-name: cuspatial - skbuild-configure-options: "-DCUSPATIAL_BUILD_WHEELS=ON" + script: ci/build_wheel.sh wheel-tests: needs: wheel-build secrets: inherit - uses: rapidsai/shared-action-workflows/.github/workflows/wheels-manylinux-test.yml@branch-23.08 + uses: rapidsai/shared-action-workflows/.github/workflows/wheels-test.yaml@branch-23.08 with: build_type: pull-request - package-name: cuspatial - test-smoketest: "python ./ci/wheel_smoke_test.py" - test-unittest: "python -m pytest -n 8 ./python/cuspatial/cuspatial/tests" - test-before-amd64: "apt update && DEBIAN_FRONTEND=noninteractive apt install -y --no-install-recommends libgdal-dev && python -m pip install --no-binary fiona 'fiona>=1.8.19,<1.9'" - test-before-arm64: "apt update && DEBIAN_FRONTEND=noninteractive apt install -y --no-install-recommends libgdal-dev && python -m pip install --no-binary fiona 'fiona>=1.8.19,<1.9'" + script: ci/test_wheel.sh diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 00a4c0446..5c058f38c 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -32,13 +32,10 @@ jobs: sha: ${{ inputs.sha }} wheel-tests: secrets: inherit - uses: rapidsai/shared-action-workflows/.github/workflows/wheels-manylinux-test.yml@branch-23.08 + uses: rapidsai/shared-action-workflows/.github/workflows/wheels-test.yaml@branch-23.08 with: build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} sha: ${{ inputs.sha }} - package-name: cuspatial - test-unittest: "python -m pytest -n 8 ./python/cuspatial/cuspatial/tests" - test-before-amd64: "apt update && DEBIAN_FRONTEND=noninteractive apt install -y --no-install-recommends libgdal-dev && python -m pip install --no-binary fiona 'fiona>=1.8.19,<1.9'" - test-before-arm64: "apt update && DEBIAN_FRONTEND=noninteractive apt install -y --no-install-recommends libgdal-dev && python -m pip install --no-binary fiona 'fiona>=1.8.19,<1.9'" + script: ci/test_wheel.sh diff --git a/ci/build_wheel.sh b/ci/build_wheel.sh new file mode 100755 index 000000000..34924f48b --- /dev/null +++ b/ci/build_wheel.sh @@ -0,0 +1,26 @@ +#!/bin/bash +# Copyright (c) 2023, NVIDIA CORPORATION. + +set -euo pipefail + +source rapids-configure-sccache +source rapids-date-string + +# Use gha-tools rapids-pip-wheel-version to generate wheel version then +# update the necessary files +version_override="$(rapids-pip-wheel-version ${RAPIDS_DATE_STRING})" + +RAPIDS_PY_CUDA_SUFFIX="$(rapids-wheel-ctk-name-gen ${RAPIDS_CUDA_VERSION})" + +ci/release/apply_wheel_modifications.sh ${version_override} "-${RAPIDS_PY_CUDA_SUFFIX}" +echo "The package name and/or version was modified in the package source. The git diff is:" +git diff + +cd python/cuspatial + +SKBUILD_CONFIGURE_OPTIONS="-DCUSPATIAL_BUILD_WHEELS=ON" python -m pip wheel . -w dist -vvv --no-deps --disable-pip-version-check + +mkdir -p final_dist +python -m auditwheel repair -w final_dist dist/* + +RAPIDS_PY_WHEEL_NAME="cuspatial_${RAPIDS_PY_CUDA_SUFFIX}" rapids-upload-wheels-to-s3 final_dist diff --git a/ci/test_wheel.sh b/ci/test_wheel.sh new file mode 100755 index 000000000..8efc79545 --- /dev/null +++ b/ci/test_wheel.sh @@ -0,0 +1,22 @@ +#!/bin/bash +# Copyright (c) 2023, NVIDIA CORPORATION. + +set -eou pipefail + +mkdir -p ./dist +RAPIDS_PY_CUDA_SUFFIX="$(rapids-wheel-ctk-name-gen ${RAPIDS_CUDA_VERSION})" +RAPIDS_PY_WHEEL_NAME="cuspatial_${RAPIDS_PY_CUDA_SUFFIX}" rapids-download-wheels-from-s3 ./dist + +# Install additional dependencies +apt update +DEBIAN_FRONTEND=noninteractive apt install -y --no-install-recommends libgdal-dev +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.py +else + python -m pytest -n 8 ./python/cuspatial/cuspatial/tests +fi diff --git a/conda/environments/all_cuda-120_arch-x86_64.yaml b/conda/environments/all_cuda-120_arch-x86_64.yaml index b385a9e14..a31779877 100644 --- a/conda/environments/all_cuda-120_arch-x86_64.yaml +++ b/conda/environments/all_cuda-120_arch-x86_64.yaml @@ -9,6 +9,7 @@ dependencies: - c-compiler - cmake>=3.26.4 - cuda-cudart-dev +- cuda-cupti-dev - cuda-nvcc - cuda-nvrtc-dev - cuda-version=12.0 diff --git a/cpp/cuproj/CMakeLists.txt b/cpp/cuproj/CMakeLists.txt index 57be97a19..585ffe0aa 100644 --- a/cpp/cuproj/CMakeLists.txt +++ b/cpp/cuproj/CMakeLists.txt @@ -175,13 +175,8 @@ endif() if(CUPROJ_BUILD_BENCHMARKS) # Find or install GoogleBench - CPMFindPackage(NAME benchmark - VERSION 1.5.3 - GIT_REPOSITORY https://github.com/google/benchmark.git - GIT_TAG v1.5.3 - GIT_SHALLOW TRUE - OPTIONS "BENCHMARK_ENABLE_TESTING OFF" - "BENCHMARK_ENABLE_INSTALL OFF") + include(${rapids-cmake-dir}/cpm/gbench.cmake) + rapids_cpm_gbench() # 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/benchmarks/CMakeLists.txt b/cpp/cuproj/benchmarks/CMakeLists.txt index fd8412f60..7c1f7146d 100644 --- a/cpp/cuproj/benchmarks/CMakeLists.txt +++ b/cpp/cuproj/benchmarks/CMakeLists.txt @@ -24,32 +24,31 @@ add_library(cuproj_benchmark_common OBJECT target_compile_features(cuproj_benchmark_common PUBLIC cxx_std_17 cuda_std_17) target_link_libraries(cuproj_benchmark_common - PUBLIC benchmark::benchmark - cudf::cudftestutil - cuproj) + PUBLIC benchmark::benchmark cudf::cudftestutil cuproj) target_compile_options(cuproj_benchmark_common - PUBLIC "$<$:${CUPROJ_CXX_FLAGS}>" - "$<$:${CUPROJ_CUDA_FLAGS}>") + PUBLIC "$<$:${CUPROJ_CXX_FLAGS}>" + "$<$:${CUPROJ_CUDA_FLAGS}>") target_include_directories(cuproj_benchmark_common - PUBLIC "$" - "$" - "$") + PUBLIC "$" + "$" + "$" + "$") function(ConfigureBench CMAKE_BENCH_NAME) - add_executable(${CMAKE_BENCH_NAME} ${ARGN}) - set_target_properties(${CMAKE_BENCH_NAME} - PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" - INSTALL_RPATH "\$ORIGIN/../../../lib" - ) - target_link_libraries(${CMAKE_BENCH_NAME} PRIVATE benchmark::benchmark_main cuproj_benchmark_common) - install( - TARGETS ${CMAKE_BENCH_NAME} - COMPONENT benchmark - DESTINATION bin/benchmarks/libcuproj - EXCLUDE_FROM_ALL + add_executable(${CMAKE_BENCH_NAME} ${ARGN}) + set_target_properties(${CMAKE_BENCH_NAME} + PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" + INSTALL_RPATH "\$ORIGIN/../../../lib" ) + target_link_libraries(${CMAKE_BENCH_NAME} PRIVATE benchmark::benchmark_main cuproj_benchmark_common PROJ::proj) + install( + TARGETS ${CMAKE_BENCH_NAME} + COMPONENT benchmark + DESTINATION bin/benchmarks/libcuproj + EXCLUDE_FROM_ALL + ) endfunction() # This function takes in a benchmark name and benchmark source for nvbench benchmarks and handles @@ -61,8 +60,20 @@ function(ConfigureNVBench CMAKE_BENCH_NAME) PROPERTIES RUNTIME_OUTPUT_DIRECTORY "$" INSTALL_RPATH "\$ORIGIN/../../../lib" ) + target_compile_features(${CMAKE_BENCH_NAME} PUBLIC cxx_std_17 cuda_std_17) + + target_compile_options(${CMAKE_BENCH_NAME} + PUBLIC "$<$:${CUPROJ_CXX_FLAGS}>" + "$<$:${CUPROJ_CUDA_FLAGS}>") + + target_include_directories(${CMAKE_BENCH_NAME} + PUBLIC "$" + "$" + "$" + "$") + target_link_libraries( - ${CMAKE_BENCH_NAME} PRIVATE cuproj_benchmark_common nvbench::main + ${CMAKE_BENCH_NAME} PRIVATE cuproj nvbench::main ) install( TARGETS ${CMAKE_BENCH_NAME} @@ -76,5 +87,4 @@ endfunction() ### benchmark sources ############################################################################# ################################################################################################### -ConfigureBench(TEST_BENCH - test.cu) +ConfigureBench(WGS_TO_UTM_BENCH wgs_to_utm_bench.cu) diff --git a/cpp/cuproj/benchmarks/fixture/benchmark_fixture.hpp b/cpp/cuproj/benchmarks/fixture/benchmark_fixture.hpp new file mode 100644 index 000000000..08e5dff51 --- /dev/null +++ b/cpp/cuproj/benchmarks/fixture/benchmark_fixture.hpp @@ -0,0 +1,91 @@ +/* + * Copyright (c) 2019-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. + */ + +#include +#include +#include +#include +#include + +namespace cuspatial { + +namespace { +// memory resource factory helpers +inline auto make_cuda() { return std::make_shared(); } + +inline auto make_pool() +{ + return rmm::mr::make_owning_wrapper(make_cuda()); +} +} // namespace + +/** + * @brief Google Benchmark fixture for libcuspatial benchmarks + * + * libcuspatial benchmarks should use a fixture derived from this fixture class to + * ensure that the RAPIDS Memory Manager pool mode is used in benchmarks, which + * eliminates memory allocation / deallocation performance overhead from the + * benchmark. + * + * The SetUp and TearDown methods of this fixture initialize RMM into pool mode + * and finalize it, respectively. These methods are called automatically by + * Google Benchmark + * + * Example: + * + * template + * class my_benchmark : public cuspatial::benchmark { + * public: + * using TypeParam = T; + * }; + * + * Then: + * + * BENCHMARK_TEMPLATE_DEFINE_F(my_benchmark, my_test_name, int) + * (::benchmark::State& state) { + * for (auto _ : state) { + * // benchmark stuff + * } + * } + * + * BENCHMARK_REGISTER_F(my_benchmark, my_test_name)->Range(128, 512); + */ +class benchmark : public ::benchmark::Fixture { + public: + virtual void SetUp(const ::benchmark::State& state) override + { + mr = make_pool(); + rmm::mr::set_current_device_resource(mr.get()); // set default resource to pool + } + + virtual void TearDown(const ::benchmark::State& state) override + { + // reset default resource to the initial resource + rmm::mr::set_current_device_resource(nullptr); + mr.reset(); + } + + // eliminate partial override warnings (see benchmark/benchmark.h) + void SetUp(::benchmark::State& st) override { SetUp(const_cast(st)); } + void TearDown(::benchmark::State& st) override + { + TearDown(const_cast(st)); + } + + std::shared_ptr mr; +}; + +}; // namespace cuspatial diff --git a/cpp/cuproj/benchmarks/test.cu b/cpp/cuproj/benchmarks/test.cu deleted file mode 100644 index e69de29bb..000000000 diff --git a/cpp/cuproj/benchmarks/wgs_to_utm_bench.cu b/cpp/cuproj/benchmarks/wgs_to_utm_bench.cu new file mode 100644 index 000000000..faefac005 --- /dev/null +++ b/cpp/cuproj/benchmarks/wgs_to_utm_bench.cu @@ -0,0 +1,135 @@ +/* + * 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. + */ + +#include +#include + +#include +#include + +#include +#include + +#include +#include + +#include + +#include + +template +using coordinate = typename cuspatial::vec_2d; + +static char const* epsg_src = "EPSG:4326"; +static char const* epsg_dst = "EPSG:32756"; + +template +auto make_input(std::size_t grid_side) +{ + // Sydney Harbour + coordinate min_corner{-33.9, 151.2}; + coordinate max_corner{-33.7, 151.3}; + + auto input = cuproj_test::make_grid_array, rmm::device_vector>>( + min_corner, max_corner, grid_side, grid_side); + + return input; +} + +template +static void cuproj_wgs_to_utm_benchmark(benchmark::State& state) +{ + auto const num_points = state.range(0); + + auto const grid_side{static_cast(sqrt(num_points))}; + + auto input = make_input(grid_side); + + rmm::device_vector> output(input.size()); + + auto proj = cuproj::make_projection>(epsg_src, epsg_dst); + + for (auto _ : state) { + cuda_event_timer raii(state, true); + proj.transform(input.begin(), + input.end(), + output.begin(), + cuproj::direction::FORWARD, + rmm::cuda_stream_default); + } + + state.SetItemsProcessed(num_points * state.iterations()); +} + +void proj_wgs_to_utm_benchmark(benchmark::State& state) +{ + using T = double; + auto const num_points = state.range(0); + + auto const grid_side{static_cast(sqrt(num_points))}; + + auto d_input = make_input(grid_side); + auto input = thrust::host_vector>(d_input); + + std::vector pj_input(input.size()); + + PJ_CONTEXT* C = proj_context_create(); + PJ* P = proj_create_crs_to_crs(C, epsg_src, epsg_dst, nullptr); + + for (auto _ : state) { + state.PauseTiming(); + cuproj_test::convert_coordinates(input, pj_input); + state.ResumeTiming(); + proj_trans_array(P, PJ_FWD, pj_input.size(), pj_input.data()); + } + + state.SetItemsProcessed(num_points * state.iterations()); +} + +class proj_utm_benchmark : public ::benchmark::Fixture {}; + +// Edit these for GPUs/CPUs with larger or smaller memory. +// For double precision, its' 16 bytes per (x,y) point, x2 for input and output +// 10^8 points -> 3.2GB+, 10^9 points -> 32GB+ +// H100 80GB is plenty for 10^9 points + +constexpr int range_min = 100; +constexpr int range_max = 100'000'000; + +BENCHMARK_DEFINE_F(proj_utm_benchmark, forward_double)(::benchmark::State& state) +{ + proj_wgs_to_utm_benchmark(state); +} +BENCHMARK_REGISTER_F(proj_utm_benchmark, forward_double) + ->RangeMultiplier(10) + ->Range(range_min, range_max) + ->Unit(benchmark::kMillisecond); + +class cuproj_utm_benchmark : public cuspatial::benchmark {}; + +#define UTM_CUPROJ_BENCHMARK_DEFINE(name, type) \ + BENCHMARK_DEFINE_F(cuproj_utm_benchmark, name)(::benchmark::State & state) \ + { \ + cuproj_wgs_to_utm_benchmark(state); \ + } \ + BENCHMARK_REGISTER_F(cuproj_utm_benchmark, name) \ + ->RangeMultiplier(10) \ + ->Range(range_min, range_max) \ + ->UseManualTime() \ + ->Unit(benchmark::kMillisecond); + +UTM_CUPROJ_BENCHMARK_DEFINE(forward_float, float); +UTM_CUPROJ_BENCHMARK_DEFINE(forward_double, double); diff --git a/cpp/cuproj/include/cuproj_test/convert_coordinates.hpp b/cpp/cuproj/include/cuproj_test/convert_coordinates.hpp new file mode 100644 index 000000000..f4bce42fb --- /dev/null +++ b/cpp/cuproj/include/cuproj_test/convert_coordinates.hpp @@ -0,0 +1,49 @@ +/* + * 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. + */ + +#pragma once + +#include + +#include +#include + +namespace cuproj_test { + +// Convert coordinates from a x-y struct to a PJ_COORD struct or vice versa +template +void convert_coordinates(InVector const& in, OutVector& out) +{ + using in_coord_type = typename InVector::value_type; + using out_coord_type = typename OutVector::value_type; + + static_assert( + (std::is_same_v != std::is_same_v), + "Invalid coordinate vector conversion"); + + if constexpr (std::is_same_v) { + using T = typename out_coord_type::value_type; + auto proj_coord_to_coordinate = [](auto const& c) { + return out_coord_type{static_cast(c.xy.x), static_cast(c.xy.y)}; + }; + std::transform(in.begin(), in.end(), out.begin(), proj_coord_to_coordinate); + } else if constexpr (std::is_same_v) { + auto coordinate_to_proj_coord = [](auto const& c) { return PJ_COORD{c.x, c.y, 0, 0}; }; + std::transform(in.begin(), in.end(), out.begin(), coordinate_to_proj_coord); + } +} + +} // namespace cuproj_test diff --git a/cpp/cuproj/include/cuproj_test/coordinate_generator.cuh b/cpp/cuproj/include/cuproj_test/coordinate_generator.cuh index a6569bd28..e892a9d70 100644 --- a/cpp/cuproj/include/cuproj_test/coordinate_generator.cuh +++ b/cpp/cuproj/include/cuproj_test/coordinate_generator.cuh @@ -48,6 +48,7 @@ struct grid_generator { } }; +// Create a Vector containing a grid of coordinates between the min and max corners template auto make_grid_array(Coord const& min_corner, Coord const& max_corner, diff --git a/cpp/cuproj/tests/wgs_to_utm_test.cu b/cpp/cuproj/tests/wgs_to_utm_test.cu index 4df935fe6..6d1751a28 100644 --- a/cpp/cuproj/tests/wgs_to_utm_test.cu +++ b/cpp/cuproj/tests/wgs_to_utm_test.cu @@ -14,6 +14,7 @@ * limitations under the License. */ +#include #include #include @@ -47,28 +48,6 @@ TYPED_TEST_CASE(ProjectionTest, TestTypes); template using coordinate = typename cuspatial::vec_2d; -template -void convert_coordinates(InVector const& in, OutVector& out) -{ - using in_coord_type = typename InVector::value_type; - using out_coord_type = typename OutVector::value_type; - - static_assert( - (std::is_same_v != std::is_same_v), - "Invalid coordinate vector conversion"); - - if constexpr (std::is_same_v) { - using T = typename out_coord_type::value_type; - auto proj_coord_to_coordinate = [](auto const& c) { - return out_coord_type{static_cast(c.xy.x), static_cast(c.xy.y)}; - }; - thrust::transform(in.begin(), in.end(), out.begin(), proj_coord_to_coordinate); - } else if constexpr (std::is_same_v) { - auto coordinate_to_proj_coord = [](auto const& c) { return PJ_COORD{c.x, c.y, 0, 0}; }; - thrust::transform(in.begin(), in.end(), out.begin(), coordinate_to_proj_coord); - } -} - // run a test using the cuproj library template void run_cuproj_test(thrust::host_vector> const& input, @@ -119,7 +98,7 @@ void run_forward_and_inverse(DeviceVector const& input, // projection, and we want to test both directions for each. thrust::host_vector> h_input(input.begin(), input.end()); thrust::host_vector pj_input{input.size()}; - convert_coordinates(h_input, pj_input); + cuproj_test::convert_coordinates(h_input, pj_input); thrust::host_vector pj_expected(pj_input); char const* epsg_src = "EPSG:4326"; @@ -131,7 +110,7 @@ void run_forward_and_inverse(DeviceVector const& input, run_proj_test(pj_expected, epsg_src, epsg_dst); thrust::host_vector> h_expected{pj_expected.size()}; - convert_coordinates(pj_expected, h_expected); + cuproj_test::convert_coordinates(pj_expected, h_expected); auto proj = cuproj::make_projection>(epsg_src, epsg_dst); @@ -143,7 +122,7 @@ void run_forward_and_inverse(DeviceVector const& input, run(); // invert construction pj_input = pj_expected; - convert_coordinates(pj_input, h_input); + cuproj_test::convert_coordinates(pj_input, h_input); std::swap(epsg_src, epsg_dst); run(); } diff --git a/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh b/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh index 54ad6b499..9e8d6583c 100644 --- a/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh +++ b/cpp/include/cuspatial/detail/algorithm/is_point_in_polygon.cuh @@ -68,8 +68,8 @@ __device__ inline bool is_point_in_polygon(vec_2d const& test_point, PolygonR T run_to_point = test_point.x - a.x; // point-on-edge test - bool is_colinear = float_equal(run * rise_to_point, run_to_point * rise); - if (is_colinear) { + bool is_collinear = float_equal(run * rise_to_point, run_to_point * rise); + if (is_collinear) { T minx = a.x; T maxx = b.x; if (minx > maxx) thrust::swap(minx, maxx); diff --git a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh index 0142f680e..5694088ad 100644 --- a/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh +++ b/cpp/include/cuspatial/detail/find/find_and_combine_segment.cuh @@ -19,12 +19,14 @@ #include #include #include +#include #include #include #include +#include #include namespace cuspatial { @@ -32,7 +34,18 @@ namespace detail { /** * @internal - * @brief Kernel to merge segments, naive n^2 algorithm. + * @pre All segments in range @p segments , it is presorted using `segment_comparator`. + * + * @brief Kernel to merge segments. Each thread works on one segment space, within each space, + * `segment_comparator` guarantees that segments with same slope are grouped together. We call + * each of such group is a "mergeable group". Within each mergeable group, the first segment is + * the "leading segment". The algorithm behave as follows: + * + * 1. For each mergeable group, loop over the rest of the segments in the group and see if it is + * mergeable with the leading segment. If it is, overwrite the leading segment with the merged + * result. Then mark the segment as merged by setting the flag to 1. This makes sure the inner loop + * for each merged segment is not run again. + * 2. Repeat 1 until all mergeable group is processed. */ template void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, @@ -45,6 +58,9 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, merged_flag[i] = 0; } + // For each of the segment, loop over the rest of the segment in the space and see + // if it is mergeable with the current segment. + // Note that if the current segment is already merged. Skip checking. for (auto i = offsets[pair_idx]; i < offsets[pair_idx + 1] && merged_flag[i] != 1; i++) { for (auto j = i + 1; j < offsets[pair_idx + 1]; j++) { auto res = maybe_merge_segments(segments[i], segments[j]); @@ -58,10 +74,44 @@ void __global__ simple_find_and_combine_segments_kernel(OffsetRange offsets, } } +/** + * @brief Comparator for sorting the segment range. + * + * This comparator makes sure that the segment range is sorted such that: + * 1. Segments with the same space id are grouped together. + * 2. Segments within the same space are grouped by their slope. + * 3. Within each slope group, segments are sorted by their lower left point. + */ +template +struct segment_comparator { + bool __device__ operator()(thrust::tuple> const& lhs, + thrust::tuple> const& rhs) const + { + auto lhs_index = thrust::get<0>(lhs); + auto rhs_index = thrust::get<0>(rhs); + auto lhs_segment = thrust::get<1>(lhs); + auto rhs_segment = thrust::get<1>(rhs); + + // Compare space id + if (lhs_index == rhs_index) { + // Compare slope + if (lhs_segment.collinear(rhs_segment)) { + // Sort by the lower left point of the segment. + return lhs_segment.lower_left() < rhs_segment.lower_left(); + } + return lhs_segment.slope() < rhs_segment.slope(); + } + return lhs_index < rhs_index; + } +}; + /** * @internal * @brief For each pair of mergeable segment, overwrites the first segment with merged result, * sets the flag for the second segment as 1. + * + * @note This function will alter the input segment range by rearranging the order of the segments + * within each space so that merging kernel can take place. */ template void find_and_combine_segment(OffsetRange offsets, @@ -69,9 +119,21 @@ void find_and_combine_segment(OffsetRange offsets, OutputIt merged_flag, rmm::cuda_stream_view stream) { + using index_t = typename OffsetRange::value_type; + using T = typename SegmentRange::value_type::value_type; auto num_spaces = offsets.size() - 1; if (num_spaces == 0) return; + // Construct a key iterator using the offsets of the segment and the segment itself. + auto space_id_iter = make_geometry_id_iterator(offsets.begin(), offsets.end()); + auto space_id_segment_iter = thrust::make_zip_iterator(space_id_iter, segments.begin()); + + thrust::sort_by_key(rmm::exec_policy(stream), + space_id_segment_iter, + space_id_segment_iter + segments.size(), + segments.begin(), + segment_comparator{}); + auto [threads_per_block, num_blocks] = grid_1d(num_spaces); simple_find_and_combine_segments_kernel<<>>( offsets, segments, merged_flag); diff --git a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh index 882ae22ae..66ad89003 100644 --- a/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry/polygon_ref.cuh @@ -59,13 +59,13 @@ CUSPATIAL_HOST_DEVICE auto polygon_ref::ring_end() co template CUSPATIAL_HOST_DEVICE auto polygon_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *_ring_begin); } template CUSPATIAL_HOST_DEVICE auto polygon_ref::point_end() const { - return _point_end; + return thrust::next(_point_begin, *thrust::prev(_ring_end)); } template diff --git a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh index 7a76f92b8..15472485f 100644 --- a/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry_collection/multilinestring_ref.cuh @@ -74,13 +74,15 @@ CUSPATIAL_HOST_DEVICE auto multilinestring_ref::part_ template CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *_part_begin); } template CUSPATIAL_HOST_DEVICE auto multilinestring_ref::point_end() const { - return _point_end; + // _part_end refers to the one past the last part index to the points of this multilinestring. + // So prior to computing the end point index, we need to decrement _part_end. + return thrust::next(_point_begin, *thrust::prev(_part_end)); } template diff --git a/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh b/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh index a663a2dde..e13b786ba 100644 --- a/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh +++ b/cpp/include/cuspatial/detail/geometry_collection/multipolygon_ref.cuh @@ -112,14 +112,14 @@ template CUSPATIAL_HOST_DEVICE auto multipolygon_ref::point_begin() const { - return _point_begin; + return thrust::next(_point_begin, *thrust::next(_ring_begin, *_part_begin)); } template CUSPATIAL_HOST_DEVICE auto multipolygon_ref::point_end() const { - return _point_end; + return thrust::next(_point_begin, *thrust::next(_ring_begin, *thrust::prev(_part_end))); } template diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh index 796365bde..2db5492ed 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh @@ -258,7 +258,12 @@ linestring_intersection_result pairwise_linestring_intersection( stream); points.remove_if(range(point_flags.begin(), point_flags.end()), stream); + + rmm::device_uvector point_flags_int(point_flags.size(), stream); + thrust::copy( + rmm::exec_policy(stream), point_flags.begin(), point_flags.end(), point_flags_int.begin()); } + // Phase 4: Assemble results as union column auto num_union_column_rows = points.geoms->size() + segments.geoms->size(); auto geometry_collection_offsets = 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 941efd426..3279df2f9 100644 --- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh +++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh @@ -50,6 +50,16 @@ namespace detail { namespace intersection_functors { +/** + * @brief Cast `uint8_t` to `X` + * + * @tparam X The type to cast to + */ +template +struct uchar_to_x { + X __device__ operator()(uint8_t c) { return static_cast(c); } +}; + /** @brief Functor to compute the updated offset buffer after `remove_if` operation. * * Given the `i`th row in the `geometry_collection_offset`, find the number of all @@ -292,11 +302,13 @@ struct linestring_intersection_intermediates { rmm::device_uvector reduced_flags(num_pairs(), stream); auto keys_begin = make_geometry_id_iterator(offsets->begin(), offsets->end()); + auto iflags = + thrust::make_transform_iterator(flags.begin(), intersection_functors::uchar_to_x{}); auto [keys_end, flags_end] = thrust::reduce_by_key(rmm::exec_policy(stream), keys_begin, keys_begin + flags.size(), - flags.begin(), + iflags, reduced_keys.begin(), reduced_flags.begin(), thrust::equal_to(), diff --git a/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh b/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh index fc768ec1c..614a3f2d4 100644 --- a/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh +++ b/cpp/include/cuspatial/detail/pairwise_multipoint_equals_count.cuh @@ -97,12 +97,14 @@ OutputIt pairwise_multipoint_equals_count(MultiPointRangeA lhs, rhs.offsets_begin(), rhs.offsets_end(), rhs_point_sorted.begin(), rhs_point_sorted.end()}; detail::zero_data_async(output, output + lhs.size(), stream); - auto [tpb, n_blocks] = grid_1d(lhs.num_points()); - detail::pairwise_multipoint_equals_count_kernel<<>>( - lhs, rhs_sorted, output); - CUSPATIAL_CHECK_CUDA(stream.value()); + if (lhs.num_points() > 0) { + auto [tpb, n_blocks] = grid_1d(lhs.num_points()); + detail::pairwise_multipoint_equals_count_kernel<<>>( + lhs, rhs_sorted, output); + CUSPATIAL_CHECK_CUDA(stream.value()); + } return output + lhs.size(); } diff --git a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh index a73d423ac..b9b53bfc0 100644 --- a/cpp/include/cuspatial/detail/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/detail/range/multilinestring_range.cuh @@ -200,12 +200,11 @@ multilinestring_range::is_valid_seg template template -CUSPATIAL_HOST_DEVICE thrust::pair< - vec_2d::element_t>, - vec_2d::element_t>> +CUSPATIAL_HOST_DEVICE auto multilinestring_range::segment(IndexType segment_idx) { - return thrust::make_pair(_point_begin[segment_idx], _point_begin[segment_idx + 1]); + using T = iterator_vec_base_type; + return cuspatial::segment{_point_begin[segment_idx], _point_begin[segment_idx + 1]}; } template diff --git a/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh b/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh index 028756347..b684146d3 100644 --- a/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh +++ b/cpp/include/cuspatial/detail/range/multilinestring_segment_range.cuh @@ -200,7 +200,7 @@ class multilinestring_segment_range { CUSPATIAL_HOST_DEVICE auto multigeometry_offset_begin() { return thrust::make_permutation_iterator(_per_linestring_offset_begin(), - _parent.geometry_offsets_begin()); + _parent.geometry_offset_begin()); } /// Returns end iterator to the range of the starting segment index per multilinestring diff --git a/cpp/include/cuspatial/detail/range/multipolygon_range.cuh b/cpp/include/cuspatial/detail/range/multipolygon_range.cuh index a76803229..1e126fb89 100644 --- a/cpp/include/cuspatial/detail/range/multipolygon_range.cuh +++ b/cpp/include/cuspatial/detail/range/multipolygon_range.cuh @@ -241,23 +241,6 @@ multipolygon_range:: thrust::prev(thrust::upper_bound(thrust::seq, _geometry_begin, _geometry_end, part_idx))); } -template -template -CUSPATIAL_HOST_DEVICE auto -multipolygon_range:: - geometry_idx_from_segment_idx(IndexType segment_idx) -{ - auto ring_idx = ring_idx_from_point_idx(segment_idx); - if (!is_valid_segment_id(segment_idx, ring_idx)) - return multipolygon_range:: - INVALID_INDEX; - - return geometry_idx_from_part_idx(part_idx_from_ring_idx(ring_idx)); -} - template ::o return multipolygon_begin()[multipolygon_idx]; } -template -template -CUSPATIAL_HOST_DEVICE auto -multipolygon_range::get_segment( - IndexType segment_idx) -{ - return segment{_point_begin[segment_idx], _point_begin[segment_idx + 1]}; -} - template -template -CUSPATIAL_HOST_DEVICE bool -multipolygon_range:: - is_first_point_of_multipolygon(IndexType1 point_idx, IndexType2 geometry_idx) -{ - return point_idx == _ring_begin[_part_begin[_geometry_begin[geometry_idx]]]; -} - template const& * @brief Computes shortest distance between two segments (ab and cd) that don't intersect. */ template -__forceinline__ T __device__ segment_distance_no_intersect_or_colinear(vec_2d const& a, - vec_2d const& b, - vec_2d const& c, - vec_2d const& d) +__forceinline__ T __device__ segment_distance_no_intersect_or_collinear(vec_2d const& a, + vec_2d const& b, + vec_2d const& c, + vec_2d const& d) { auto dist_sqr = min( min(point_to_segment_distance_squared(a, c, d), point_to_segment_distance_squared(b, c, d)), @@ -123,7 +123,7 @@ __forceinline__ T __device__ squared_segment_distance(vec_2d const& a, if (float_equal(denom, T{0})) { // Segments parallel or collinear - return segment_distance_no_intersect_or_colinear(a, b, c, d); + return segment_distance_no_intersect_or_collinear(a, b, c, d); } auto ac = c - a; @@ -132,7 +132,7 @@ __forceinline__ T __device__ squared_segment_distance(vec_2d const& a, auto r = r_numer * denom_reciprocal; auto s = det(ac, ab) * denom_reciprocal; if (r >= 0 and r <= 1 and s >= 0 and s <= 1) { return 0.0; } - return segment_distance_no_intersect_or_colinear(a, b, c, d); + return segment_distance_no_intersect_or_collinear(a, b, c, d); } /** diff --git a/cpp/include/cuspatial/detail/utility/zero_data.cuh b/cpp/include/cuspatial/detail/utility/zero_data.cuh index c6196a4c2..1a4f72a54 100644 --- a/cpp/include/cuspatial/detail/utility/zero_data.cuh +++ b/cpp/include/cuspatial/detail/utility/zero_data.cuh @@ -19,6 +19,8 @@ #include +#include + namespace cuspatial { namespace detail { diff --git a/cpp/include/cuspatial/geometry/segment.cuh b/cpp/include/cuspatial/geometry/segment.cuh index 9cfaf3b8e..bc2ce86e8 100644 --- a/cpp/include/cuspatial/geometry/segment.cuh +++ b/cpp/include/cuspatial/geometry/segment.cuh @@ -55,6 +55,20 @@ class alignas(sizeof(Vertex)) segment { /// Return the length squared of segment. T CUSPATIAL_HOST_DEVICE length2() const { return dot(v2 - v1, v2 - v1); } + /// Return slope of segment. + T CUSPATIAL_HOST_DEVICE slope() { return (v2.y - v1.y) / (v2.x - v1.x); } + + /// Return the lower left vertex of segment. + Vertex CUSPATIAL_HOST_DEVICE lower_left() { return v1 < v2 ? v1 : v2; } + + /// Returns true if two segments are on the same line + /// Test if the determinant of the matrix, of which the column vector is constructed from the + /// segments is 0. + bool CUSPATIAL_HOST_DEVICE collinear(segment const& other) + { + return (v1.x - v2.x) * (other.v1.y - other.v2.y) == (v1.y - v2.y) * (other.v1.x - other.v2.x); + } + private: friend std::ostream& operator<<(std::ostream& os, segment const& seg) { diff --git a/cpp/include/cuspatial/iterator_factory.cuh b/cpp/include/cuspatial/iterator_factory.cuh index 1f026512c..9b22ba9bc 100644 --- a/cpp/include/cuspatial/iterator_factory.cuh +++ b/cpp/include/cuspatial/iterator_factory.cuh @@ -16,6 +16,7 @@ #pragma once +#include #include #include #include @@ -424,6 +425,13 @@ auto make_geometry_id_iterator(GeometryIter geometry_offsets_begin, std::distance(geometry_offsets_begin, geometry_offsets_end))); } +template +auto make_count_iterator_from_offset_iterator(OffsetIterator it) +{ + auto zipped_offsets_it = thrust::make_zip_iterator(it, thrust::next(it)); + return thrust::make_transform_iterator(zipped_offsets_it, detail::offset_pair_to_count_functor{}); +} + /** * @} // end of doxygen group */ diff --git a/cpp/include/cuspatial/range/multilinestring_range.cuh b/cpp/include/cuspatial/range/multilinestring_range.cuh index b4bb6eaf5..b4b396ee5 100644 --- a/cpp/include/cuspatial/range/multilinestring_range.cuh +++ b/cpp/include/cuspatial/range/multilinestring_range.cuh @@ -101,6 +101,12 @@ class multilinestring_range { /// Return the iterator to the one past the last point in the range. CUSPATIAL_HOST_DEVICE auto point_end() { return _point_end; } + /// Return the iterator to the first geometry offset in the range. + CUSPATIAL_HOST_DEVICE auto geometry_offset_begin() { return _geometry_begin; } + + /// Return the iterator to the one past the last geometry offset in the range. + CUSPATIAL_HOST_DEVICE auto geometry_offset_end() { return _geometry_end; } + /// Return the iterator to the first part offset in the range. CUSPATIAL_HOST_DEVICE auto part_offset_begin() { return _part_begin; } @@ -144,8 +150,7 @@ class multilinestring_range { /// Returns the segment given a segment index. template - CUSPATIAL_HOST_DEVICE thrust::pair, vec_2d> segment( - IndexType segment_idx); + CUSPATIAL_HOST_DEVICE auto segment(IndexType segment_idx); /// Returns an iterator to the counts of points per multilinestring CUSPATIAL_HOST_DEVICE auto multilinestring_point_count_begin(); @@ -168,13 +173,6 @@ class multilinestring_range { template CUSPATIAL_HOST_DEVICE auto operator[](IndexType multilinestring_idx); - /// Raw offsets iterator - - CUSPATIAL_HOST_DEVICE auto geometry_offsets_begin() { return _geometry_begin; } - CUSPATIAL_HOST_DEVICE auto geometry_offsets_end() { return _geometry_end; } - CUSPATIAL_HOST_DEVICE auto part_offsets_begin() { return _part_begin; } - CUSPATIAL_HOST_DEVICE auto part_offsets_end() { return _part_end; } - /// Range Casts /// Casts the multilinestring range into a multipoint range. diff --git a/cpp/include/cuspatial/range/multipoint_range.cuh b/cpp/include/cuspatial/range/multipoint_range.cuh index 92f7bff76..caee16ad4 100644 --- a/cpp/include/cuspatial/range/multipoint_range.cuh +++ b/cpp/include/cuspatial/range/multipoint_range.cuh @@ -149,6 +149,7 @@ class multipoint_range { /** * @brief Returns `true` if the range contains only single points + * Undefined behavior if the range is an empty range. */ CUSPATIAL_HOST_DEVICE bool is_single_point_range(); diff --git a/cpp/include/cuspatial/range/multipolygon_range.cuh b/cpp/include/cuspatial/range/multipolygon_range.cuh index 2af70e066..1aae07346 100644 --- a/cpp/include/cuspatial/range/multipolygon_range.cuh +++ b/cpp/include/cuspatial/range/multipolygon_range.cuh @@ -72,8 +72,6 @@ class multipolygon_range { using index_t = iterator_value_type; using element_t = iterator_vec_base_type; - int64_t static constexpr INVALID_INDEX = -1; - multipolygon_range(GeometryIterator geometry_begin, GeometryIterator geometry_end, PartIterator part_begin, @@ -117,10 +115,10 @@ class multipolygon_range { CUSPATIAL_HOST_DEVICE auto point_end(); /// Return the iterator to the first geometry offset in the range. - CUSPATIAL_HOST_DEVICE auto geometry_offsets_begin() { return _part_begin; } + CUSPATIAL_HOST_DEVICE auto geometry_offset_begin() { return _part_begin; } /// Return the iterator to the one past the last geometry offset in the range. - CUSPATIAL_HOST_DEVICE auto geometry_offsets_end() { return _part_end; } + CUSPATIAL_HOST_DEVICE auto geometry_offset_end() { return _part_end; } /// Return the iterator to the first part offset in the range. CUSPATIAL_HOST_DEVICE auto part_offset_begin() { return _part_begin; } @@ -134,13 +132,6 @@ class multipolygon_range { /// Return the iterator to the one past the last ring offset in the range. CUSPATIAL_HOST_DEVICE auto ring_offset_end() { return _ring_end; } - /// Given the index of a segment, return the index of the geometry (multipolygon) that contains - /// the segment. Segment index is the index to the starting point of the segment. If the index is - /// the last point of the ring, then it is not a valid index. This function returns - /// multipolygon_range::INVALID_INDEX if the index is invalid. - template - CUSPATIAL_HOST_DEVICE auto geometry_idx_from_segment_idx(IndexType segment_idx); - /// Given the index of a point, return the index of the ring that contains the point. template CUSPATIAL_HOST_DEVICE auto ring_idx_from_point_idx(IndexType point_idx); @@ -158,16 +149,6 @@ class multipolygon_range { template CUSPATIAL_HOST_DEVICE auto operator[](IndexType multipolygon_idx); - /// Returns the `segment_idx`th segment in the multipolygon range. - template - CUSPATIAL_HOST_DEVICE auto get_segment(IndexType segment_idx); - - /// Returns `true` if `point_idx`th point is the first point of `geometry_idx`th - /// multipolygon - template - CUSPATIAL_HOST_DEVICE bool is_first_point_of_multipolygon(IndexType1 point_idx, - IndexType2 geometry_idx); - /// Returns an iterator to the number of points of the first multipolygon /// @note The count includes the duplicate first and last point of the ring. CUSPATIAL_HOST_DEVICE auto multipolygon_point_count_begin(); diff --git a/cpp/include/cuspatial_test/geometry_generator.cuh b/cpp/include/cuspatial_test/geometry_generator.cuh index 2fd3be28c..f966664fb 100644 --- a/cpp/include/cuspatial_test/geometry_generator.cuh +++ b/cpp/include/cuspatial_test/geometry_generator.cuh @@ -354,7 +354,7 @@ auto generate_multilinestring_array(multilinestring_generator_parameter param params.origin, detail::random_walk_functor{params.segment_length}); - return make_multilinestring_array( + return make_multilinestring_array( std::move(geometry_offset), std::move(part_offset), std::move(points)); } diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh index 91d2c04d7..cbf1d8f56 100644 --- a/cpp/include/cuspatial_test/vector_factories.cuh +++ b/cpp/include/cuspatial_test/vector_factories.cuh @@ -101,6 +101,17 @@ class multipolygon_array { { } + multipolygon_array(rmm::device_vector&& geometry_offsets_array, + rmm::device_vector&& part_offsets_array, + rmm::device_vector&& ring_offsets_array, + rmm::device_vector&& coordinates_array) + : _geometry_offsets_array(std::move(geometry_offsets_array)), + _part_offsets_array(std::move(part_offsets_array)), + _ring_offsets_array(std::move(ring_offsets_array)), + _coordinates_array(std::move(coordinates_array)) + { + } + multipolygon_array(rmm::device_uvector&& geometry_offsets_array, rmm::device_uvector&& part_offsets_array, rmm::device_uvector&& ring_offsets_array, @@ -230,9 +241,9 @@ class multilinestring_array { multilinestring_array(GeometryArray geometry_offsets_array, PartArray part_offsets_array, CoordinateArray coordinate_array) - : _geometry_offset_array(geometry_offsets_array), - _part_offset_array(part_offsets_array), - _coordinate_array(coordinate_array) + : _geometry_offset_array(std::move(geometry_offsets_array)), + _part_offset_array(std::move(part_offsets_array)), + _coordinate_array(std::move(coordinate_array)) { } @@ -264,29 +275,41 @@ class multilinestring_array { }; /** - * @brief Construct an owning object of a multilinestring array from ranges + * @brief Construct an owning object of a multilinestring array from `device_uvectors` * * @param geometry_inl Range of geometry offsets * @param part_inl Range of part offsets * @param coord_inl Ramge of coordinate * @return multilinestring array object */ -template -auto make_multilinestring_array(IndexRangeA geometry_inl, - IndexRangeB part_inl, - CoordRange coord_inl) +template +auto make_multilinestring_array(rmm::device_uvector&& geometry_inl, + rmm::device_uvector&& part_inl, + rmm::device_uvector>&& coord_inl) { - using CoordType = typename CoordRange::value_type; - using DeviceIndexVector = thrust::device_vector; - using DeviceCoordVector = thrust::device_vector; + return multilinestring_array, + rmm::device_uvector, + rmm::device_uvector>>( + std::move(geometry_inl), std::move(part_inl), std::move(coord_inl)); +} - return multilinestring_array( - make_device_vector(std::move(geometry_inl)), - make_device_vector(std::move(part_inl)), - make_device_vector(std::move(coord_inl))); +/** + * @brief Construct an owning object of a multilinestring array from `device_vectors` + * + * @param geometry_inl Range of geometry offsets + * @param part_inl Range of part offsets + * @param coord_inl Ramge of coordinate + * @return multilinestring array object + */ +template +auto make_multilinestring_array(rmm::device_vector&& geometry_inl, + rmm::device_vector&& part_inl, + rmm::device_vector>&& coord_inl) +{ + return multilinestring_array, + rmm::device_vector, + rmm::device_vector>>( + std::move(geometry_inl), std::move(part_inl), std::move(coord_inl)); } /** @@ -414,5 +437,17 @@ auto make_multipoint_array(rmm::device_uvector geometry_offsets, std::move(geometry_offsets), std::move(coords)}; } +/** + * @brief Factory method to construct multipoint array by moving the offsets and coordinates from + * `rmm::device_vector`. + */ +template +auto make_multipoint_array(rmm::device_vector geometry_offsets, + rmm::device_vector> coords) +{ + return multipoint_array, rmm::device_vector>>{ + std::move(geometry_offsets), std::move(coords)}; +} + } // namespace test } // namespace cuspatial diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index b67b214ce..7f7cb112a 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -171,6 +171,7 @@ ConfigureTest(SINUSOIDAL_PROJECTION_TEST_EXP # range ConfigureTest(RANGE_TEST_EXP + range/multipoint_range_test.cu range/multilinestring_range_test.cu range/multipolygon_range_test.cu) diff --git a/cpp/tests/find/find_and_combine_segments_test.cu b/cpp/tests/find/find_and_combine_segments_test.cu index ea01d7e53..f7a584f51 100644 --- a/cpp/tests/find/find_and_combine_segments_test.cu +++ b/cpp/tests/find/find_and_combine_segments_test.cu @@ -279,5 +279,67 @@ TYPED_TEST(FindAndCombineSegmentsTest, nooverlap3) CUSPATIAL_RUN_TEST(this->run_single_test, segments, {0, 0}, - {S{P{0.0, 0.0}, P{1.0, 1.0}}, S{P{0.0, 1.0}, P{1.0, 0.0}}}); + {S{P{0.0, 1.0}, P{1.0, 0.0}}, S{P{0.0, 0.0}, P{1.0, 1.0}}}); +} + +TYPED_TEST(FindAndCombineSegmentsTest, twospaces) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array({0, 2, 4}, + {S{P{0.0, 0.0}, P{1.0, 1.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{1.0, 1.0}, P{0.0, 0.0}}, + S{P{2.0, 2.0}, P{1.0, 1.0}}}); + + CUSPATIAL_RUN_TEST(this->run_single_test, + segments, + {0, 1, 0, 1}, + {S{P{0.0, 0.0}, P{2.0, 2.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{0.0, 0.0}, P{2.0, 2.0}}, + S{P{2.0, 2.0}, P{1.0, 1.0}}}); +} + +TYPED_TEST(FindAndCombineSegmentsTest, twospaces_non_contiguous_segments_with_empty) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array({0, 4, 4}, + {S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{3.0, 3.0}, P{4.0, 4.0}}, + S{P{0.0, 0.0}, P{1.0, 1.0}}, + S{P{2.0, 2.0}, P{3.0, 3.0}}}); + + CUSPATIAL_RUN_TEST(this->run_single_test, + segments, + {0, 1, 1, 1}, + {S{P{0.0, 0.0}, P{4.0, 4.0}}, + S{P{1.0, 1.0}, P{2.0, 2.0}}, + S{P{2.0, 2.0}, P{3.0, 3.0}}, + S{P{3.0, 3.0}, P{4.0, 4.0}}}); +} + +TYPED_TEST(FindAndCombineSegmentsTest, onespace_non_contiguous_segments_overlaps) +{ + using T = TypeParam; + using index_t = std::size_t; + using P = vec_2d; + using S = segment; + + auto segments = make_segment_array( + {0, 3}, + {S{P{1.0, 1.0}, P{2.0, 2.0}}, S{P{4.0, 4.0}, P{5.0, 5.0}}, S{P{-1.0, -1.0}, P{4.0, 4.0}}}); + + CUSPATIAL_RUN_TEST( + this->run_single_test, + segments, + {0, 1, 1}, + {S{P{-1.0, -1.0}, P{5.0, 5.0}}, S{P{1.0, 1.0}, P{2.0, 2.0}}, S{P{4.0, 4.0}, P{5.0, 5.0}}}); } diff --git a/cpp/tests/intersection/linestring_intersection_large_test.cu b/cpp/tests/intersection/linestring_intersection_large_test.cu index aed4c00f3..51c571392 100644 --- a/cpp/tests/intersection/linestring_intersection_large_test.cu +++ b/cpp/tests/intersection/linestring_intersection_large_test.cu @@ -27,6 +27,8 @@ #include #include +#include + #include template @@ -2025,3 +2027,32 @@ TYPED_TEST(LinestringIntersectionLargeTest, LongInput) CUSPATIAL_RUN_TEST( this->template verify_legal_result, multilinestrings1.range(), multilinestrings2.range()); } + +template +struct coordinate_functor { + cuspatial::vec_2d __device__ operator()(std::size_t i) + { + return cuspatial::vec_2d{static_cast(i), static_cast(i)}; + } +}; + +TYPED_TEST(LinestringIntersectionLargeTest, LongInput_2) +{ + using P = cuspatial::vec_2d; + auto geometry_offset = cuspatial::test::make_device_vector({0, 1}); + auto part_offset = cuspatial::test::make_device_vector({0, 130}); + auto coordinates = rmm::device_uvector

(260, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + thrust::next(coordinates.begin(), 128), + coordinate_functor{}); + + coordinates.set_element(128, P{127.0, 0.0}, this->stream()); + coordinates.set_element(129, P{0.0, 0.0}, this->stream()); + + auto rng = cuspatial::make_multilinestring_range( + 1, geometry_offset.begin(), 1, part_offset.begin(), 130, coordinates.begin()); + + CUSPATIAL_RUN_TEST(this->template verify_legal_result, rng, rng); +} diff --git a/cpp/tests/range/multilinestring_range_test.cu b/cpp/tests/range/multilinestring_range_test.cu index 250e88355..11c17e5ac 100644 --- a/cpp/tests/range/multilinestring_range_test.cu +++ b/cpp/tests/range/multilinestring_range_test.cu @@ -19,17 +19,31 @@ #include #include +#include #include #include +#include +#include +#include #include #include +#include +#include + #include +#include using namespace cuspatial; using namespace cuspatial::test; +template +void __global__ array_access_tester(MultiLineStringRange mls, std::size_t i, OutputIt output_points) +{ + thrust::copy(thrust::seq, mls[i].point_begin(), mls[i].point_end(), output_points); +} + template struct MultilinestringRangeTest : public BaseFixture { void run_segment_test_single(std::initializer_list geometry_offset, @@ -533,3 +547,807 @@ TYPED_TEST(MultilinestringRangeTest, MultilinestringAsMultipointTest6) {P{1, 1}, P{0, 0}, P{6, 6}, P{6, 7}}, {{P{1, 1}, P{0, 0}, P{6, 6}, P{6, 7}}}); } + +template +class MultilinestringRangeTestBase : public BaseFixture { + public: + struct copy_leading_point_per_multilinestring { + template + vec_2d __device__ operator()(MultiLineStringRef m) + { + return m.size() > 0 ? m[0].point_begin()[0] : vec_2d{-1, -1}; + } + }; + + template + struct part_idx_from_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t point_idx) + { + return _rng.part_idx_from_point_idx(point_idx); + } + }; + + template + struct part_idx_from_segment_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t segment_idx) + { + auto opt = _rng.part_idx_from_segment_idx(segment_idx); + if (opt.has_value()) { + return opt.value(); + } else { + return std::numeric_limits::max(); + } + } + }; + + template + struct geometry_idx_from_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t point_idx) + { + return _rng.geometry_idx_from_point_idx(point_idx); + } + }; + + template + struct intra_part_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t i) { return _rng.intra_part_idx(i); } + }; + + template + struct intra_point_idx_functor { + MultiLineStringRange _rng; + std::size_t __device__ operator()(std::size_t i) { return _rng.intra_point_idx(i); } + }; + + template + struct is_valid_segment_id_functor { + MultiLineStringRange _rng; + bool __device__ operator()(std::size_t i) + { + auto part_idx = _rng.part_idx_from_point_idx(i); + return _rng.is_valid_segment_id(i, part_idx); + } + }; + + template + struct segment_functor { + MultiLineStringRange _rng; + segment __device__ operator()(std::size_t i) + { + auto part_idx = _rng.part_idx_from_point_idx(i); + return _rng.is_valid_segment_id(i, part_idx) + ? _rng.segment(i) + : segment{vec_2d{-1, -1}, vec_2d{-1, -1}}; + } + }; + + void SetUp() { make_test_multilinestring(); } + + virtual void make_test_multilinestring() = 0; + + auto range() { return test_multilinestring->range(); } + + void run_test() + { + test_size(); + + test_num_multilinestrings(); + + test_num_linestrings(); + + test_num_points(); + + test_multilinestring_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_part_offset_it(); + + test_part_idx_from_point_idx(); + + test_part_idx_from_segment_idx(); + + test_geometry_idx_from_point_idx(); + + test_intra_part_idx(); + + test_intra_point_idx(); + + test_is_valid_segment_id(); + + test_segment(); + + test_multilinestring_point_count_it(); + + test_multilinestring_linestring_count_it(); + + test_array_access_operator(); + + test_geometry_offset_it(); + } + + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multilinestrings()); } + + virtual void test_num_multilinestrings() = 0; + + virtual void test_num_linestrings() = 0; + + virtual void test_num_points() = 0; + + virtual void test_multilinestring_it() = 0; + + void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multilinestring_begin()); } + + void test_end() { EXPECT_EQ(this->range().end(), this->range().multilinestring_end()); } + + virtual void test_point_it() = 0; + + virtual void test_geometry_offset_it() = 0; + + virtual void test_part_offset_it() = 0; + + virtual void test_part_idx_from_point_idx() = 0; + + virtual void test_part_idx_from_segment_idx() = 0; + + virtual void test_geometry_idx_from_point_idx() = 0; + + virtual void test_intra_part_idx() = 0; + + virtual void test_intra_point_idx() = 0; + + virtual void test_is_valid_segment_id() = 0; + + virtual void test_segment() = 0; + + virtual void test_multilinestring_point_count_it() = 0; + + virtual void test_multilinestring_linestring_count_it() = 0; + + virtual void test_array_access_operator() = 0; + + // Helper functions to be used by all subclass (test cases). + rmm::device_uvector> copy_leading_points() + { + auto rng = this->range(); + auto d_leading_point = rmm::device_uvector>(rng.num_multilinestrings(), stream()); + thrust::transform(rmm::exec_policy(stream()), + rng.begin(), + rng.end(), + d_leading_point.begin(), + copy_leading_point_per_multilinestring()); + return d_leading_point; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = this->range(); + auto d_all_points = rmm::device_uvector>(rng.num_points(), stream()); + thrust::copy( + rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_all_points.begin()); + return d_all_points; + } + + rmm::device_uvector copy_geometry_offsets() + { + auto rng = this->range(); + auto d_geometry_offsets = + rmm::device_uvector(rng.num_multilinestrings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.geometry_offset_begin(), + rng.geometry_offset_end(), + d_geometry_offsets.begin()); + return d_geometry_offsets; + } + + rmm::device_uvector copy_part_offset() + { + auto rng = this->range(); + auto d_part_offset = rmm::device_uvector(rng.num_linestrings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.part_offset_begin(), + rng.part_offset_end(), + d_part_offset.begin()); + return d_part_offset; + } + + rmm::device_uvector copy_part_idx_from_point_idx() + { + auto rng = this->range(); + auto d_part_idx = rmm::device_uvector(rng.num_points(), stream()); + auto f = part_idx_from_point_idx_functor{rng}; + thrust::tabulate(rmm::exec_policy(stream()), d_part_idx.begin(), d_part_idx.end(), f); + return d_part_idx; + } + + rmm::device_uvector copy_part_idx_from_segment_idx() + { + auto rng = this->range(); + auto d_part_idx = rmm::device_uvector(rng.num_points(), stream()); + auto f = part_idx_from_segment_idx_functor{rng}; + + thrust::tabulate(rmm::exec_policy(stream()), d_part_idx.begin(), d_part_idx.end(), f); + return d_part_idx; + } + + rmm::device_uvector copy_geometry_idx_from_point_idx() + { + auto rng = this->range(); + auto d_geometry_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_geometry_idx.begin(), + d_geometry_idx.end(), + geometry_idx_from_point_idx_functor{rng}); + return d_geometry_idx; + } + + rmm::device_uvector copy_intra_part_idx() + { + auto rng = this->range(); + auto d_intra_part_idx = rmm::device_uvector(rng.num_linestrings(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_intra_part_idx.begin(), + d_intra_part_idx.end(), + intra_part_idx_functor{rng}); + return d_intra_part_idx; + } + + rmm::device_uvector copy_intra_point_idx() + { + auto rng = this->range(); + auto d_intra_point_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_intra_point_idx.begin(), + d_intra_point_idx.end(), + intra_point_idx_functor{rng}); + return d_intra_point_idx; + } + + rmm::device_uvector copy_is_valid_segment_id() + { + auto rng = this->range(); + auto d_is_valid_segment_id = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_is_valid_segment_id.begin(), + d_is_valid_segment_id.end(), + is_valid_segment_id_functor{rng}); + return d_is_valid_segment_id; + } + + rmm::device_uvector> copy_segments() + { + auto rng = this->range(); + auto d_segments = rmm::device_uvector>(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_segments.begin(), + d_segments.end(), + segment_functor{rng}); + return d_segments; + } + + rmm::device_uvector copy_multilinestring_point_count() + { + auto rng = this->range(); + auto d_multilinestring_point_count = + rmm::device_uvector(rng.num_multilinestrings(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multilinestring_point_count_begin(), + rng.multilinestring_point_count_end(), + d_multilinestring_point_count.begin()); + return d_multilinestring_point_count; + } + + rmm::device_uvector copy_multilinestring_linestring_count() + { + auto rng = this->range(); + auto d_multilinestring_linestring_count = + rmm::device_uvector(rng.num_multilinestrings(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multilinestring_linestring_count_begin(), + rng.multilinestring_linestring_count_end(), + d_multilinestring_linestring_count.begin()); + return d_multilinestring_linestring_count; + } + + rmm::device_uvector> copy_all_points_of_ith_multilinestring(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(stream()); + + thrust::copy_n(rmm::exec_policy(stream()), + rng.multilinestring_point_count_begin() + i, + 1, + num_points.data()); + + auto d_all_points = rmm::device_uvector>(num_points.value(stream()), stream()); + + array_access_tester<<<1, 1, 0, stream()>>>(rng, i, d_all_points.data()); + return d_all_points; + } + + protected: + std::unique_ptr, + rmm::device_vector, + rmm::device_vector>>> + test_multilinestring; +}; + +template +class MultilinestringRangeEmptyTest : public MultilinestringRangeTestBase { + void make_test_multilinestring() + { + auto array = make_multilinestring_array({0}, {0}, {}); + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 0); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = rmm::device_vector>(0); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = rmm::device_vector(0); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + // Nothing to access + SUCCEED(); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeEmptyTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeEmptyTest, EmptyTest) { this->run_test(); } + +template +class MultilinestringRangeOneTest : public MultilinestringRangeTestBase { + void make_test_multilinestring() + { + auto array = make_multilinestring_array( + {0, 2}, {0, 2, 5}, {{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 1); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 2); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 5); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>({{10, 10}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = + make_device_vector>({{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = make_device_vector({0, 2, 5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 1, 1, 1}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = make_device_vector( + {0, std::numeric_limits::max(), 1, 1, std::numeric_limits::max()}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 0, 0, 0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = make_device_vector({0, 1}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = make_device_vector({0, 1, 0, 1, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = make_device_vector({1, 0, 1, 1, 0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = make_device_vector>({{{10, 10}, {20, 20}}, + {{-1, -1}, {-1, -1}}, + {{100, 100}, {200, 200}}, + {{200, 200}, {300, 300}}, + {{-1, -1}, {-1, -1}}}); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = make_device_vector({5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = make_device_vector({2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multilinestring(0); + auto expected = + make_device_vector>({{10, 10}, {20, 20}, {100, 100}, {200, 200}, {300, 300}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0, 2, 5}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeOneTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeOneTest, OneTest) { this->run_test(); } + +template +class MultilinestringRangeOneThousandTest : public MultilinestringRangeTestBase { + public: + struct make_points_functor { + vec_2d __device__ operator()(std::size_t i) + { + auto part_idx = i / 2; + auto intra_point_idx = i % 2; + return vec_2d{static_cast(part_idx * 10 + intra_point_idx), + static_cast(part_idx * 10 + intra_point_idx)}; + } + }; + + void make_test_multilinestring() + { + rmm::device_vector geometry_offset(1001); + rmm::device_vector part_offset(1001); + rmm::device_vector> points(2000); + + thrust::sequence( + rmm::exec_policy(this->stream()), geometry_offset.begin(), geometry_offset.end()); + + thrust::sequence( + rmm::exec_policy(this->stream()), part_offset.begin(), part_offset.end(), 0, 2); + + thrust::tabulate( + rmm::exec_policy(this->stream()), points.begin(), points.end(), make_points_functor{}); + + auto array = make_multilinestring_array( + std::move(geometry_offset), std::move(part_offset), std::move(points)); + + this->test_multilinestring = std::make_unique(std::move(array)); + } + + void test_num_multilinestrings() { EXPECT_EQ(this->range().num_multilinestrings(), 1000); } + + void test_num_linestrings() { EXPECT_EQ(this->range().num_linestrings(), 1000); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 2000); } + + void test_multilinestring_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(1000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return vec_2d{i * T{10.}, i * T{10.}}; + }); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expected = rmm::device_uvector>(2000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expected.begin(), expected.end(), make_points_functor{}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_part_offset_it() + { + auto part_offset = this->copy_part_offset(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 2); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offset, expected); + } + + void test_part_idx_from_point_idx() + { + auto part_idx = this->copy_part_idx_from_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_part_idx_from_segment_idx() + { + auto part_idx = this->copy_part_idx_from_segment_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return i % 2 == 0 ? i / 2 : std::numeric_limits::max(); + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_idx, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_idx = this->copy_geometry_idx_from_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_idx, expected); + } + + void test_intra_part_idx() + { + auto intra_part_idx = this->copy_intra_part_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + detail::zero_data_async(expected.begin(), expected.end(), this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_part_idx, expected); + } + + void test_intra_point_idx() + { + auto intra_point_idx = this->copy_intra_point_idx(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i % 2; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(intra_point_idx, expected); + } + + void test_is_valid_segment_id() + { + auto is_valid_segment_id = this->copy_is_valid_segment_id(); + auto expected = rmm::device_uvector(2000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return (i + 1) % 2; }); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(is_valid_segment_id, expected); + } + + void test_segment() + { + auto segments = this->copy_segments(); + auto expected = rmm::device_uvector>(2000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + auto part_idx = i / 2; + auto intra_point_idx = i % 2; + return i % 2 == 0 + ? segment{vec_2d{static_cast(part_idx * 10 + intra_point_idx), + static_cast(part_idx * 10 + intra_point_idx)}, + vec_2d{static_cast(part_idx * 10 + intra_point_idx + 1), + static_cast(part_idx * 10 + intra_point_idx + 1)}} + : segment{vec_2d{-1, -1}, vec_2d{-1, -1}}; + }); + CUSPATIAL_EXPECT_VEC2D_PAIRS_EQUIVALENT(segments, expected); + } + + void test_multilinestring_point_count_it() + { + auto multilinestring_point_count = this->copy_multilinestring_point_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 2); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_point_count, expected); + } + + void test_multilinestring_linestring_count_it() + { + auto multilinestring_linestring_count = this->copy_multilinestring_linestring_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 1); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multilinestring_linestring_count, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multilinestring(513); + auto expected = make_device_vector>({{5130, 5130}, {5131, 5131}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_geometry_offset_it() + { + auto geometry_offsets = this->copy_geometry_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_offsets, expected); + } + + void test_part_offsets_it() + { + auto part_offsets = this->copy_part_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 2); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(part_offsets, expected); + } +}; + +TYPED_TEST_CASE(MultilinestringRangeOneThousandTest, FloatingPointTypes); +TYPED_TEST(MultilinestringRangeOneThousandTest, OneThousandTest) { this->run_test(); } diff --git a/cpp/tests/range/multipoint_range_test.cu b/cpp/tests/range/multipoint_range_test.cu new file mode 100644 index 000000000..501ddd6cf --- /dev/null +++ b/cpp/tests/range/multipoint_range_test.cu @@ -0,0 +1,514 @@ +/* + * 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. + */ + +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include + +#include + +using namespace cuspatial; +using namespace cuspatial::test; + +template +void __global__ array_access_tester(MultiPointRange multipoints, + std::size_t i, + OutputIt output_points) +{ + using T = typename MultiPointRange::element_t; + thrust::copy(thrust::seq, multipoints[i].begin(), multipoints[i].end(), output_points); +} + +template +void __global__ point_accessor_tester(MultiPointRange multipoints, std::size_t i, OutputIt point) +{ + using T = typename MultiPointRange::element_t; + point[0] = multipoints.point(i); +} + +template +class MultipointRangeTest : public BaseFixture { + public: + struct copy_leading_point_per_multipoint { + template + vec_2d __device__ operator()(MultiPointRef multipoint) + { + return multipoint.size() > 0 ? multipoint[0] : vec_2d{-1, -1}; + } + }; + + template + struct point_idx_to_geometry_idx { + MultiPointRange rng; + + point_idx_to_geometry_idx(MultiPointRange r) : rng(r) {} + + std::size_t __device__ operator()(std::size_t pidx) + { + return rng.geometry_idx_from_point_idx(pidx); + } + }; + + void SetUp() { make_test_multipoints(); } + auto range() { return test_multipoints->range(); } + + virtual void make_test_multipoints() = 0; + + void run_test() + { + test_num_multipoints(); + + test_num_points(); + + test_size(); + + test_multipoint_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_offsets_it(); + + test_geometry_idx_from_point_idx(); + + test_subscript_operator(); + + test_point_accessor(); + + test_is_single_point_range(); + } + + virtual void test_num_multipoints() = 0; + + virtual void test_num_points() = 0; + + void test_size() { EXPECT_EQ(this->range().size(), this->range().num_multipoints()); } + + virtual void test_multipoint_it() = 0; + + void test_begin() { EXPECT_EQ(this->range().begin(), this->range().multipoint_begin()); } + + void test_end() { EXPECT_EQ(this->range().end(), this->range().multipoint_end()); } + + virtual void test_point_it() = 0; + + virtual void test_offsets_it() = 0; + + virtual void test_geometry_idx_from_point_idx() = 0; + + virtual void test_subscript_operator() = 0; + + virtual void test_point_accessor() = 0; + + virtual void test_is_single_point_range() = 0; + + protected: + rmm::device_uvector> copy_leading_points() + { + auto rng = this->range(); + rmm::device_uvector> leading_points(rng.num_multipoints(), this->stream()); + thrust::transform(rmm::exec_policy(this->stream()), + rng.multipoint_begin(), + rng.multipoint_end(), + leading_points.begin(), + copy_leading_point_per_multipoint{}); + + return leading_points; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = this->range(); + rmm::device_uvector> points(rng.num_points(), this->stream()); + thrust::copy( + rmm::exec_policy(this->stream()), rng.point_begin(), rng.point_end(), points.begin()); + return points; + }; + + rmm::device_uvector copy_offsets() + { + auto rng = this->range(); + rmm::device_uvector offsets(rng.num_multipoints() + 1, this->stream()); + thrust::copy( + rmm::exec_policy(this->stream()), rng.offsets_begin(), rng.offsets_end(), offsets.begin()); + return offsets; + }; + + rmm::device_uvector copy_geometry_idx() + { + auto rng = this->range(); + rmm::device_uvector idx(rng.num_points(), this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), idx.begin(), idx.end(), point_idx_to_geometry_idx{rng}); + return idx; + } + + rmm::device_scalar> copy_ith_point(std::size_t i) + { + auto rng = this->range(); + + rmm::device_scalar> point(this->stream()); + point_accessor_tester<<<1, 1, 0, this->stream()>>>(rng, i, point.data()); + CUSPATIAL_CHECK_CUDA(this->stream()); + + return point; + } + + rmm::device_uvector> copy_ith_multipoint(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(this->stream()); + auto count_iterator = make_count_iterator_from_offset_iterator(this->range().offsets_begin()); + thrust::copy_n(rmm::exec_policy(this->stream()), count_iterator, 1, num_points.data()); + + rmm::device_uvector> multipoint(num_points.value(this->stream()), this->stream()); + array_access_tester<<<1, 1, 0, this->stream()>>>(rng, i, multipoint.begin()); + CUSPATIAL_CHECK_CUDA(this->stream()); + + return multipoint; + } + + std::unique_ptr, rmm::device_vector>>> + test_multipoints; +}; + +template +class EmptyMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = rmm::device_uvector>(0, this->stream()); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = rmm::device_uvector(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + // Range is empty, nothing to test. + SUCCEED(); + } + + void test_point_accessor() + { + // Range is empty, nothing to test. + SUCCEED(); + } + + void test_is_single_point_range() + { + // Range is empty, undefined behavior. + SUCCEED(); + } +}; + +TYPED_TEST_CASE(EmptyMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(EmptyMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthOneMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({{{1.0, 1.0}, {10.0, 10.0}}}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 1); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 2); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>({{1.0, 1.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = make_device_vector>({{1.0, 1.0}, {10.0, 10.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0, 2}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = make_device_vector({0, 0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + auto multipoint = this->copy_ith_multipoint(0); + auto expected = make_device_vector>({{1.0, 1.0}, {10.0, 10.0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multipoint, expected); + } + + void test_point_accessor() + { + auto point = this->copy_ith_point(1); + auto expected = vec_2d{10.0, 10.0}; + EXPECT_EQ(point.value(this->stream()), expected); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthOneMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(LengthOneMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthFiveMultiPointRangeTest : public MultipointRangeTest { + public: + void make_test_multipoints() + { + auto array = make_multipoint_array({{{0.0, 0.0}, {1.0, 1.0}}, + {{10.0, 10.0}}, + {{20.0, 21.0}, {22.0, 23.0}}, + {{30.0, 31.0}, {32.0, 33.0}, {34.0, 35.0}}, + {}}); + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), 5); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 8); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expected = make_device_vector>( + {{0.0, 0.0}, {10.0, 10.0}, {20.0, 21.0}, {30.0, 31.0}, {-1, -1}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expected); + } + + void test_point_it() + { + auto points = this->copy_all_points(); + auto expected = make_device_vector>({{0.0, 0.0}, + {1.0, 1.0}, + {10.0, 10.0}, + {20.0, 21.0}, + {22.0, 23.0}, + {30.0, 31.0}, + {32.0, 33.0}, + {34.0, 35.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(points, expected); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expected = make_device_vector({0, 2, 3, 5, 8, 8}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expected); + } + + void test_geometry_idx_from_point_idx() + { + auto geometry_indices = this->copy_geometry_idx(); + auto expected = make_device_vector({0, 0, 1, 2, 2, 3, 3, 3}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(geometry_indices, expected); + } + + void test_subscript_operator() + { + auto second_multipoint = this->copy_ith_multipoint(2); + auto expected = make_device_vector>({{20.0, 21.0}, {22.0, 23.0}}); + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(second_multipoint, expected); + } + + void test_point_accessor() + { + auto third_point = this->copy_ith_point(3); + auto expected = vec_2d{20.0, 21.0}; + EXPECT_EQ(third_point.value(this->stream()), expected); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthFiveMultiPointRangeTest, FloatingPointTypes); +TYPED_TEST(LengthFiveMultiPointRangeTest, Test) { this->run_test(); } + +template +class LengthOneThousandRangeTest : public MultipointRangeTest { + public: + std::size_t static constexpr num_multipoints = 1000; + std::size_t static constexpr num_point_per_multipoint = 3; + std::size_t static constexpr num_points = num_multipoints * num_point_per_multipoint; + void make_test_multipoints() + { + rmm::device_vector geometry_offsets(num_multipoints + 1); + rmm::device_vector> coordinates(num_points); + + thrust::sequence(rmm::exec_policy(this->stream()), + geometry_offsets.begin(), + geometry_offsets.end(), + 0ul, + num_point_per_multipoint); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + coordinates.end(), + [] __device__(auto i) { + return vec_2d{static_cast(i), 10.0}; + }); + + auto array = + make_multipoint_array(std::move(geometry_offsets), std::move(coordinates)); + + this->test_multipoints = std::make_unique(std::move(array)); + } + + void test_num_multipoints() { EXPECT_EQ(this->range().num_multipoints(), num_multipoints); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), num_points); } + + void test_multipoint_it() + { + auto leading_points = this->copy_leading_points(); + auto expect = rmm::device_uvector>(num_multipoints, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return vec_2d{static_cast(i) * num_point_per_multipoint, 10.0}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(leading_points, expect); + } + + void test_point_it() + { + auto all_points = this->copy_all_points(); + auto expect = rmm::device_uvector>(num_points, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return vec_2d{static_cast(i), 10.0}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expect); + } + + void test_offsets_it() + { + auto offsets = this->copy_offsets(); + auto expect = rmm::device_uvector(num_multipoints + 1, this->stream()); + thrust::sequence(rmm::exec_policy(this->stream()), + expect.begin(), + expect.end(), + 0ul, + num_point_per_multipoint); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(offsets, expect); + } + + void test_geometry_idx_from_point_idx() + { + auto indices = this->copy_geometry_idx(); + auto expect = rmm::device_uvector(3000, this->stream()); + thrust::tabulate( + rmm::exec_policy(this->stream()), expect.begin(), expect.end(), [] __device__(auto i) { + return i / num_point_per_multipoint; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(indices, expect); + } + + void test_subscript_operator() + { + auto multipoint_five_hundred_thirty_third = this->copy_ith_multipoint(533); + auto expect = make_device_vector>({{533 * num_point_per_multipoint, 10.0}, + {533 * num_point_per_multipoint + 1, 10.0}, + {533 * num_point_per_multipoint + 2, 10.0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(multipoint_five_hundred_thirty_third, expect); + } + + void test_point_accessor() + { + auto point_seventeen_hundred_seventy_six = this->copy_ith_point(1776); + auto expect = vec_2d{1776, 10.0}; + + EXPECT_EQ(point_seventeen_hundred_seventy_six.value(this->stream()), expect); + } + + void test_is_single_point_range() { EXPECT_FALSE(this->range().is_single_point_range()); } +}; + +TYPED_TEST_CASE(LengthOneThousandRangeTest, FloatingPointTypes); +TYPED_TEST(LengthOneThousandRangeTest, Test) { this->run_test(); } diff --git a/cpp/tests/range/multipolygon_range_test.cu b/cpp/tests/range/multipolygon_range_test.cu index 2fe133945..1ba2d9935 100644 --- a/cpp/tests/range/multipolygon_range_test.cu +++ b/cpp/tests/range/multipolygon_range_test.cu @@ -22,10 +22,14 @@ #include #include +#include #include #include #include +#include +#include + #include using namespace cuspatial; @@ -121,10 +125,14 @@ struct MultipolygonRangeTest : public BaseFixture { multipolygon_coordinates); auto rng = multipolygon_array.range().as_multilinestring_range(); - auto got = - make_multilinestring_array(range(rng.geometry_offsets_begin(), rng.geometry_offsets_end()), - range(rng.part_offsets_begin(), rng.part_offsets_end()), - range(rng.point_begin(), rng.point_end())); + auto geometry_offsets = + rmm::device_vector(rng.geometry_offset_begin(), rng.geometry_offset_end()); + auto part_offsets = + rmm::device_vector(rng.part_offset_begin(), rng.part_offset_end()); + auto points = rmm::device_vector>(rng.point_begin(), rng.point_end()); + + auto got = make_multilinestring_array( + std::move(geometry_offsets), std::move(part_offsets), std::move(points)); auto expected = make_multilinestring_array( multilinestring_geometry_offset, multilinestring_part_offset, multilinestring_coordinates); @@ -482,7 +490,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonSegmentCount_ContainsEmptyPart) {6, 3}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring1) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon1) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -494,7 +502,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring1) {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring2) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon2) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -528,7 +536,7 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring2) {20, 20}}); } -TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultilinestring3) +TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultipolygon3) { CUSPATIAL_RUN_TEST(this->test_multipolygon_as_multilinestring, {0, 1, 2}, @@ -638,3 +646,677 @@ TYPED_TEST(MultipolygonRangeTest, MultipolygonAsMultiPoint3) {21, 21}, {20, 20}}); } + +template +__global__ void array_access_tester(MultiPolygonRange rng, std::size_t i, PointOutputIt output) +{ + thrust::copy(thrust::seq, rng[i].point_begin(), rng[i].point_end(), output); +} + +template +class MultipolygonRangeTestBase : public BaseFixture { + public: + struct copy_leading_point_functor { + template + __device__ vec_2d operator()(MultiPolygonRef mpolygon) + { + return mpolygon.size() > 0 ? mpolygon.point_begin()[0] : vec_2d{-1, -1}; + } + }; + + template + struct ring_idx_from_point_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t point_idx) + { + return mpolygons.ring_idx_from_point_idx(point_idx); + } + }; + + template + struct part_idx_from_ring_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t ring_idx) + { + return mpolygons.part_idx_from_ring_idx(ring_idx); + } + }; + + template + struct geometry_idx_from_part_idx_functor { + MultiPolygonRange mpolygons; + __device__ std::size_t operator()(std::size_t part_idx) + { + return mpolygons.geometry_idx_from_part_idx(part_idx); + } + }; + + void SetUp() { make_test_multipolygon(); } + + virtual void make_test_multipolygon() = 0; + + auto range() { return test_multipolygon->range(); } + + void run_test() + { + test_size(); + + test_num_multipolygons(); + + test_num_polygons(); + + test_num_rings(); + + test_num_points(); + + test_multipolygon_it(); + + test_begin(); + + test_end(); + + test_point_it(); + + test_geometry_offsets_it(); + + test_part_offset_it(); + + test_ring_offset_it(); + + test_ring_idx_from_point_idx(); + + test_part_idx_from_ring_idx(); + + test_geometry_idx_from_part_idx(); + + test_array_access_operator(); + + test_multipolygon_point_count_it(); + + test_multipolygon_ring_count_it(); + } + + void test_size() { EXPECT_EQ(range().size(), range().num_multipolygons()); }; + + virtual void test_num_multipolygons() = 0; + + virtual void test_num_polygons() = 0; + + virtual void test_num_rings() = 0; + + virtual void test_num_points() = 0; + + virtual void test_multipolygon_it() = 0; + + void test_begin() { EXPECT_EQ(range().begin(), range().multipolygon_begin()); } + + void test_end() { EXPECT_EQ(range().end(), range().multipolygon_end()); } + + virtual void test_point_it() = 0; + + virtual void test_geometry_offsets_it() = 0; + + virtual void test_part_offset_it() = 0; + + virtual void test_ring_offset_it() = 0; + + virtual void test_ring_idx_from_point_idx() = 0; + + virtual void test_part_idx_from_ring_idx() = 0; + + virtual void test_geometry_idx_from_part_idx() = 0; + + virtual void test_array_access_operator() = 0; + + virtual void test_multipolygon_point_count_it() = 0; + + virtual void test_multipolygon_ring_count_it() = 0; + + // helper method to access multipolygon range + rmm::device_uvector> copy_leading_point_multipolygon() + { + auto rng = range(); + auto d_points = rmm::device_uvector>(rng.num_multipolygons(), stream()); + thrust::transform(rmm::exec_policy(stream()), + rng.multipolygon_begin(), + rng.multipolygon_end(), + d_points.begin(), + copy_leading_point_functor{}); + return d_points; + } + + rmm::device_uvector> copy_all_points() + { + auto rng = range(); + auto d_points = rmm::device_uvector>(rng.num_points(), stream()); + thrust::copy(rmm::exec_policy(stream()), rng.point_begin(), rng.point_end(), d_points.begin()); + return d_points; + } + + rmm::device_uvector copy_geometry_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_multipolygons() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.geometry_offset_begin(), + rng.geometry_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_part_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_polygons() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.part_offset_begin(), + rng.part_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_ring_offsets() + { + auto rng = range(); + auto d_offsets = rmm::device_uvector(rng.num_rings() + 1, stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.ring_offset_begin(), + rng.ring_offset_end(), + d_offsets.begin()); + return d_offsets; + } + + rmm::device_uvector copy_ring_idx_from_point_idx() + { + auto rng = range(); + auto d_ring_idx = rmm::device_uvector(rng.num_points(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_ring_idx.begin(), + d_ring_idx.end(), + ring_idx_from_point_idx_functor{rng}); + return d_ring_idx; + } + + rmm::device_uvector copy_part_idx_from_ring_idx() + { + auto rng = range(); + auto d_part_idx = rmm::device_uvector(rng.num_rings(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_part_idx.begin(), + d_part_idx.end(), + part_idx_from_ring_idx_functor{rng}); + return d_part_idx; + } + + rmm::device_uvector copy_geometry_idx_from_part_idx() + { + auto rng = range(); + auto d_geometry_idx = rmm::device_uvector(rng.num_polygons(), stream()); + thrust::tabulate(rmm::exec_policy(stream()), + d_geometry_idx.begin(), + d_geometry_idx.end(), + geometry_idx_from_part_idx_functor{rng}); + return d_geometry_idx; + } + + rmm::device_uvector> copy_all_points_of_ith_multipolygon(std::size_t i) + { + auto rng = this->range(); + rmm::device_scalar num_points(stream()); + + thrust::copy_n( + rmm::exec_policy(stream()), rng.multipolygon_point_count_begin() + i, 1, num_points.data()); + + auto d_all_points = rmm::device_uvector>(num_points.value(stream()), stream()); + + array_access_tester<<<1, 1, 0, stream()>>>(rng, i, d_all_points.data()); + return d_all_points; + } + + rmm::device_uvector copy_multipolygon_point_count() + { + auto rng = this->range(); + auto d_point_count = rmm::device_uvector(rng.num_multipolygons(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multipolygon_point_count_begin(), + rng.multipolygon_point_count_end(), + d_point_count.begin()); + return d_point_count; + } + + rmm::device_uvector copy_multipolygon_ring_count() + { + auto rng = this->range(); + auto d_ring_count = rmm::device_uvector(rng.num_multipolygons(), stream()); + thrust::copy(rmm::exec_policy(stream()), + rng.multipolygon_ring_count_begin(), + rng.multipolygon_ring_count_end(), + d_ring_count.begin()); + return d_ring_count; + } + + protected: + std::unique_ptr, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>> + test_multipolygon; +}; + +template +class MultipolygonRangeEmptyTest : public MultipolygonRangeTestBase { + void make_test_multipolygon() + { + auto geometry_offsets = make_device_vector({0}); + auto part_offsets = make_device_vector({0}); + auto ring_offsets = make_device_vector({0}); + auto coordinates = make_device_vector>({}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 0); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 0); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 0); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 0); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + rmm::device_uvector> expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + rmm::device_uvector> expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = make_device_vector({0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = make_device_vector({}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + // Nothing to access + SUCCEED(); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + rmm::device_uvector expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + rmm::device_uvector expected(0, this->stream()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeEmptyTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeEmptyTest, EmptyMultipolygonRange) { this->run_test(); } + +template +class MultipolygonRangeOneTest : public MultipolygonRangeTestBase { + void make_test_multipolygon() + { + auto geometry_offsets = make_device_vector({0, 2}); + auto part_offsets = make_device_vector({0, 1, 2}); + auto ring_offsets = make_device_vector({0, 4, 8}); + auto coordinates = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 1); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 2); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 2); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 8); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + auto expected = make_device_vector>({{0, 0}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + auto expected = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = make_device_vector({0, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = make_device_vector({0, 1, 2}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = make_device_vector({0, 4, 8}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = make_device_vector({0, 0, 0, 0, 1, 1, 1, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = make_device_vector({0, 1}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = make_device_vector({0, 0}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multipolygon(0); + auto expected = make_device_vector>( + {{0, 0}, {1, 0}, {1, 1}, {0, 0}, {10, 10}, {11, 10}, {11, 11}, {10, 10}}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + auto expected = make_device_vector({8}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + auto expected = make_device_vector({2}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeOneTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeOneTest, OneMultipolygonRange) { this->run_test(); } + +template +class MultipolygonRangeOneThousandTest : public MultipolygonRangeTestBase { + public: + struct make_points_functor { + __device__ auto operator()(std::size_t i) + { + auto geometry_idx = i / 4; + auto intra_point_idx = i % 4; + return vec_2d{geometry_idx * T{10.} + intra_point_idx, + geometry_idx * T{10.} + intra_point_idx}; + } + }; + + void make_test_multipolygon() + { + auto geometry_offsets = rmm::device_vector(1001); + auto part_offsets = rmm::device_vector(1001); + auto ring_offsets = rmm::device_vector(1001); + auto coordinates = rmm::device_vector>(4000); + + thrust::sequence( + rmm::exec_policy(this->stream()), geometry_offsets.begin(), geometry_offsets.end()); + + thrust::sequence(rmm::exec_policy(this->stream()), part_offsets.begin(), part_offsets.end()); + + thrust::sequence( + rmm::exec_policy(this->stream()), ring_offsets.begin(), ring_offsets.end(), 0, 4); + + thrust::tabulate(rmm::exec_policy(this->stream()), + coordinates.begin(), + coordinates.end(), + make_points_functor{}); + + this->test_multipolygon = std::make_unique, + rmm::device_vector, + rmm::device_vector, + rmm::device_vector>>>( + std::move(geometry_offsets), + std::move(part_offsets), + std::move(ring_offsets), + std::move(coordinates)); + } + + void test_num_multipolygons() { EXPECT_EQ(this->range().num_multipolygons(), 1000); } + + void test_num_polygons() { EXPECT_EQ(this->range().num_polygons(), 1000); } + + void test_num_rings() { EXPECT_EQ(this->range().num_rings(), 1000); } + + void test_num_points() { EXPECT_EQ(this->range().num_points(), 4000); } + + void test_multipolygon_it() + { + rmm::device_uvector> d_points = this->copy_leading_point_multipolygon(); + rmm::device_uvector> expected(1000, this->stream()); + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { + return vec_2d{i * T{10.}, i * T{10.}}; + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_point_it() + { + rmm::device_uvector> d_points = this->copy_all_points(); + rmm::device_uvector> expected(4000, this->stream()); + + thrust::tabulate( + rmm::exec_policy(this->stream()), expected.begin(), expected.end(), make_points_functor{}); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_points, expected); + } + + void test_geometry_offsets_it() + { + rmm::device_uvector d_offsets = this->copy_geometry_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_part_offset_it() + { + rmm::device_uvector d_offsets = this->copy_part_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_offset_it() + { + rmm::device_uvector d_offsets = this->copy_ring_offsets(); + auto expected = rmm::device_uvector(1001, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 0, 4); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_offsets, expected); + } + + void test_ring_idx_from_point_idx() + { + rmm::device_uvector d_ring_idx = this->copy_ring_idx_from_point_idx(); + auto expected = rmm::device_uvector(4000, this->stream()); + + thrust::tabulate(rmm::exec_policy(this->stream()), + expected.begin(), + expected.end(), + [] __device__(std::size_t i) { return i / 4; }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_idx, expected); + } + + void test_part_idx_from_ring_idx() + { + rmm::device_uvector d_part_idx = this->copy_part_idx_from_ring_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_part_idx, expected); + } + + void test_geometry_idx_from_part_idx() + { + rmm::device_uvector d_geometry_idx = this->copy_geometry_idx_from_part_idx(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::sequence(rmm::exec_policy(this->stream()), expected.begin(), expected.end()); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_geometry_idx, expected); + } + + void test_array_access_operator() + { + auto all_points = this->copy_all_points_of_ith_multipolygon(777); + auto expected = make_device_vector>({ + {7770, 7770}, + {7771, 7771}, + {7772, 7772}, + {7773, 7773}, + }); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(all_points, expected); + } + + void test_multipolygon_point_count_it() + { + rmm::device_uvector d_point_count = this->copy_multipolygon_point_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 4); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_point_count, expected); + } + + void test_multipolygon_ring_count_it() + { + rmm::device_uvector d_ring_count = this->copy_multipolygon_ring_count(); + auto expected = rmm::device_uvector(1000, this->stream()); + + thrust::fill(rmm::exec_policy(this->stream()), expected.begin(), expected.end(), 1); + + CUSPATIAL_EXPECT_VECTORS_EQUIVALENT(d_ring_count, expected); + } +}; + +TYPED_TEST_CASE(MultipolygonRangeOneThousandTest, FloatingPointTypes); + +TYPED_TEST(MultipolygonRangeOneThousandTest, OneThousandMultipolygonRange) { this->run_test(); } diff --git a/dependencies.yaml b/dependencies.yaml index 66f19f6a1..b09ddeed5 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -176,6 +176,7 @@ dependencies: - cuda-version=12.0 - cuda-cudart-dev - cuda-nvrtc-dev + - cuda-cupti-dev - matrix: cuda: "11.8" packages: diff --git a/fetch_rapids.cmake b/fetch_rapids.cmake index 99da99888..6de15d0e6 100644 --- a/fetch_rapids.cmake +++ b/fetch_rapids.cmake @@ -11,7 +11,9 @@ # or implied. See the License for the specific language governing permissions and limitations under # the License. # ============================================================================= -file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-23.08/RAPIDS.cmake - ${CMAKE_BINARY_DIR}/RAPIDS.cmake -) -include(${CMAKE_BINARY_DIR}/RAPIDS.cmake) +if(NOT EXISTS ${CMAKE_CURRENT_BINARY_DIR}/CUSPATIAL_RAPIDS.cmake) + file(DOWNLOAD https://raw.githubusercontent.com/rapidsai/rapids-cmake/branch-23.08/RAPIDS.cmake + ${CMAKE_BINARY_DIR}/CUSPATIAL_RAPIDS.cmake + ) +endif() +include(${CMAKE_BINARY_DIR}/CUSPATIAL_RAPIDS.cmake) diff --git a/python/cuspatial/cuspatial/core/binops/distance_dispatch.py b/python/cuspatial/cuspatial/core/binops/distance_dispatch.py new file mode 100644 index 000000000..15e3aaa81 --- /dev/null +++ b/python/cuspatial/cuspatial/core/binops/distance_dispatch.py @@ -0,0 +1,202 @@ +import cudf +from cudf.core.column import arange, full + +from cuspatial._lib.distance import ( + pairwise_linestring_distance, + pairwise_linestring_polygon_distance, + pairwise_point_distance, + pairwise_point_linestring_distance, + pairwise_point_polygon_distance, + pairwise_polygon_distance, +) +from cuspatial._lib.types import CollectionType +from cuspatial.core._column.geometa import Feature_Enum +from cuspatial.utils.column_utils import ( + contains_only_linestrings, + contains_only_multipoints, + contains_only_points, + contains_only_polygons, +) + +# Maps from type combinations to a tuple of (function, reverse, +# point_collection_types). +# +# If reverse is True, the arguments need to be swapped. +# Due to the way the functions are written, certain combinations of types +# requires that the arguments be swapped. For example, +# `point_linestring_distance` requires that the first argument be a point and +# the second argument be a linestring. In this case, when lhs is a linestring +# and rhs is a point, the arguments need to be swapped. The results holds true +# thanks to that cartesian distance is symmetric. +# +# `point_collection_types` is a tuple of the types of the point column type. +# For example, if the first argument is a `MultiPoint` and the second is a +# `Point`, then the `point_collection_types` is (`CollectionType.MULTI`, +# `CollectionType.SINGLE`). They are only needed for point/multipoint columns, +# because the cython APIs are designed to handle both point and multipoint +# columns based on their collection types. + +type_to_func = { + (Feature_Enum.POINT, Feature_Enum.POINT): ( + pairwise_point_distance, + False, + (CollectionType.SINGLE, CollectionType.SINGLE), + ), + (Feature_Enum.POINT, Feature_Enum.MULTIPOINT): ( + pairwise_point_distance, + False, + (CollectionType.SINGLE, CollectionType.MULTI), + ), + (Feature_Enum.POINT, Feature_Enum.LINESTRING): ( + pairwise_point_linestring_distance, + False, + (CollectionType.SINGLE,), + ), + (Feature_Enum.POINT, Feature_Enum.POLYGON): ( + pairwise_point_polygon_distance, + False, + (CollectionType.SINGLE,), + ), + (Feature_Enum.LINESTRING, Feature_Enum.POINT): ( + pairwise_point_linestring_distance, + True, + (CollectionType.SINGLE,), + ), + (Feature_Enum.LINESTRING, Feature_Enum.MULTIPOINT): ( + pairwise_point_linestring_distance, + True, + (CollectionType.MULTI,), + ), + (Feature_Enum.LINESTRING, Feature_Enum.LINESTRING): ( + pairwise_linestring_distance, + False, + (), + ), + (Feature_Enum.LINESTRING, Feature_Enum.POLYGON): ( + pairwise_linestring_polygon_distance, + False, + (), + ), + (Feature_Enum.POLYGON, Feature_Enum.POINT): ( + pairwise_point_polygon_distance, + True, + (CollectionType.SINGLE,), + ), + (Feature_Enum.POLYGON, Feature_Enum.MULTIPOINT): ( + pairwise_point_polygon_distance, + True, + (CollectionType.MULTI,), + ), + (Feature_Enum.POLYGON, Feature_Enum.LINESTRING): ( + pairwise_linestring_polygon_distance, + True, + (), + ), + (Feature_Enum.POLYGON, Feature_Enum.POLYGON): ( + pairwise_polygon_distance, + False, + (), + ), + (Feature_Enum.MULTIPOINT, Feature_Enum.POINT): ( + pairwise_point_distance, + False, + (CollectionType.MULTI, CollectionType.SINGLE), + ), + (Feature_Enum.MULTIPOINT, Feature_Enum.MULTIPOINT): ( + pairwise_point_distance, + False, + (CollectionType.MULTI, CollectionType.MULTI), + ), + (Feature_Enum.MULTIPOINT, Feature_Enum.LINESTRING): ( + pairwise_point_linestring_distance, + False, + (CollectionType.MULTI,), + ), + (Feature_Enum.MULTIPOINT, Feature_Enum.POLYGON): ( + pairwise_point_polygon_distance, + False, + (CollectionType.MULTI,), + ), +} + + +class DistanceDispatch: + """Dispatches distance operations between two GeoSeries""" + + def __init__(self, lhs, rhs, align): + if align: + self._lhs, self._rhs = lhs.align(rhs) + else: + self._lhs, self._rhs = lhs, rhs + + self._align = align + self._res_index = lhs.index + self._non_null_mask = self._lhs.notna() & self._rhs.notna() + self._lhs = self._lhs[self._non_null_mask] + self._rhs = self._rhs[self._non_null_mask] + + # TODO: This test is expensive, so would be nice if we can cache it + self._lhs_type = self._determine_series_type(self._lhs) + self._rhs_type = self._determine_series_type(self._rhs) + + def _determine_series_type(self, s): + """Check single geometry type of `s`.""" + if contains_only_multipoints(s): + typ = Feature_Enum.MULTIPOINT + elif contains_only_points(s): + typ = Feature_Enum.POINT + elif contains_only_linestrings(s): + typ = Feature_Enum.LINESTRING + elif contains_only_polygons(s): + typ = Feature_Enum.POLYGON + else: + raise NotImplementedError( + "Geoseries with mixed geometry types are not supported" + ) + return typ + + def _column(self, s, typ): + """Get column of `s` based on `typ`.""" + if typ == Feature_Enum.POINT: + return s.points.column() + elif typ == Feature_Enum.MULTIPOINT: + return s.multipoints.column() + elif typ == Feature_Enum.LINESTRING: + return s.lines.column() + elif typ == Feature_Enum.POLYGON: + return s.polygons.column() + + @property + def _lhs_column(self): + return self._column(self._lhs, self._lhs_type) + + @property + def _rhs_column(self): + return self._column(self._rhs, self._rhs_type) + + def __call__(self): + func, reverse, collection_types = type_to_func[ + (self._lhs_type, self._rhs_type) + ] + if reverse: + dist = func(*collection_types, self._rhs_column, self._lhs_column) + else: + dist = func(*collection_types, self._lhs_column, self._rhs_column) + + # Rows with misaligned indices contains nan. Here we scatter the + # distance values to the correct indices. + result = full( + len(self._res_index), + float("nan"), + dtype="float64", + ) + scatter_map = arange( + len(self._res_index), dtype="int32" + ).apply_boolean_mask(self._non_null_mask) + + result[scatter_map] = dist + + # If `align==False`, geopandas preserves lhs index. + index = None if self._align else self._res_index + + return cudf.Series(result, index=index, nan_as_null=False) diff --git a/python/cuspatial/cuspatial/core/binpreds/feature_contains.py b/python/cuspatial/cuspatial/core/binpreds/feature_contains.py index 04636fa19..e551313ec 100644 --- a/python/cuspatial/cuspatial/core/binpreds/feature_contains.py +++ b/python/cuspatial/cuspatial/core/binpreds/feature_contains.py @@ -25,10 +25,9 @@ MultiPoint, Point, Polygon, - _false_series, - _linestrings_to_center_point, _open_polygon_rings, _pli_lines_to_multipoints, + _pli_points_to_multipoints, _points_and_lines_to_multipoints, _zero_series, ) @@ -51,7 +50,23 @@ def _preprocess(self, lhs, rhs): preprocessor_result = super()._preprocess_multipoint_rhs(lhs, rhs) return self._compute_predicate(lhs, rhs, preprocessor_result) - def _intersection_results_for_contains(self, lhs, rhs): + def _intersection_results_for_contains_linestring(self, lhs, rhs): + pli = _basic_intersects_pli(lhs, rhs) + + # Convert the pli into points and multipoint intersections. + multipoint_points = _pli_points_to_multipoints(pli) + multipoint_lines = _pli_lines_to_multipoints(pli) + + # Count the point intersections that are equal to points in the + # LineString + # Count the linestring intersections that are equal to points in + # the LineString + return ( + _basic_equals_count(rhs, multipoint_points), + _basic_equals_count(rhs, multipoint_lines), + ) + + def _intersection_results_for_contains_polygon(self, lhs, rhs): pli = _basic_intersects_pli(lhs, rhs) pli_features = pli[1] if len(pli_features) == 0: @@ -64,19 +79,15 @@ def _intersection_results_for_contains(self, lhs, rhs): pli_features, pli_offsets ) - # A point in the rhs can be one of three possible states: - # 1. It is in the interior of the lhs - # 2. It is in the exterior of the lhs - # 3. It is on the boundary of the lhs - # This function tests if the point in the rhs is in the boundary - # of the lhs intersect_equals_count = _basic_equals_count(rhs, multipoints) return intersect_equals_count def _compute_polygon_polygon_contains(self, lhs, rhs, preprocessor_result): lines_rhs = _open_polygon_rings(rhs) contains = _basic_contains_count(lhs, lines_rhs).reset_index(drop=True) - intersects = self._intersection_results_for_contains(lhs, lines_rhs) + intersects = self._intersection_results_for_contains_polygon( + lhs, lines_rhs + ) # A closed polygon has an extra line segment that is not used in # counting the number of points. We need to subtract this from the # number of points in the polygon. @@ -89,47 +100,35 @@ def _compute_polygon_polygon_contains(self, lhs, rhs, preprocessor_result): result = contains + intersects >= rhs.sizes - polygon_size_reduction return result - def _test_interior(self, lhs, rhs): - # We only need to test linestrings that are length 2. - # Divide the linestring in half and test the point for containment - # in the polygon. - size_two = rhs.sizes == 2 - if (size_two).any(): - center_points = _linestrings_to_center_point(rhs[size_two]) - size_two_results = _false_series(len(lhs)) - size_two_results.iloc[rhs.index[size_two]] = ( - _basic_contains_count(lhs, center_points) > 0 - ) - return size_two_results - else: - return _false_series(len(lhs)) - def _compute_polygon_linestring_contains( self, lhs, rhs, preprocessor_result ): + # Count the number of points in lhs that are properly contained by + # rhs contains = _basic_contains_count(lhs, rhs).reset_index(drop=True) - intersects = self._intersection_results_for_contains(lhs, rhs) - - # If a linestring has intersection but not containment, we need to - # test if the linestring is in the interior of the polygon. - final_result = _false_series(len(lhs)) - intersection_with_no_containment = (contains == 0) & (intersects != 0) - interior_tests = self._test_interior( - lhs[intersection_with_no_containment].reset_index(drop=True), - rhs[intersection_with_no_containment].reset_index(drop=True), + + # Count the number of point intersections (line crossings) between + # lhs and rhs. + # Also count the number of perfectly overlapping linestring sections. + # Each linestring overlap counts as two point overlaps. + ( + point_intersects_count, + linestring_intersects_count, + ) = self._intersection_results_for_contains_linestring(lhs, rhs) + + # Subtract the length of the linestring intersections from the length + # of the rhs linestring, then test that the sum of contained points + # is equal to that adjusted rhs length. + rhs_sizes_less_line_intersection_size = ( + rhs.sizes - linestring_intersects_count ) - interior_tests.index = intersection_with_no_containment[ - intersection_with_no_containment - ].index - # LineStrings that have intersection but no containment are set - # according to the `intersection_with_no_containment` mask. - final_result[intersection_with_no_containment] = interior_tests - # LineStrings that do not are contained if the sum of intersecting - # and containing points is greater than or equal to the number of - # points that make up the linestring. - final_result[~intersection_with_no_containment] = ( - contains + intersects >= rhs.sizes + rhs_sizes_less_line_intersection_size[ + rhs_sizes_less_line_intersection_size <= 0 + ] = 1 + final_result = contains + point_intersects_count == ( + rhs_sizes_less_line_intersection_size ) + return final_result def _compute_predicate(self, lhs, rhs, preprocessor_result): diff --git a/python/cuspatial/cuspatial/core/binpreds/feature_crosses.py b/python/cuspatial/cuspatial/core/binpreds/feature_crosses.py index 4650408e2..a8c869c79 100644 --- a/python/cuspatial/cuspatial/core/binpreds/feature_crosses.py +++ b/python/cuspatial/cuspatial/core/binpreds/feature_crosses.py @@ -66,7 +66,7 @@ def _compute_predicate(self, lhs, rhs, preprocessor_result): class LineStringPolygonCrosses(BinPred): def _preprocess(self, lhs, rhs): - intersects = _basic_intersects_count(rhs, lhs) > 1 + intersects = _basic_intersects_count(rhs, lhs) > 0 touches = rhs.touches(lhs) contains = rhs.contains(lhs) return ~touches & intersects & ~contains diff --git a/python/cuspatial/cuspatial/core/binpreds/feature_touches.py b/python/cuspatial/cuspatial/core/binpreds/feature_touches.py index d1b7b1686..f1493e495 100644 --- a/python/cuspatial/cuspatial/core/binpreds/feature_touches.py +++ b/python/cuspatial/cuspatial/core/binpreds/feature_touches.py @@ -24,6 +24,7 @@ Point, Polygon, _false_series, + _pli_points_to_multipoints, _points_and_lines_to_multipoints, ) @@ -101,18 +102,19 @@ def _preprocess(self, lhs, rhs): class LineStringPolygonTouches(BinPred): def _preprocess(self, lhs, rhs): + intersects = _basic_intersects_count(lhs, rhs) > 0 + contains = rhs.contains(lhs) + contains_any = _basic_contains_properly_any(rhs, lhs) + pli = _basic_intersects_pli(lhs, rhs) if len(pli[1]) == 0: return _false_series(len(lhs)) - intersections = _points_and_lines_to_multipoints(pli[1], pli[0]) + points = _pli_points_to_multipoints(pli) # A touch can only occur if the point in the intersection # is equal to a point in the linestring: it must # terminate in the boundary of the polygon. - equals = _basic_equals_count(intersections, lhs) > 0 - intersects = _basic_intersects_count(lhs, rhs) - intersects = (intersects == 1) | (intersects == 2) - contains = rhs.contains(lhs) - contains_any = _basic_contains_properly_any(rhs, lhs) + equals = _basic_equals_count(points, lhs) == points.sizes + return equals & intersects & ~contains & ~contains_any diff --git a/python/cuspatial/cuspatial/core/geoseries.py b/python/cuspatial/cuspatial/core/geoseries.py index 4f6b2c3ea..9fa1e316b 100644 --- a/python/cuspatial/cuspatial/core/geoseries.py +++ b/python/cuspatial/cuspatial/core/geoseries.py @@ -18,6 +18,7 @@ Point, Polygon, ) +from shapely.geometry.base import BaseGeometry, BaseMultipartGeometry import cudf from cudf._typing import ColumnLike @@ -27,6 +28,7 @@ import cuspatial.io.pygeoarrow as pygeoarrow from cuspatial.core._column.geocolumn import ColumnType, GeoColumn from cuspatial.core._column.geometa import Feature_Enum, GeoMeta +from cuspatial.core.binops.distance_dispatch import DistanceDispatch from cuspatial.core.binpreds.binpred_dispatch import ( CONTAINS_DISPATCH, CONTAINS_PROPERLY_DISPATCH, @@ -1391,3 +1393,93 @@ def touches(self, other, align=True): align=align ) return predicate(self, other) + + def isna(self): + """Detect missing values.""" + + c = self._column._meta.input_types == Feature_Enum.NONE.value + return cudf.Series(c, index=self.index) + + def notna(self): + """Detect non-missing values.""" + + c = self._column._meta.input_types != Feature_Enum.NONE.value + return cudf.Series(c, index=self.index) + + def distance(self, other, align=True): + """Returns a `Series` containing the distance to aligned other. + + The operation works on a 1-to-1 row-wise manner. See + `geopandas.GeoSeries.distance` documentation for details. + + Parameters + ---------- + other + The GeoSeries (elementwise) or geometric object to find the + distance to. + align : bool, default True + If True, automatically aligns GeoSeries based on their indices. + If False, the order of the elements is preserved. + + Returns + ------- + Series (float) + + Notes + ----- + Unlike GeoPandas, this API currently only supports geoseries that + contain only single type geometries. + + Examples + -------- + >>> from shapely.geometry import Point + >>> point = GeoSeries([Point(0, 0)]) + >>> point2 = GeoSeries([Point(1, 1)]) + >>> print(point.distance(point2)) + 0 1.414214 + dtype: float64 + + By default, geoseries are aligned before computing: + + >>> from shapely.geometry import Point + >>> point = GeoSeries([Point(0, 0)]) + >>> point2 = GeoSeries([Point(1, 1), Point(2, 2)]) + >>> print(point.distance(point2)) + 0 1.414214 + 1 NaN + dtype: float64 + + This can be overridden by setting `align=False`: + + >>> lines = GeoSeries([ + LineString([(0, 0), (1, 1)]), LineString([(2, 2), (3, 3)])]) + >>> polys = GeoSeries([ + Polygon([(0, 0), (1, 1), (1, 0)]), + Polygon([(2, 2), (3, 3), (3, 2)])], + index=[1, 0]) + >>> lines.distance(polys), align=False) + 0 0.0 + 1 0.0 + dtype: float64 + >>> lines.distance(polys, align=True) + 0 1.414214 + 1 1.414214 + dtype: float64 + """ + + other_is_scalar = False + if issubclass(type(other), (BaseGeometry, BaseMultipartGeometry)): + other_is_scalar = True + other = GeoSeries([other] * len(self), index=self.index) + + if not align: + if len(self) != len(other): + raise ValueError( + f"Lengths of inputs do not match. Left: {len(self)}, " + f"Right: {len(other)}" + ) + + res = DistanceDispatch(self, other, align)() + if other_is_scalar: + res.index = self.index + return res diff --git a/python/cuspatial/cuspatial/tests/binpreds/binpred_test_dispatch.py b/python/cuspatial/cuspatial/tests/binpreds/binpred_test_dispatch.py index 1d6fd6d61..7469a0325 100644 --- a/python/cuspatial/cuspatial/tests/binpreds/binpred_test_dispatch.py +++ b/python/cuspatial/cuspatial/tests/binpreds/binpred_test_dispatch.py @@ -359,6 +359,78 @@ def predicate(request): LineString([(0.5, 1.25), (0.5, -0.25)]), point_polygon, ), + "linestring-polygon-cross-concave-edge": ( + """ + x x x + |\\ | /| + | xx- | + | | | + ---x--- + """, + LineString([(0.5, 0.0), (0.5, 1.0)]), + Polygon([(0, 0), (0, 1), (0.3, 0.4), (1, 1), (1, 0)]), + ), + "linestring-polygon-half-in": ( + """ + ----- + | | + | x | + |/ \\| + xx-xx + """, + LineString( + [(0.0, 0.0), (0.25, 0.0), (0.5, 0.5), (0.75, 0.0), (1.0, 0.0)] + ), + point_polygon, + ), + "linestring-polygon-half-out": ( + """ + ----- + | | + | | + | | + xx-xx + \\/ + x + """, + LineString( + [(0.0, 0.0), (0.25, 0.0), (0.5, -0.5), (0.75, 0.0), (1.0, 0.0)] + ), + point_polygon, + ), + "linestring-polygon-two-edges": ( + """ + x---- + | | + | | + | | + x---x + """, + LineString([(0.0, 1.0), (0.0, 0.0), (1.0, 0.0)]), + point_polygon, + ), + "linestring-polygon-edge-to-interior": ( + """ + x---- + | | + | -x + |-/ | + x---- + """, + LineString([(0.0, 1.0), (0.0, 0.0), (1.0, 0.5)]), + point_polygon, + ), + "linestring-polygon-edge-cross-to-exterior": ( + """ + x------ + | | + | ---x + | --- | + x------ + """, + LineString([(0.0, 1.0), (0.0, 0.0), (1.5, 0.5)]), + point_polygon, + ), "polygon-polygon-disjoint": ( """ Polygon polygon tests use a triangle for the lhs and a square for the rhs. @@ -555,6 +627,12 @@ def predicate(request): "linestring-polygon-edge-interior", "linestring-polygon-in", "linestring-polygon-crosses", + "linestring-polygon-cross-concave-edge", + "linestring-polygon-half-in", + "linestring-polygon-half-out", + "linestring-polygon-two-edges", + "linestring-polygon-edge-to-interior", + "linestring-polygon-edge-cross-to-exterior", ] polygon_polygon_dispatch_list = [ diff --git a/python/cuspatial/cuspatial/tests/binpreds/test_binpred_internals.py b/python/cuspatial/cuspatial/tests/binpreds/test_binpred_internals.py index 9b87f821f..501b94e51 100644 --- a/python/cuspatial/cuspatial/tests/binpreds/test_binpred_internals.py +++ b/python/cuspatial/cuspatial/tests/binpreds/test_binpred_internals.py @@ -3,11 +3,14 @@ import pandas as pd from shapely.geometry import LineString, MultiPoint, Point, Polygon +import cudf + import cuspatial from cuspatial.core.binpreds.binpred_dispatch import EQUALS_DISPATCH from cuspatial.utils.binpred_utils import ( - _linestrings_to_center_point, _open_polygon_rings, + _pli_lines_to_multipoints, + _pli_points_to_multipoints, _points_and_lines_to_multipoints, ) @@ -318,22 +321,74 @@ def test_points_and_lines_to_multipoints_real_example(): assert (got.multipoints.xy == expected.multipoints.xy).all() -def test_linestrings_to_center_point(): - linestrings = cuspatial.GeoSeries( +def test_pli_points_to_multipoints_no_points(): + points = cuspatial.GeoSeries([]) + offsets = cudf.Series([0, 1, 2, 3]) + mpoints = _pli_points_to_multipoints((offsets, points)) + assert len(mpoints) == 3 + + +def test_pli_points_to_multipoints_first(): + points = cuspatial.GeoSeries([Point(1, 2)]) + offsets = cudf.Series([0, 1, 1, 1]) + mpoints = _pli_points_to_multipoints((offsets, points)) + assert (mpoints.sizes.values_host == [1, 0, 0]).all() + assert (mpoints.multipoints.xy.values_host == [1, 2]).all() + + +def test_pli_points_to_multipoints_two(): + points = cuspatial.GeoSeries( [ - LineString([(0, 0), (10, 10)]), - LineString([(5, 5), (6, 6)]), - LineString([(10, 10), (9, 9)]), - LineString([(11, 11), (1, 1)]), + Point(1, 2), + Point(3, 4), ] ) - expected = cuspatial.GeoSeries( + offsets = cudf.Series([0, 2, 2, 2]) + mpoints = _pli_points_to_multipoints((offsets, points)) + assert (mpoints.sizes.values_host == [2, 0, 0]).all() + assert (mpoints.multipoints.xy.values_host == [1, 2, 3, 4]).all() + + +def test_pli_points_to_multipoints_split(): + points = cuspatial.GeoSeries( + [ + Point(1, 2), + Point(3, 4), + Point(5, 6), + Point(7, 8), + ] + ) + offsets = cudf.Series([0, 2, 2, 4]) + mpoints = _pli_points_to_multipoints((offsets, points)) + assert (mpoints.sizes.values_host == [2, 0, 2]).all() + assert ( + mpoints.multipoints.xy.values_host == [1, 2, 3, 4, 5, 6, 7, 8] + ).all() + + +def test_pli_points_to_multipoints_drop_linestring(): + mixed = cuspatial.GeoSeries( + [ + Point(1, 2), + LineString([(3, 4), (5, 6)]), + Point(7, 8), + ] + ) + offsets = cudf.Series([0, 1, 2, 3]) + mpoints = _pli_points_to_multipoints((offsets, mixed)) + assert (mpoints.sizes.values_host == [1, 0, 1]).all() + assert (mpoints.multipoints.xy.values_host == [1, 2, 7, 8]).all() + + +def test_pli_lines_to_multipoints_drop_point(): + mixed = cuspatial.GeoSeries( [ - Point(5, 5), - Point(5.5, 5.5), - Point(9.5, 9.5), - Point(6, 6), + Point(1, 2), + LineString([(3, 4), (5, 6)]), + Point(7, 8), ] ) - got = _linestrings_to_center_point(linestrings) - assert (got.points.xy == expected.points.xy).all() + offsets = cudf.Series([0, 1, 2, 3]) + mpoints = _pli_lines_to_multipoints((offsets, mixed)) + assert (mpoints.sizes.values_host == [0, 2, 0]).all() + assert (mpoints[1:2].multipoints.xy.values_host == [3, 4, 5, 6]).all() diff --git a/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py b/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py index e3a67df6c..a50094737 100644 --- a/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py +++ b/python/cuspatial/cuspatial/tests/binpreds/test_contains_properly.py @@ -131,7 +131,7 @@ def test_float_precision_limits(point, polygon, expects): Unique success cases identified by @mharris. These go in a pair with test_float_precision_limits_failures because these are inconsistent results, where 0.6 fails above (as True, within the - polygon) and 0.66 below succeeds, though they are colinear. + polygon) and 0.66 below succeeds, though they are collinear. """ gpdpoint = gpd.GeoSeries(point) gpdpolygon = gpd.GeoSeries(polygon) diff --git a/python/cuspatial/cuspatial/tests/conftest.py b/python/cuspatial/cuspatial/tests/conftest.py index b204efde2..7b6ae2d1b 100644 --- a/python/cuspatial/cuspatial/tests/conftest.py +++ b/python/cuspatial/cuspatial/tests/conftest.py @@ -257,7 +257,7 @@ def generator(n, distance_from_origin, radius=1.0): @pytest.fixture -def multipolygon_generator(): +def multipolygon_generator(polygon_generator): """Generator for multi complex polygons. Usage: multipolygon_generator(n, max_per_multi) """ diff --git a/python/cuspatial/cuspatial/tests/spatial/distance/test_distance.py b/python/cuspatial/cuspatial/tests/spatial/distance/test_distance.py new file mode 100644 index 000000000..2ca6f7965 --- /dev/null +++ b/python/cuspatial/cuspatial/tests/spatial/distance/test_distance.py @@ -0,0 +1,150 @@ +from functools import partial +from itertools import product + +import geopandas +import pytest +from shapely.geometry import ( + LineString, + MultiLineString, + MultiPoint, + MultiPolygon, + Point, + Polygon, +) + +import cudf +from cudf.testing import assert_series_equal + +import cuspatial + +geometry_types = set( + [Point, LineString, Polygon, MultiPoint, MultiLineString, MultiPolygon] +) + + +@pytest.fixture(params=set(product(geometry_types, geometry_types))) +def _binary_op_combinations( + request, + point_generator, + multipoint_generator, + linestring_generator, + multilinestring_generator, + polygon_generator, + multipolygon_generator, +): + type_to_generator = { + Point: point_generator, + MultiPoint: multipoint_generator, + LineString: linestring_generator, + MultiLineString: multilinestring_generator, + Polygon: polygon_generator, + MultiPolygon: multipolygon_generator, + } + + generator_parameters = { + point_generator: {}, + multipoint_generator: {"max_num_geometries": 10}, + linestring_generator: {"max_num_segments": 10}, + multilinestring_generator: { + "max_num_geometries": 10, + "max_num_segments": 10, + }, + polygon_generator: {"distance_from_origin": 10, "radius": 10}, + multipolygon_generator: { + "max_per_multi": 10, + "distance_from_origin": 10, + "radius": 10, + }, + } + + ty0, ty1 = request.param + g0, g1 = type_to_generator[ty0], type_to_generator[ty1] + p0, p1 = generator_parameters[g0], generator_parameters[g1] + return (partial(g0, **p0), partial(g1, **p1)) + + +def test_geoseires_distance_empty(): + expected = geopandas.GeoSeries([]).distance(geopandas.GeoSeries([])) + actual = cuspatial.GeoSeries([]).distance(cuspatial.GeoSeries([])) + + assert_series_equal(actual, cudf.Series(expected)) + + +@pytest.mark.parametrize("n", [1, 100]) +@pytest.mark.parametrize("align", [True, False]) +@pytest.mark.parametrize( + "index_shuffle", + [ + (lambda x: x, lambda x: reversed(x)), + (lambda x: reversed(x), lambda x: x), + ], +) +def test_geoseries_distance_non_empty( + _binary_op_combinations, n, align, index_shuffle +): + g0, g1 = _binary_op_combinations + sfl0, sfl1 = index_shuffle + + h0 = geopandas.GeoSeries([*g0(n=n)], index=sfl0(range(n))) + h1 = geopandas.GeoSeries([*g1(n=n)], index=sfl1(range(n))) + + expected = h0.distance(h1, align=align) + + d0 = cuspatial.GeoSeries(h0) + d1 = cuspatial.GeoSeries(h1) + + actual = d0.distance(d1, align=align) + + assert_series_equal(actual, cudf.Series(expected)) + + +def test_geoseries_distance_indices_different(): + h0 = geopandas.GeoSeries([Point(0, 0), Point(1, 1)], index=[0, 1]) + h1 = geopandas.GeoSeries([Point(0, 0)], index=[0]) + + expected = h0.distance(h1) + + d0 = cuspatial.GeoSeries(h0) + d1 = cuspatial.GeoSeries(h1) + + actual = d0.distance(d1) + + assert_series_equal(actual, cudf.Series(expected, nan_as_null=False)) + + +def test_geoseries_distance_indices_different_not_aligned(): + h0 = geopandas.GeoSeries([Point(0, 0), Point(1, 1)], index=[0, 1]) + h1 = geopandas.GeoSeries([Point(0, 0)], index=[0]) + + d0 = cuspatial.GeoSeries(h0) + d1 = cuspatial.GeoSeries(h1) + + try: + h0.distance(h1, align=False) + except Exception as e: + with pytest.raises(e.__class__, match=e.__str__()): + d0.distance(d1, align=False) + + +@pytest.mark.parametrize("align", [True, False]) +@pytest.mark.parametrize( + "index", + [ + [0, 1, 2, 3, 4, 5, 6, 7, 8, 9], + [9, 8, 7, 6, 5, 4, 3, 2, 1, 0], + ], +) +def test_geoseries_distance_to_shapely_object( + _binary_op_combinations, align, index +): + g0, g1 = _binary_op_combinations + h0 = geopandas.GeoSeries([*g0(n=10)], index=index) + obj = [*g1(n=1)][0] + + expected = h0.distance(obj, align=align) + + d0 = cuspatial.GeoSeries(h0) + + actual = d0.distance(obj, align=align) + + assert_series_equal(actual, cudf.Series(expected)) diff --git a/python/cuspatial/cuspatial/tests/test_geoseries.py b/python/cuspatial/cuspatial/tests/test_geoseries.py index 1a66d4457..ce2f8ff3a 100644 --- a/python/cuspatial/cuspatial/tests/test_geoseries.py +++ b/python/cuspatial/cuspatial/tests/test_geoseries.py @@ -20,6 +20,7 @@ ) import cudf +from cudf.testing import assert_series_equal import cuspatial @@ -757,3 +758,31 @@ def test_from_polygons_xy_example(): [Polygon([(0, 0), (1, 1), (2, 2), (3, 3), (4, 4), (0, 0)])] ) gpd.testing.assert_geoseries_equal(gpolygon.to_geopandas(), hpolygon) + + +@pytest.mark.parametrize( + "s", + [ + gpd.GeoSeries(), + gpd.GeoSeries([Point(0, 0)]), + gpd.GeoSeries([None]), + gpd.GeoSeries([Point(0, 0), None, Point(1, 1)]), + gpd.GeoSeries([Point(0, 0), None, LineString([(1, 1), (2, 2)]), None]), + ], +) +def test_isna(s): + assert_series_equal(cudf.Series(s.isna()), cuspatial.GeoSeries(s).isna()) + + +@pytest.mark.parametrize( + "s", + [ + gpd.GeoSeries(), + gpd.GeoSeries([Point(0, 0)]), + gpd.GeoSeries([None]), + gpd.GeoSeries([Point(0, 0), None, Point(1, 1)]), + gpd.GeoSeries([Point(0, 0), None, LineString([(1, 1), (2, 2)]), None]), + ], +) +def test_notna(s): + assert_series_equal(cudf.Series(s.notna()), cuspatial.GeoSeries(s).notna()) diff --git a/python/cuspatial/cuspatial/utils/binpred_utils.py b/python/cuspatial/cuspatial/utils/binpred_utils.py index ec9cdedd9..7945e7837 100644 --- a/python/cuspatial/cuspatial/utils/binpred_utils.py +++ b/python/cuspatial/cuspatial/utils/binpred_utils.py @@ -310,6 +310,8 @@ def _open_polygon_rings(geoseries): def _points_and_lines_to_multipoints(geoseries, offsets): + # TODO: This modification drops linestrings which are not used + # in geopandas intersection results for contains. """Converts a geoseries of points and lines into a geoseries of multipoints. @@ -384,31 +386,6 @@ def _points_and_lines_to_multipoints(geoseries, offsets): return result -def _linestrings_to_center_point(geoseries): - if (geoseries.sizes != 2).any(): - raise ValueError( - "Geoseries must contain only linestrings with two points" - ) - x = geoseries.lines.x - y = geoseries.lines.y - return cuspatial.GeoSeries.from_points_xy( - cudf.DataFrame( - { - "x": ( - x[::2].reset_index(drop=True) - + x[1::2].reset_index(drop=True) - ) - / 2, - "y": ( - y[::2].reset_index(drop=True) - + y[1::2].reset_index(drop=True) - ) - / 2, - } - ).interleave_columns() - ) - - def _multipoints_is_degenerate(geoseries): """Only tests if the first two points are degenerate.""" offsets = geoseries.multipoints.geometry_offset[:-1] @@ -447,7 +424,10 @@ def _pli_features_rebuild_offsets(pli, features): if len(features) == 0: return _zero_series(len(pli[0])) - in_sizes = features.sizes + in_sizes = ( + features.sizes if len(features) > 0 else _zero_series(len(pli[0]) - 1) + ) + offsets = cudf.Series(pli[0]) offset_sizes = offsets[1:].reset_index(drop=True) - offsets[ :-1 @@ -507,7 +487,7 @@ def _pli_points_to_multipoints(pli): xy = ( points.points.xy if len(points.points.xy) > 0 - else cudf.Series([0.0, 0.0]) + else cudf.Series([], dtype=cp.float64) ) multipoints = cuspatial.GeoSeries.from_multipoints_xy(xy, offsets) return multipoints @@ -521,7 +501,11 @@ def _pli_lines_to_multipoints(pli): pairwise_linestring_intersection.""" lines = pli[1][pli[1].feature_types == Feature_Enum.LINESTRING.value] offsets = _pli_features_rebuild_offsets(pli, lines) - xy = lines.lines.xy if len(lines.lines.xy) > 0 else cudf.Series([0.0, 0.0]) + xy = ( + lines.lines.xy + if len(lines.lines.xy) > 0 + else cudf.Series([], dtype=cp.float64) + ) multipoints = cuspatial.GeoSeries.from_multipoints_xy(xy, offsets) return multipoints