From 1b716b1b52afc43e3cbe9718c78c81377f2d43f7 Mon Sep 17 00:00:00 2001 From: Jianting Zhang Date: Wed, 14 Aug 2019 16:37:08 -0700 Subject: [PATCH 001/102] init check-in (cpp/python) --- README.md | 156 +- cpp/CMakeLists.txt | 265 +++ .../Modules/ConfigureGoogleBenchmark.cmake | 59 + cpp/cmake/Modules/ConfigureGoogleTest.cmake | 59 + .../GoogleBenchmark.CMakeLists.txt.cmake | 19 + .../Templates/GoogleTest.CMakeLists.txt.cmake | 19 + cpp/include/cuspatial/coor_trans.hpp | 20 + cpp/include/cuspatial/cuspatial.h | 31 + cpp/include/cuspatial/hausdorff.hpp | 11 + cpp/include/cuspatial/haversine.hpp | 22 + cpp/include/cuspatial/pip2.hpp | 31 + cpp/include/cuspatial/pip_util.h | 95 + cpp/include/cuspatial/shared_util.h | 212 ++ cpp/include/cuspatial/soa_readers.hpp | 52 + cpp/include/cuspatial/stq2.hpp | 23 + cpp/include/cuspatial/stq_thrust.h | 24 + cpp/include/cuspatial/traj2.hpp | 74 + cpp/include/cuspatial/traj_thrust.h | 72 + cpp/include/cuspatial/traj_util.h | 30 + cpp/src/io/soa/ply_soa_reader.cu | 87 + cpp/src/io/soa/pnt_soa_reader.cu | 77 + cpp/src/io/soa/ts_soa_reader.cu | 39 + cpp/src/io/soa/uint_soa_reader.cu | 35 + cpp/src/shared/shared_util.cpp | 83 + cpp/src/spatial/hausdorff.cu | 279 +++ cpp/src/spatial/haversine.cu | 124 ++ cpp/src/spatial/ll2coor.cu | 137 ++ cpp/src/spatial/pip2.cu | 179 ++ cpp/src/stq/sw_xy.cu | 136 ++ cpp/src/traj/coor2traj.cu | 171 ++ cpp/src/traj/traj_distspeed.cu | 188 ++ cpp/src/traj/traj_sbbox.cu | 186 ++ cpp/src/traj/traj_util.cpp | 101 + cpp/tests/CMakeLists.txt | 93 + cpp/tests/spatial/hausdorff_test.cu | 169 ++ cpp/tests/spatial/hausdorff_util.h | 47 + cpp/tests/spatial/pip2_test.cu | 204 ++ cpp/tests/spatial/pip2_toy.cu | 115 + data/README.md | 35 + data/its.cat | 1 + data/json2soa.cpp | 230 ++ data/poly2soa.cpp | 262 +++ python/cuspatial/cuspatial/__init__.py | 0 python/cuspatial/cuspatial/_version.py | 520 +++++ .../cuspatial/cuspatial/bindings/GDFError.py | 7 + .../cuspatial/cuspatial/bindings/__init__.py | 0 .../cuspatial/cuspatial/bindings/cudf_cpp.pxd | 313 +++ .../cuspatial/cuspatial/bindings/cudf_cpp.pyx | 477 +++++ .../cuspatial/cuspatial/bindings/dlpack.pxd | 19 + .../cuspatial/cuspatial/bindings/dlpack.pyx | 134 ++ .../cuspatial/bindings/soa_readers.pxd | 20 + .../cuspatial/bindings/soa_readers.pyx | 79 + .../cuspatial/cuspatial/bindings/spatial.pxd | 30 + .../cuspatial/cuspatial/bindings/spatial.pyx | 95 + python/cuspatial/cuspatial/bindings/stq.pxd | 19 + python/cuspatial/cuspatial/bindings/stq.pyx | 30 + python/cuspatial/cuspatial/bindings/traj.pxd | 22 + python/cuspatial/cuspatial/bindings/traj.pyx | 72 + .../tests/hausdorff_clustering_test_toy.py | 44 + .../tests/hausdorff_distance_test_toy.py | 14 + .../tests/hausdorff_distance_verify.py | 104 + .../tests/haversine_distance_test.py | 25 + .../tests/haversine_distance_test_toy.py | 30 + .../cuspatial/tests/ll2coor_test_toy.py | 20 + .../cuspatial/tests/pip2_test_gdal.py | 36 + .../cuspatial/tests/pip2_test_shapely.py | 31 + .../cuspatial/tests/pip2_test_soa1.py | 20 + .../cuspatial/tests/pip2_test_soa2.py | 18 + .../cuspatial/tests/pip2_test_soa_toy.py | 17 + .../cuspatial/cuspatial/tests/pip2_verify.py | 51 + .../cuspatial/tests/stq_test_soa1.py | 21 + .../cuspatial/tests/traj2_test_soa1.py | 35 + .../cuspatial/tests/traj2_test_soa2.py | 26 + .../cuspatial/tests/traj2_test_soa3.py | 32 + .../cuspatial/tests/traj2_test_tools.py | 23 + python/cuspatial/setup.cfg | 61 + python/cuspatial/setup.py | 56 + python/cuspatial/versioneer.py | 1904 +++++++++++++++++ 78 files changed, 8585 insertions(+), 72 deletions(-) create mode 100644 cpp/CMakeLists.txt create mode 100644 cpp/cmake/Modules/ConfigureGoogleBenchmark.cmake create mode 100644 cpp/cmake/Modules/ConfigureGoogleTest.cmake create mode 100644 cpp/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake create mode 100644 cpp/cmake/Templates/GoogleTest.CMakeLists.txt.cmake create mode 100644 cpp/include/cuspatial/coor_trans.hpp create mode 100644 cpp/include/cuspatial/cuspatial.h create mode 100644 cpp/include/cuspatial/hausdorff.hpp create mode 100644 cpp/include/cuspatial/haversine.hpp create mode 100644 cpp/include/cuspatial/pip2.hpp create mode 100644 cpp/include/cuspatial/pip_util.h create mode 100644 cpp/include/cuspatial/shared_util.h create mode 100644 cpp/include/cuspatial/soa_readers.hpp create mode 100644 cpp/include/cuspatial/stq2.hpp create mode 100644 cpp/include/cuspatial/stq_thrust.h create mode 100644 cpp/include/cuspatial/traj2.hpp create mode 100644 cpp/include/cuspatial/traj_thrust.h create mode 100644 cpp/include/cuspatial/traj_util.h create mode 100644 cpp/src/io/soa/ply_soa_reader.cu create mode 100644 cpp/src/io/soa/pnt_soa_reader.cu create mode 100644 cpp/src/io/soa/ts_soa_reader.cu create mode 100644 cpp/src/io/soa/uint_soa_reader.cu create mode 100644 cpp/src/shared/shared_util.cpp create mode 100644 cpp/src/spatial/hausdorff.cu create mode 100644 cpp/src/spatial/haversine.cu create mode 100644 cpp/src/spatial/ll2coor.cu create mode 100644 cpp/src/spatial/pip2.cu create mode 100644 cpp/src/stq/sw_xy.cu create mode 100644 cpp/src/traj/coor2traj.cu create mode 100644 cpp/src/traj/traj_distspeed.cu create mode 100644 cpp/src/traj/traj_sbbox.cu create mode 100644 cpp/src/traj/traj_util.cpp create mode 100644 cpp/tests/CMakeLists.txt create mode 100644 cpp/tests/spatial/hausdorff_test.cu create mode 100644 cpp/tests/spatial/hausdorff_util.h create mode 100644 cpp/tests/spatial/pip2_test.cu create mode 100644 cpp/tests/spatial/pip2_toy.cu create mode 100644 data/README.md create mode 100644 data/its.cat create mode 100644 data/json2soa.cpp create mode 100644 data/poly2soa.cpp create mode 100644 python/cuspatial/cuspatial/__init__.py create mode 100644 python/cuspatial/cuspatial/_version.py create mode 100644 python/cuspatial/cuspatial/bindings/GDFError.py create mode 100644 python/cuspatial/cuspatial/bindings/__init__.py create mode 100644 python/cuspatial/cuspatial/bindings/cudf_cpp.pxd create mode 100644 python/cuspatial/cuspatial/bindings/cudf_cpp.pyx create mode 100644 python/cuspatial/cuspatial/bindings/dlpack.pxd create mode 100644 python/cuspatial/cuspatial/bindings/dlpack.pyx create mode 100644 python/cuspatial/cuspatial/bindings/soa_readers.pxd create mode 100644 python/cuspatial/cuspatial/bindings/soa_readers.pyx create mode 100644 python/cuspatial/cuspatial/bindings/spatial.pxd create mode 100644 python/cuspatial/cuspatial/bindings/spatial.pyx create mode 100644 python/cuspatial/cuspatial/bindings/stq.pxd create mode 100644 python/cuspatial/cuspatial/bindings/stq.pyx create mode 100644 python/cuspatial/cuspatial/bindings/traj.pxd create mode 100644 python/cuspatial/cuspatial/bindings/traj.pyx create mode 100644 python/cuspatial/cuspatial/tests/hausdorff_clustering_test_toy.py create mode 100644 python/cuspatial/cuspatial/tests/hausdorff_distance_test_toy.py create mode 100644 python/cuspatial/cuspatial/tests/hausdorff_distance_verify.py create mode 100644 python/cuspatial/cuspatial/tests/haversine_distance_test.py create mode 100644 python/cuspatial/cuspatial/tests/haversine_distance_test_toy.py create mode 100644 python/cuspatial/cuspatial/tests/ll2coor_test_toy.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_test_gdal.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_test_shapely.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_test_soa1.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_test_soa2.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_test_soa_toy.py create mode 100644 python/cuspatial/cuspatial/tests/pip2_verify.py create mode 100644 python/cuspatial/cuspatial/tests/stq_test_soa1.py create mode 100644 python/cuspatial/cuspatial/tests/traj2_test_soa1.py create mode 100644 python/cuspatial/cuspatial/tests/traj2_test_soa2.py create mode 100644 python/cuspatial/cuspatial/tests/traj2_test_soa3.py create mode 100644 python/cuspatial/cuspatial/tests/traj2_test_tools.py create mode 100644 python/cuspatial/setup.cfg create mode 100644 python/cuspatial/setup.py create mode 100644 python/cuspatial/versioneer.py diff --git a/README.md b/README.md index e9e296682..6c7fe9039 100644 --- a/README.md +++ b/README.md @@ -1,76 +1,88 @@ -# RAPIDS Standard Repo Template +# cuSpatial +GPU-Accelerated Spatial and Trajectory Data Management and Analytics Library
+**NOTE:** cuSpatial depends on [cudf](https://github.com/rapidsai/cudf) and [rmm](https://github.com/rapidsai/rmm) under [RAPDIS](https://rapids.ai/) framework
+See [here](https://nvidia-my.sharepoint.com/:p:/r/personal/jiantingz_nvidia_com/Documents/GPU4STA_V5.pptx?d=wa5b5d6d397074ea9a1600e74fd8a6345&csf=1&e=h7MdRq) +for a brief summary/discussion on C++ backend performance (as standalone components and with comparions to serial/mutli-core implementations on CPUs and/or legacy code)
+See the [deisgn documentation](doc/design.md) for a breif description on how spatial and trajectory data are represented in cuSpatial and the graph of operations on them. + +

Implemented operations:

+Currently, cuSpatial supports a subset of operations for spatial and trajectory data:
+1) [Spatial window query](cpp/src/stq)
+2) [Point-in-polygon test](cpp/src/spatial)
+3) [Harversine distance](cpp/src/spatial)
+4) [Hausdorff distance](cpp/src/spatial)
+5) [Deriving trajectories from point location data](cpp/src/traj)
+6) [Computing distance/speed of trajectories](cpp/src/traj)
+7) [Computing spatial bounding boxes of trajectories](cpp/src/traj)
+ +
+Another subset of operations will be added shortly:
+1) Temporal window query (cpp/src/stq)
+2) Temporal point query (year+month+day+hour+minute+second+millisecond)(cpp/src/stq)
+3) Point-to-polyline nearest neighbor distance](cpp/src/spatial)
+4) Grid based indexing for points and polygons (cpp/src/idx)
+5) Quadtree based indexing for large-scale point data (cpp/src/idx)
+6) R-Tree based indexing for Polygons/Polylines (cpp/src/idx)
+
+More operations are being planned/developed. + +

Compile/Install C++ backend

+To compile and run cuSpatial, use the following steps
+export CUSPATIAL_HOME=$(pwd)/cuspatial
+Step 1: clone a copy of cuSpatial (using your nvidia git-lab username/password)
+git clone https://github.com/zhangjianting/cuspatial.git -b fea-initial-code ${CUSPATIAL_HOME}
+
+ +Step 2: install cudf by following the [instructions](https://github.com/rapidsai/cudf/blob/branch-0.9/CONTRIBUTING.md)
+As a shortcut, one can just run the scripts "./build.sh libcudf" and "./build.sh cudf" under {CUSPATIAL_HOME}
+The rest of steps assume "export CUDACXX=/usr/local/cuda/bin/nvcc" and "export CUDF_HOME=$(pwd)/cudf are executed, and conda environment cudf_dev is activated after Step 2.
+Please note that, on some environments (OS+BASH), CUDF_HOME set by using "CUDF_HOME=..." when installing cuDF can not be used in cuSpatial. +Please use export CUDF_HOME=$(pwd)/cudf instead. + +Step 3: comile and install C++ backend
+ +cd $CUSPATIAL_HOME/cpp
+mkdir build
+cd build
+cmake .. -DCMAKE_INSTALL_PREFIX=$CONDA_PREFIX
+make (or make -j [n])
+make install
+ +cuSpatial should be installed at $CONDA_PREFIX, e.g., /home/jianting/anaconda3/envs/cudf_dev
+For cuspatial, the include path is $CONDA_PREFIX/include/cuspatial/ and the library path $CONDA_PREFIX/lib/libcuspatial.so, respetively. + +

Compile/Install Python wrapper and run Python test code

+ +Step 4: build and install python wrapper
+ +cd $CUSPATIAL_HOME/python/cuspatial
+python setup.py build_ext --inplace
+python setup.py install
+ +Step 5: Run python test code
+ +First,cuSpatial Python API path to PYTHONPATH (there are tools under tests subdir), i.e.,
+export PYTHONPATH=$CUSPATIAL_HOME/python/cuspatial
+ +Some test code using toy data can be run directly, e.g.,
+python $CUSPATIAL_HOME/python/cuspatial/cuspatial/tests/pip2_test_soa_toy.py
+ +However, many test code uses real data from an ITS (Intelligent Transportation System) application. +You will need to follow instructions at [data/README.md](./data/README.md) to generate data for these test code.
+Alternatively, you can download the preprocessed data("locust.*", "its_4326_roi.*", "itsroi.ply" and "its_camera_2.csv") from [here](https://nvidia-my.sharepoint.com/:u:/p/jiantingz/EdHR7qlaRSVPtw46XYVR9sQBjCcnUHygCuPUC3Hf8gW73A?e=LCr9nK), +extrat the files and put them directly under $CUSPATIAL_HOME/data for quick demos.
+A brief discription of these data files and their semantic roles in the ITS application can be found [here](doc/itsdata.md) + +After data are dowloaded and/or pre-processed, you can run the [python test code](python/cuspatial/cuspatial/tests), e.g.,
+python $CUSPATIAL_HOME/python/cuspatial/cuspatial/tests/pip2_verify.py
+python $CUSPATIAL_HOME/python/cuspatial/cuspatial/tests/traj2_test_soa3.py
+python $CUSPATIAL_HOME/python/cuspatial/cuspatial/tests/stq_test_soa1.py
+ +
+**NOTE:** Currently, cuSpatial supports reading point/polyine/polygon data using Structure of Array (SoA) format (more readers are being developed)
+Alternatively, python users can read any point/polyine/polygon data using existing python packages, e.g., [Shapely](https://pypi.org/project/Shapely/), +to generate numpy arrays and feed them to [cuSpatial python APIs](python/cuspatial/cuspatial/bindings).
-This is repo is the standard RAPIDS repo with the following items to make all RAPIDS repos consistent: -- GitHub File Templates - - Issue templates - - PR template -- GitHub Repo Templates - - Issue/PR labels - - Project tracking and release board templates -- Files - - `CHANGELOG.md` skeleton - - `CONTRIBUTING.md` skeleton - - `LICENSE` file with Apache 2.0 License - - `README.md` skeleton -## Usage for new RAPIDS repos - -1. Clone this repo -2. Find/replace all in the clone of `___PROJECT___` and replace with the name of the new library -3. Inspect files to make sure all replacements work and update text as needed -4. Customize issue/PR templates to fit the repo -5. Update `CHANGELOG.md` with next release version, see [changelog format](https://rapidsai.github.io/devdocs/docs/resources/changelog/) for more info -6. Add developer documentation to the end of the `CONTRIBUTING.md` that is project specific and useful for developers contributing to the project - - The goal here is to keep the `README.md` light, so the development/debugging information should go in `CONTRIBUTING.md` -7. Complete `README.md` with project description, quick start, install, and contribution information -8. Remove everything above the RAPIDS logo below from the `README.md` -9. Check `LICENSE` file is correct -10. Change git origin to point to new repo and push -11. Alert OPS team to copy labels and project boards to new repo - -## Usage for existing RAPIDS repos - -1. Follow the steps 1-9 above, but add the files to your existing repo and merge -2. Alert OPS team to copy labels and project boards to new repo - -## Useful docs to review - -- Issue triage & release planning - - [Issue triage process with GitHub projects](https://rapidsai.github.io/devdocs/docs/releases/triage/) - - [Release planning with GitHub projects](https://rapidsai.github.io/devdocs/docs/releases/planning/) -- Code release process - - [Hotfix process](https://rapidsai.github.io/devdocs/docs/releases/hotfix/) - - [Release process](https://rapidsai.github.io/devdocs/docs/releases/process/) -- Code contributions - - [Code contribution guide](https://rapidsai.github.io/devdocs/docs/contributing/code/) - - [Filing issues](https://rapidsai.github.io/devdocs/docs/contributing/issues/) - - [Filing PRs](https://rapidsai.github.io/devdocs/docs/contributing/prs/) - - [Code of conduct](https://rapidsai.github.io/devdocs/docs/resources/conduct/) -- Development process - - [Git branching and merging methodology](https://rapidsai.github.io/devdocs/docs/resources/git/) - - [Versions and tags](https://rapidsai.github.io/devdocs/docs/resources/versions/) - - [Changelog format](https://rapidsai.github.io/devdocs/docs/resources/changelog/) - - [Style guide](https://rapidsai.github.io/devdocs/docs/resources/style/) - - [Labels](https://rapidsai.github.io/devdocs/docs/maintainers/labels/) - ---- - -#
 ___PROJECT___
- -The [RAPIDS](https://rapids.ai) ___PROJECT___ ..._insert project description_... - -**NOTE:** For the latest stable [README.md](https://github.com/rapidsai/___PROJECT___/blob/master/README.md) ensure you are on the `master` branch. - -## Quick Start - -## Install ___PROJECT___ - -### Conda - -### Pip - -## Contributing Guide - -Review the [CONTRIBUTING.md](https://github.com/rapidsai/___PROJECT___/blob/master/CONTRIBUTING.md) file for information on how to contribute code and issues to the project. diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt new file mode 100644 index 000000000..6ffd2c1ce --- /dev/null +++ b/cpp/CMakeLists.txt @@ -0,0 +1,265 @@ +#============================================================================= +# Copyright (c) 2018-2019, 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. +#============================================================================= +cmake_minimum_required(VERSION 3.12 FATAL_ERROR) + +project(CUDA_SPATIAL VERSION 0.9.0 LANGUAGES C CXX CUDA) + +if(NOT CMAKE_CUDA_COMPILER) + message(SEND_ERROR "CMake cannot locate a CUDA compiler") +endif() + +################################################################################################### +# - build type ------------------------------------------------------------------------------------ + +# Set a default build type if none was specified +set(DEFAULT_BUILD_TYPE "Release") + +if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES) + message(STATUS "Setting build type to '${DEFAULT_BUILD_TYPE}' since none specified.") + set(CMAKE_BUILD_TYPE "${DEFAULT_BUILD_TYPE}" CACHE + STRING "Choose the type of build." FORCE) + # Set the possible values of build type for cmake-gui + set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS + "Debug" "Release" "MinSizeRel" "RelWithDebInfo") +endif() + +################################################################################################### +# - compiler options ------------------------------------------------------------------------------ + +set(CMAKE_CXX_STANDARD 14) +set(CMAKE_C_COMPILER $ENV{CC}) +set(CMAKE_CXX_COMPILER $ENV{CXX}) +set(CMAKE_CXX_STANDARD_REQUIRED ON) + +set(CMAKE_CUDA_STANDARD 14) +set(CMAKE_CUDA_STANDARD_REQUIRED ON) + +if(CMAKE_COMPILER_IS_GNUCXX) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Werror") + + option(CMAKE_CXX11_ABI "Enable the GLIBCXX11 ABI" ON) + if(CMAKE_CXX11_ABI) + message(STATUS "CUDF: Enabling the GLIBCXX11 ABI") + else() + message(STATUS "CUDF: Disabling the GLIBCXX11 ABI") + set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -D_GLIBCXX_USE_CXX11_ABI=0") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D_GLIBCXX_USE_CXX11_ABI=0") + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler -D_GLIBCXX_USE_CXX11_ABI=0") + endif(CMAKE_CXX11_ABI) +endif(CMAKE_COMPILER_IS_GNUCXX) + +#set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_60,code=sm_60 -gencode=arch=compute_61,code=sm_61") +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_60,code=sm_60") +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_70,code=sm_70 -gencode=arch=compute_70,code=compute_70") + +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-extended-lambda --expt-relaxed-constexpr") + +# set warnings as errors +# TODO: remove `no-maybe-unitialized` used to suppress warnings in rmm::exec_policy +set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Werror cross-execution-space-call -Xcompiler -Wall,-Werror") + +# Option to enable line info in CUDA device compilation to allow introspection when profiling / memchecking +option(CMAKE_CUDA_LINEINFO "Enable the -lineinfo option for nvcc (useful for cuda-memcheck / profiler" OFF) +if (CMAKE_CUDA_LINEINFO) + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -lineinfo") +endif(CMAKE_CUDA_LINEINFO) + +# Debug options +if(CMAKE_BUILD_TYPE MATCHES Debug) + message(STATUS "Building with debugging flags") + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -G -Xcompiler -rdynamic") +endif(CMAKE_BUILD_TYPE MATCHES Debug) + +# To apply RUNPATH to transitive dependencies (this is a temporary solution) +set(CMAKE_SHARED_LINKER_FLAGS "-Wl,--disable-new-dtags") +set(CMAKE_EXE_LINKER_FLAGS "-Wl,--disable-new-dtags") + +option(BUILD_TESTS "Configure CMake to build tests" + ON) + +option(BUILD_BENCHMARKS "Configure CMake to build (google) benchmarks" OFF) + +################################################################################################### +# - cmake modules --------------------------------------------------------------------------------- + +set(CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/Modules/" ${CMAKE_MODULE_PATH}) + +include(FeatureSummary) +include(CheckIncludeFiles) +include(CheckLibraryExists) + +################################################################################################### +# - RMM ------------------------------------------------------------------------------------------- + +find_path(RMM_INCLUDE "rmm" + HINTS "$ENV{RMM_ROOT}/include" + "$ENV{CONDA_PREFIX}/include/rmm" + "$ENV{CONDA_PREFIX}/include") + +find_library(RMM_LIBRARY "rmm" + HINTS "$ENV{RMM_ROOT}/lib" + "$ENV{CONDA_PREFIX}/lib") + +message(STATUS "RMM: RMM_ROOT set to $ENV{RMM_ROOT}") +message(STATUS "RMM: CONDA_PREFIX set to $ENV{CONDA_PREFIX}") +message(STATUS "RMM: RMM_LIBRARY set to ${RMM_LIBRARY}") +message(STATUS "RMM: RMM_INCLUDE set to ${RMM_INCLUDE}") + +add_library(rmm SHARED IMPORTED ${RMM_LIBRARY}) +if (RMM_INCLUDE AND RMM_LIBRARY) + set_target_properties(rmm PROPERTIES IMPORTED_LOCATION ${RMM_LIBRARY}) +endif (RMM_INCLUDE AND RMM_LIBRARY) + +# - CUDF ------------------------------------------------------------------------------------------- + +#find_path(CUDF_INCLUDE "cudf" +# "$ENV{CONDA_PREFIX}/include/cudf" +# "$ENV{CONDA_PREFIX}/include") + +#find_library(CUDF_LIBRARY "cudf" +# "$ENV{CONDA_PREFIX}/lib") + +SET(CUDF_INCLUDE "$ENV{CONDA_PREFIX}/include") +SET(CUDF_LIBRARY "$ENV{CONDA_PREFIX}/lib/libcudf.so") +SET(CUDF_SRC_INCLUDE "$ENV{CUDF_HOME}/cpp/src") +SET(CUB_INCLUDE "$ENV{CUDF_HOME}/cpp/thirdparty/cub") +SET(JITIFY_INCLUDE "$ENV{CUDF_HOME}/cpp/thirdparty/jitify") +SET(CUDF_TEST_INCLUDE "$ENV{CUDF_HOME}/cpp/") + +message(STATUS "CUDF: CUDF_HOME set to $ENV{CUDF_HOME}") +message(STATUS "CUDF: CONDA_PREFIX set to $ENV{CONDA_PREFIX}") +message(STATUS "CUDF: CUDF_LIBRARY set to ${CUDF_LIBRARY}") +message(STATUS "CUDF: CUDF_INCLUDE set to ${CUDF_INCLUDE}") +message(STATUS "CUDF SRC: CUDF_SRC_INCLUDE set to ${CUDF_SRC_INCLUDE}") +message(STATUS "CUDF: CUB_INCLUDE set to ${CUB_INCLUDE}") +message(STATUS "CUDF: JITIFY_INCLUDE set to ${JITIFY_INCLUDE}") +message(STATUS "CUDF: CUDF_TEST_INCLUDE set to ${CUDF_TEST_INCLUDE}") + +add_library(cudf SHARED IMPORTED ${CUDF_LIBRARY}) +if (CUDF_INCLUDE AND CUDF_LIBRARY) + set_target_properties(cudf PROPERTIES IMPORTED_LOCATION ${CUDF_LIBRARY}) +endif (CUDF_INCLUDE AND CUDF_LIBRARY) + +################################################################################################### +# - add gtest ------------------------------------------------------------------------------------- + +if(BUILD_TESTS) + include(CTest) + include(ConfigureGoogleTest) + + if(GTEST_FOUND) + message(STATUS "Google C++ Testing Framework (Google Test) found in ${GTEST_ROOT}") + include_directories(${GTEST_INCLUDE_DIR}) + add_subdirectory(${CMAKE_SOURCE_DIR}/tests) + else() + message(AUTHOR_WARNING "Google C++ Testing Framework (Google Test) not found: automated tests are disabled.") + endif(GTEST_FOUND) +endif(BUILD_TESTS) + +################################################################################################### +# - add google benchmark -------------------------------------------------------------------------- + +if(BUILD_BENCHMARKS) + + include(ConfigureGoogleBenchmark) + + if(GBENCH_FOUND) + message(STATUS "Google C++ Benchmarking Framework (Google Benchmark) found in ${GBENCH_ROOT}") + include_directories(${GBENCH_INCLUDE_DIR}) + add_subdirectory(${CMAKE_SOURCE_DIR}/benchmarks) + else() + message(AUTHOR_WARNING "Google C++ Benchmarking Framework (Google Benchmark) not found: automated tests are disabled.") + endif(GBENCH_FOUND) + +endif(BUILD_BENCHMARKS) + +################################################################################################### +# - include paths --------------------------------------------------------------------------------- + +if(CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES) + include_directories("${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}") +endif(CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES) + +include_directories("${CMAKE_BINARY_DIR}/include" + "${CMAKE_SOURCE_DIR}/include" + "${CMAKE_SOURCE_DIR}/src" + "${CMAKE_SOURCE_DIR}/thirdparty/cub" + "${RMM_INCLUDE}" + "${CUDF_SRC_INCLUDE}" + "${CUDF_INCLUDE}" + "${CUB_INCLUDE}" + "${JITIFY_INCLUDE}") + + +################################################################################################### +# - library paths --------------------------------------------------------------------------------- + +link_directories("${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}" # CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES is an undocumented/unsupported variable containing the link directories for nvcc + "${CMAKE_BINARY_DIR}/lib" + "${FLATBUFFERS_LIBRARY_DIR}" + "${GTEST_LIBRARY_DIR}" + "${RMM_LIBRARY}" + "${CUDF_LIBRARY}") + +################################################################################################### +# - library targets ------------------------------------------------------------------------------- + +add_library(cuspatial SHARED + src/shared/shared_util.cpp + src/io/soa/ply_soa_reader.cu + src/io/soa/pnt_soa_reader.cu + src/io/soa/uint_soa_reader.cu + src/io/soa/ts_soa_reader.cu + src/stq/sw_xy.cu + src/spatial/pip2.cu + src/spatial/haversine.cu + src/spatial/hausdorff.cu + src/spatial/ll2coor.cu + src/traj/traj_util.cpp + src/traj/traj_sbbox.cu + src/traj/traj_distspeed.cu + src/traj/coor2traj.cu + ) + +#Override RPATH for cuspatial +SET_TARGET_PROPERTIES(cuspatial PROPERTIES BUILD_RPATH "\$ORIGIN") + +################################################################################################### +# - build options --------------------------------------------------------------------------------- + +option(USE_NVTX "Build with NVTX support" ON) +if(USE_NVTX) + message(STATUS "Using Nvidia Tools Extension") + find_library(NVTX_LIBRARY nvToolsExt PATH ${CMAKE_CUDA_IMPLICIT_LINK_DIRECTORIES}) + target_link_libraries(cuspatial ${NVTX_LIBRARY}) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DUSE_NVTX") +endif(USE_NVTX) + + +################################################################################################### +# - link libraries -------------------------------------------------------------------------------- + +target_link_libraries(cuspatial cudf rmm cudart cuda nvrtc) + + +################################################################################################### +# - install targets ------------------------------------------------------------------------------- + +install(TARGETS cuspatial + DESTINATION lib) + +install(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/include/cuspatial + DESTINATION include) diff --git a/cpp/cmake/Modules/ConfigureGoogleBenchmark.cmake b/cpp/cmake/Modules/ConfigureGoogleBenchmark.cmake new file mode 100644 index 000000000..55ed165fc --- /dev/null +++ b/cpp/cmake/Modules/ConfigureGoogleBenchmark.cmake @@ -0,0 +1,59 @@ +set(GBENCH_ROOT "${CMAKE_BINARY_DIR}/googlebenchmark") + +set(GBENCH_CMAKE_ARGS " -DCMAKE_BUILD_TYPE=Release") + #" -Dgtest_build_samples=ON" + #" -DCMAKE_VERBOSE_MAKEFILE=ON") + +if(NOT CMAKE_CXX11_ABI) + message(STATUS "GBENCH: Disabling the GLIBCXX11 ABI") + list(APPEND GBENCH_CMAKE_ARGS " -DCMAKE_C_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=0") + list(APPEND GBENCH_CMAKE_ARGS " -DCMAKE_CXX_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=0") +elseif(CMAKE_CXX11_ABI) + message(STATUS "GBENCH: Enabling the GLIBCXX11 ABI") + list(APPEND GBENCH_CMAKE_ARGS " -DCMAKE_C_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=1") + list(APPEND GBENCH_CMAKE_ARGS " -DCMAKE_CXX_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=1") +endif(NOT CMAKE_CXX11_ABI) + +configure_file("${CMAKE_SOURCE_DIR}/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake" + "${GBENCH_ROOT}/CMakeLists.txt") + +file(MAKE_DIRECTORY "${GBENCH_ROOT}/build") +file(MAKE_DIRECTORY "${GBENCH_ROOT}/install") + +execute_process(COMMAND ${CMAKE_COMMAND} -G ${CMAKE_GENERATOR} . + RESULT_VARIABLE GBENCH_CONFIG + WORKING_DIRECTORY ${GBENCH_ROOT}) + +if(GBENCH_CONFIG) + message(FATAL_ERROR "Configuring Google Benchmark failed: " ${GBENCH_CONFIG}) +endif(GBENCH_CONFIG) + +set(PARALLEL_BUILD -j) +if($ENV{PARALLEL_LEVEL}) + set(NUM_JOBS $ENV{PARALLEL_LEVEL}) + set(PARALLEL_BUILD "${PARALLEL_BUILD}${NUM_JOBS}") +endif($ENV{PARALLEL_LEVEL}) + +if(${NUM_JOBS}) + if(${NUM_JOBS} EQUAL 1) + message(STATUS "GBENCH BUILD: Enabling Sequential CMake build") + elseif(${NUM_JOBS} GREATER 1) + message(STATUS "GBENCH BUILD: Enabling Parallel CMake build with ${NUM_JOBS} jobs") + endif(${NUM_JOBS} EQUAL 1) +else() + message(STATUS "GBENCH BUILD: Enabling Parallel CMake build with all threads") +endif(${NUM_JOBS}) + +execute_process(COMMAND ${CMAKE_COMMAND} --build .. -- ${PARALLEL_BUILD} + RESULT_VARIABLE GBENCH_BUILD + WORKING_DIRECTORY ${GBENCH_ROOT}/build) + +if(GBENCH_BUILD) + message(FATAL_ERROR "Building Google Benchmark failed: " ${GBENCH_BUILD}) +endif(GBENCH_BUILD) + +message(STATUS "Google Benchmark installed here: " ${GBENCH_ROOT}/install) +set(GBENCH_INCLUDE_DIR "${GBENCH_ROOT}/install/include") +set(GBENCH_LIBRARY_DIR "${GBENCH_ROOT}/install/lib" "${GBENCH_ROOT}/install/lib64") +set(GBENCH_FOUND TRUE) + diff --git a/cpp/cmake/Modules/ConfigureGoogleTest.cmake b/cpp/cmake/Modules/ConfigureGoogleTest.cmake new file mode 100644 index 000000000..d62bee2b1 --- /dev/null +++ b/cpp/cmake/Modules/ConfigureGoogleTest.cmake @@ -0,0 +1,59 @@ +set(GTEST_ROOT "${CMAKE_BINARY_DIR}/googletest") + +set(GTEST_CMAKE_ARGS "") + #" -Dgtest_build_samples=ON" + #" -DCMAKE_VERBOSE_MAKEFILE=ON") + +if(NOT CMAKE_CXX11_ABI) + message(STATUS "GTEST: Disabling the GLIBCXX11 ABI") + list(APPEND GTEST_CMAKE_ARGS " -DCMAKE_C_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=0") + list(APPEND GTEST_CMAKE_ARGS " -DCMAKE_CXX_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=0") +elseif(CMAKE_CXX11_ABI) + message(STATUS "GTEST: Enabling the GLIBCXX11 ABI") + list(APPEND GTEST_CMAKE_ARGS " -DCMAKE_C_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=1") + list(APPEND GTEST_CMAKE_ARGS " -DCMAKE_CXX_FLAGS=-D_GLIBCXX_USE_CXX11_ABI=1") +endif(NOT CMAKE_CXX11_ABI) + +configure_file("${CMAKE_SOURCE_DIR}/cmake/Templates/GoogleTest.CMakeLists.txt.cmake" + "${GTEST_ROOT}/CMakeLists.txt") + +file(MAKE_DIRECTORY "${GTEST_ROOT}/build") +file(MAKE_DIRECTORY "${GTEST_ROOT}/install") + +execute_process(COMMAND ${CMAKE_COMMAND} -G ${CMAKE_GENERATOR} . + RESULT_VARIABLE GTEST_CONFIG + WORKING_DIRECTORY ${GTEST_ROOT}) + +if(GTEST_CONFIG) + message(FATAL_ERROR "Configuring GoogleTest failed: " ${GTEST_CONFIG}) +endif(GTEST_CONFIG) + +set(PARALLEL_BUILD -j) +if($ENV{PARALLEL_LEVEL}) + set(NUM_JOBS $ENV{PARALLEL_LEVEL}) + set(PARALLEL_BUILD "${PARALLEL_BUILD}${NUM_JOBS}") +endif($ENV{PARALLEL_LEVEL}) + +if(${NUM_JOBS}) + if(${NUM_JOBS} EQUAL 1) + message(STATUS "GTEST BUILD: Enabling Sequential CMake build") + elseif(${NUM_JOBS} GREATER 1) + message(STATUS "GTEST BUILD: Enabling Parallel CMake build with ${NUM_JOBS} jobs") + endif(${NUM_JOBS} EQUAL 1) +else() + message(STATUS "GTEST BUILD: Enabling Parallel CMake build with all threads") +endif(${NUM_JOBS}) + +execute_process(COMMAND ${CMAKE_COMMAND} --build .. -- ${PARALLEL_BUILD} + RESULT_VARIABLE GTEST_BUILD + WORKING_DIRECTORY ${GTEST_ROOT}/build) + +if(GTEST_BUILD) + message(FATAL_ERROR "Building GoogleTest failed: " ${GTEST_BUILD}) +endif(GTEST_BUILD) + +message(STATUS "GoogleTest installed here: " ${GTEST_ROOT}/install) +set(GTEST_INCLUDE_DIR "${GTEST_ROOT}/install/include") +set(GTEST_LIBRARY_DIR "${GTEST_ROOT}/install/lib") +set(GTEST_FOUND TRUE) + diff --git a/cpp/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake b/cpp/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake new file mode 100644 index 000000000..c0f6abbb0 --- /dev/null +++ b/cpp/cmake/Templates/GoogleBenchmark.CMakeLists.txt.cmake @@ -0,0 +1,19 @@ +cmake_minimum_required(VERSION 3.12) + +include(ExternalProject) + +ExternalProject_Add(GoogleBenchmark + GIT_REPOSITORY https://github.com/google/benchmark.git + GIT_TAG master + SOURCE_DIR "${GBENCH_ROOT}/googlebenchmark" + BINARY_DIR "${GBENCH_ROOT}/build" + INSTALL_DIR "${GBENCH_ROOT}/install" + CMAKE_ARGS ${GBENCH_CMAKE_ARGS} -DBENCHMARK_ENABLE_TESTING=OFF -DCMAKE_INSTALL_PREFIX=${GBENCH_ROOT}/install) + # The flag BENCHMARK_ENABLE_TESTING=OFF prevents Google Benchmark from asking for Google Test. + + + + + + + diff --git a/cpp/cmake/Templates/GoogleTest.CMakeLists.txt.cmake b/cpp/cmake/Templates/GoogleTest.CMakeLists.txt.cmake new file mode 100644 index 000000000..66e1dc85a --- /dev/null +++ b/cpp/cmake/Templates/GoogleTest.CMakeLists.txt.cmake @@ -0,0 +1,19 @@ +cmake_minimum_required(VERSION 3.12) + +include(ExternalProject) + +ExternalProject_Add(GoogleTest + GIT_REPOSITORY https://github.com/google/googletest.git + GIT_TAG release-1.8.0 + SOURCE_DIR "${GTEST_ROOT}/googletest" + BINARY_DIR "${GTEST_ROOT}/build" + INSTALL_DIR "${GTEST_ROOT}/install" + CMAKE_ARGS ${GTEST_CMAKE_ARGS} -DCMAKE_INSTALL_PREFIX=${GTEST_ROOT}/install) + + + + + + + + diff --git a/cpp/include/cuspatial/coor_trans.hpp b/cpp/include/cuspatial/coor_trans.hpp new file mode 100644 index 000000000..e4d188001 --- /dev/null +++ b/cpp/include/cuspatial/coor_trans.hpp @@ -0,0 +1,20 @@ +#pragma once +#include + +namespace cuSpatial { + + +/** + * @Brief transforming in_lon/in_lat (lon/lat defined in Coord) to out_x/out_y relative to a camera origiin + + * @param[in] cam_lon/cam_lat: x and y coordiante of camera origin (lon/lat) + + * @param[in] in_lon/in_lat: arrays of x and y coordiantes (lon/lat) of input points to be transformed + + * @param[out] out_x/out_y: arrays of x and y coordiantes after transformation (in kilometers -km) + */ +void ll2coor(const gdf_scalar & cam_lon,const gdf_scalar & cam_lat,const gdf_column & in_lon,const gdf_column & in_lat, + gdf_column & out_x,gdf_column & out_y /* ,cudaStream_t stream = 0 */); + +} // namespace cuSpatial + diff --git a/cpp/include/cuspatial/cuspatial.h b/cpp/include/cuspatial/cuspatial.h new file mode 100644 index 000000000..0321e4c8e --- /dev/null +++ b/cpp/include/cuspatial/cuspatial.h @@ -0,0 +1,31 @@ +#pragma once + +typedef unsigned int uint; + + +typedef struct Location +{ + double lat; + double lon; + double alt; +} Location; + +typedef struct Coord +{ + double x; + double y; +} Coord; + +typedef struct Time +{ + uint y : 6; + uint m : 4; + uint d : 5; + uint hh : 5; + uint mm : 6; + uint ss : 6; + uint wd: 3; + uint yd: 9; + uint ms: 10; + uint pid:10; +}Time; \ No newline at end of file diff --git a/cpp/include/cuspatial/hausdorff.hpp b/cpp/include/cuspatial/hausdorff.hpp new file mode 100644 index 000000000..5c5b54987 --- /dev/null +++ b/cpp/include/cuspatial/hausdorff.hpp @@ -0,0 +1,11 @@ +#pragma once +#include + +namespace cuSpatial { + + +gdf_column hausdorff_distance(const gdf_column& coor_x,const gdf_column& coor_y,const gdf_column& traj_cnt + /* ,cudaStream_t stream = 0 */); + +} // namespace cuSpatial + diff --git a/cpp/include/cuspatial/haversine.hpp b/cpp/include/cuspatial/haversine.hpp new file mode 100644 index 000000000..67ce51614 --- /dev/null +++ b/cpp/include/cuspatial/haversine.hpp @@ -0,0 +1,22 @@ +#pragma once +#include + +namespace cuSpatial { + +/** + * @param[in] x1: pointer/array of x coordiantes of points + + * @param[in] y1: pointer/array of y coordiantes of points + + * @param[in] x2: pointer/array to index polygons, i.e., prefix-sum of #of rings of all polygons + + * @param[in] y2: pointer/array to index rings, i.e., prefix-sum of #of vertices of all rings + + * @param[out] dist: distance array for each (x1,y1) and (x2,y2) point pari + */ + +gdf_column haversine_distance(const gdf_column& x1,const gdf_column& y1,const gdf_column& x2,const gdf_column& y2 + /* , cudaStream_t stream = 0 */); + +} // namespace cuSpatial + diff --git a/cpp/include/cuspatial/pip2.hpp b/cpp/include/cuspatial/pip2.hpp new file mode 100644 index 000000000..c479fbf9f --- /dev/null +++ b/cpp/include/cuspatial/pip2.hpp @@ -0,0 +1,31 @@ +#pragma once +#include + +namespace cuSpatial { + +/** + * @param[in] pnt_x: pointer/array of x coordiantes of points + + * @param[in] pnt_y: pointer/array of y coordiantes of points + + * @param[in] ply_fpos: pointer/array to index polygons, i.e., prefix-sum of #of rings of all polygons + + * @param[in] ply_rpos: pointer/array to index rings, i.e., prefix-sum of #of vertices of all rings + + * @param[in] ply_x: pointer/array of x coordiantes of concatenated polygons + + * @param[in] ply_y: pointer/array of x coordiantes of concatenated polygons + * + * @returns gdf_column of type GDF_INT32; the jth bit of the ith element of the returned GDF_INT32 array indicate + * whether the ith point is in the jth polygon. + + * Note: The # of polygons, i.e., ply_fpos.size can not exceed sizeof(uint)*8, i.e., 32. + */ + +gdf_column pip2_bm(const gdf_column& pnt_x,const gdf_column& pnt_y, + const gdf_column& ply_fpos, const gdf_column& ply_rpos, + const gdf_column& ply_x,const gdf_column& ply_y + /* , cudaStream_t stream = 0 */); + +} // namespace cuSpatial + diff --git a/cpp/include/cuspatial/pip_util.h b/cpp/include/cuspatial/pip_util.h new file mode 100644 index 000000000..85912a156 --- /dev/null +++ b/cpp/include/cuspatial/pip_util.h @@ -0,0 +1,95 @@ +#pragma once + +#include +#include + +/** + * @brief sequential point-in-polygon test between a single point and a single polygon; + basic unit for either sequential execution or parallellization in multi-point/multi-polygon test cases + * + * @param[in] x /y the x/y coordinate of the input point + + * @param[in] ply complete metadata for a polygon dataset (with multiple polygons) + + * @param[in] fid index of the polygon dataset to identify the polyogn to be tested + * + * @return whehter the point is in the polygon + */ +template +bool pip_test_sequential(const T& x, const T& y, const struct PolyMeta& ply,int fid) +{ + //printf("pip: x=%15.10f y=%15.10f\n",x,y); + uint *f_pos=ply.p_f_pos; + uint *r_pos=ply.p_r_pos; + T *poly_x=ply.p_x; + T *poly_y=ply.p_y; + uint r_f = (0 == fid) ? 0 : f_pos[fid-1]; + uint r_t=f_pos[fid]; + //printf("(%d %d)=>\n",r_f,r_t); + bool in_polygon = false; + for (uint r = r_f; r < r_t; r++) //for each ring + { + uint m = (r==0)?0:r_pos[r-1]; + for (;m < r_pos[r]-1; m++) //for each line segment + { + T x0, x1, y0, y1; + x0 = poly_x[m]; + y0 = poly_y[m]; + x1 = poly_x[m+1]; + y1 = poly_y[m+1]; + //printf("r=%d m=%d, x0=%15.10f y0=%15.10f x1=%15.10f y1=%15.10f\n",r, m,x0,y0,x1,y1); + if ((((y0 <= y) && (y < y1)) || + ((y1 <= y) && (y < y0))) && + (x < (x1 - x0) * (y - y0) / (y1 - y0) + x0)) + in_polygon = !in_polygon; + } + } + return (in_polygon); +} + +template bool pip_test_sequential(const double& x, const double& y, const struct PolyMeta& ply,int fid); +template bool pip_test_sequential(const float& x, const float& y, const struct PolyMeta& ply,int fid); + + +/** + * @brief multi-point/multi-polygon test on CUPs with the same interface as the GPU implementation + parallelization (e.g., OpenMP and Intel TBB) can be applied to the array/vector of points. + + @param[in] num_pnt number of points + + @param[in] x pointer/array of x coodinates + + @param[in] x pointer/array of y coodinates + + @param[in] ply complete metadata for a polygon dataset (with multiple polygons) + + @param[out] res pointer/array of unsinged integers; the jth bit of res[i] indicates whehter + a point of (x[i],y[i]) is in polygon j. + + Note: The # of polygons, i.e., poly.f_num can not exceed sizeof(uint)*8, i.e., 32. + +*/ +template +void cpu_pip_loop(int num_pnt,const T* x, const T *y ,const struct PolyMeta& poly,uint* res) +{ + //assert(res!=NULL); + for(int i=0;i(x[i],y[i],poly,j); + if(in_polygon) + { + //printf("pip: %10d %10d\n",i,j); + mask|=(0x01<& poly,uint* res); +template void cpu_pip_loop(int num_pnt,const float *x, const float *y ,const struct PolyMeta& poly,uint* res); \ No newline at end of file diff --git a/cpp/include/cuspatial/shared_util.h b/cpp/include/cuspatial/shared_util.h new file mode 100644 index 000000000..4c347b8b1 --- /dev/null +++ b/cpp/include/cuspatial/shared_util.h @@ -0,0 +1,212 @@ +#pragma once + +#include +//for partial_sum +#include +#include +#include +#include +#include + +/** + * @brief struct for SoA (Structure of Array) representation of a set of polygons with multiple rings + * The group level hierarchy variables (num_g, p_g_len and p_g_pos) are not exposed in cuSpatial but + * can be added later for applications that need to orgnaize polygons into groups for fast indexing + p_{g,f,r}_pos should be the the prefix-sum (inclusive) of p_{g,f,r}_len, + using e.g., std::partial_sum(...) and thrust::inclusive_scan (...) + If is_inplace is set to true, p_{g,f,r}_pos will be the same as p_{g,f,r}_len to reduce memory footprint + This is useful only when it can be sure that p_{g,f,r}_len are not needed anymore (e.g., in CPU for sequentail access) + is_inplace should be set to false (default) when p_{g,f,r}_len and p_{g,f,r}_pos are used together (on GPUs for efficiency) + * + **/ + +template +struct PolyMeta +{ + uint num_g,num_f,num_r,num_v; + uint * p_g_len=NULL,*p_f_len=NULL,*p_r_len=NULL; + uint * p_g_pos=NULL,*p_f_pos=NULL,*p_r_pos=NULL; + T *p_x=NULL,*p_y=NULL; + bool is_inplace=false; +}; + + +/** + * @brief output operator for Time + * + **/ +std::ostream& operator<<(std::ostream& os, const Time & t); + +/** + * @brief read a set of Location data (lon/lat/alt) into separate x,y arrays + * + **/ +template +int read_point_ll(const char *fn,T *& x, T *& y) +{ + std::cout<<"reading point data from file"< +int read_point_xy(const char *fn,T *& x, T *& y) +{ + std::cout<<"reading point data from file"< +int read_polygon_soa(const char *poly_fn,struct PolyMeta& ply) +{ + FILE *fp=fopen(poly_fn,"rb"); + CUDF_EXPECTS(fp!=NULL,"can not open the input polygon file"); + + //verify file integrity + fseek(fp, 0L, SEEK_END); + size_t sz=ftell(fp); + fseek(fp, 0L, SEEK_SET); + + size_t ln=0; + ln=fread(&(ply.num_g),sizeof(int),1,fp); + CUDF_EXPECTS(ln==1,"expect reading an integer for num_g"); + ln=fread(&(ply.num_f),sizeof(int),1,fp); + CUDF_EXPECTS(ln==1,"expect reading an integer for num_f"); + ln=fread(&(ply.num_r),sizeof(int),1,fp); + CUDF_EXPECTS(ln==1,"expect reading an integer for num_r"); + ln=fread(&(ply.num_v),sizeof(int),1,fp); + CUDF_EXPECTS(ln==1,"expect reading an integer for num_v"); + + //std::cout<<"# of groups="<< ply.num_g<& ply); +int read_polygon_soa(const char *poly_fn,struct PolyMeta& ply); + +/** + * @brief read a CSV file into a map + * + **/ +int read_csv(const char *fn, std::vector cols,int num_col,std::map>& df); + +/** + * @brief timing function to calaculate duration between t1 and t0 in milliseconds and output the duration proceeded by msg + * + **/ +float calc_time(const char *msg,timeval t0, timeval t1); + + +/** + * @brief templated function to read data in SoA format from file to array + for the repective data type or user defined struct. + * + **/ +template +size_t read_field(const char *fn,T*& field) +{ + FILE *fp=NULL; + if((fp=fopen(fn,"rb"))==NULL) + CUDF_EXPECTS(fp!=NULL,"can not open the input point file"); + + fseek (fp , 0 , SEEK_END); + size_t sz=ftell (fp); + CUDF_EXPECTS(sz%sizeof(T)==0,"sizeof(T) does not divide file length"); + size_t num_rec = sz/sizeof(T); + std::cout<<"num_rec="< +#include + +namespace cuSpatial { + +/** + * @Brief read uint (unsigned integer) data from file as column +* +*/ +void read_uint_soa(const char *ID_fn, gdf_column& ids); + +/** + * @Brief read timestamp (ts: Time type) data from file as column +* +*/ +void read_ts_soa(const char *ts_fn, gdf_column& ts); + +/** + * @Brief read lon/lat from file as two columns; data type is fixed to double (GDF_FLOAT64) +* +*/ +void read_pnt_lonlat_soa(const char *pnt_fn, gdf_column& pnt_lon,gdf_column& pnt_lat); + + +/** + * @Brief read x/y from file as two columns; data type is fixed to double (GDF_FLOAT64) +* +*/ +void read_pnt_xy_soa(const char *pnt_fn, gdf_column& pnt_x,gdf_column& pnt_y); + + +/** + * @Brief read poygon data from file in SoA format; data type of vertices is fixed to double (GDF_FLOAT64) + + * @param[in] ply_fn: polygon data file name + + * @param[out] ply_fpos: pointer/array to index polygons, i.e., prefix-sum of #of rings of all polygons + + * @param[out] ply_rpos: pointer/array to index rings, i.e., prefix-sum of #of vertices of all rings + + * @param[out] ply_x: pointer/array of x coordiantes of concatenated polygons + + * @param[out] ply_y: pointer/array of x coordiantes of concatenated polygons +* +*/ + +void read_ply_soa(const char *ply_fn,gdf_column& ply_fpos, gdf_column& ply_rpos, + gdf_column& ply_x,gdf_column& ply_y); + +} diff --git a/cpp/include/cuspatial/stq2.hpp b/cpp/include/cuspatial/stq2.hpp new file mode 100644 index 000000000..d9998a3be --- /dev/null +++ b/cpp/include/cuspatial/stq2.hpp @@ -0,0 +1,23 @@ +#pragma once + +#include + +namespace cuSpatial { + +/** + * @Brief retrive all points (x,y) that fall within a query window (x1,y1,x2,y2) and output the filtered points + + * @param[in] x1/x2/y1/y2: defines the query window + + * @param[in] in_x/in_y: pointer/array of x/y coordiantes of points to be queried + + * @param[out] out_x/out_y: pointer/array of x/y coordiantes of points that fall within the query window + + * @returns the number of points that satisfy the spatial window query criteria + */ + +int sw_xy(const gdf_scalar x1,const gdf_scalar y1,const gdf_scalar x2,const gdf_scalar y2, + const gdf_column& in_x,const gdf_column& in_y, gdf_column& out_x,gdf_column& out_y); + +} // namespace cuSpatial + diff --git a/cpp/include/cuspatial/stq_thrust.h b/cpp/include/cuspatial/stq_thrust.h new file mode 100644 index 000000000..c1f468478 --- /dev/null +++ b/cpp/include/cuspatial/stq_thrust.h @@ -0,0 +1,24 @@ +#pragma once + +/** + *@brief Thrust functor for spatial window query on point data (x/y) + */ +template +struct sw_functor_xy +{ + T x1,y1,x2,y2; + + __host__ __device__ + sw_functor_xy(T _x1,T _x2,T _y1,T _y2): + x1(_x1),y1(_y1),x2(_x2),y2(_y2){}; + + __host__ __device__ + bool operator()(const thrust::tuple& t) + { + T x= thrust::get<0>(t); + T y= thrust::get<1>(t); + bool b1=x>x1 && xx1 && y + +namespace cuSpatial { + +/** + * @Brief deriving trajectories from points (x/y relative to an origin), timestamps and objectids + * by first sorting based on id and timestamp and then group by id. + + * @param[in/out] coor_x/coor_y: x and y coordiantes reative to a camera origin (before/after sorting) + + * @param[in/out] oid: object (e.g., vehicle) id column (before/after sorting); upon completion, unqiue ids become trajectory ids + + * @param[in/out] ts: timestamp column (before/after sorting) + + * @param[out] tid: trajectory id column (see comments on oid) + + * @param[out] len: #of points in the derived trajectories + + * @param[out] pos: position offsets of trajectories used to index coor_x/coor_y/oid/ts after completion + + * @returns the number of derived trajectory + + */ +int coor2traj(gdf_column& coor_x,gdf_column& coor_y,gdf_column& oid, gdf_column& ts, + gdf_column& tid, gdf_column& len,gdf_column& pos/* ,cudaStream_t stream = 0 */); + + +/** + * @Brief computing distance(length) and speed of trajectories after their formation (e.g., from coor2traj) + + * @param[in] coor_x/coor_y: x and y coordiantes reative to a camera origin ordered by (id,timestamp) + + * @param[in] ts: timestamp column ordered by (id,timestamp) + + * @param[in] len: number of points column ordered by (id,timestamp) + + * @param[in] pos: position offsets of trajectories used to index coor_x/coor_y/oid/ts ordered by (id,timestamp) + + * @param[out] dist: computed distances/lengths of trajectories in meters (m) + + * @param[out] speed: computed speed of trajectories in meters per second (m/s) + + * Note: we might output duration (in addtiion to distance/speed), should there is a need + * duration can be easiy computed on CPUs by fetching begining/ending timestamps of a trajectory in the timestamp array + */ + +void traj_distspeed(const gdf_column& coor_x,const gdf_column& coor_y,const gdf_column& ts, + const gdf_column& len,const gdf_column& pos,gdf_column& dist,gdf_column& speed + /* ,cudaStream_t stream = 0 */); + + + +/** + * @Brief computing spatial bounding boxes of trjectories + + * @param[in] coor_x/coor_y: x and y coordiantes reative to a camera origin ordered by (id,timestamp) + + * @param[in] len: number of points column ordered by (id,timestamp) + + * @param[in] pos: position offsets of trajectories used to index coor_x/coor_y/ ordered by (id,timestamp) + + * @param[out] bbox_x1/bbox_y1/bbox_x2/bbox_y2: computed spaital bounding boxes in four columns + + * Note: temporal 1D bounding box can be computed similary but it seems that there is no such a need; + * Similar to the dicussion in coor2traj, the temporal 1D bounding box can be retrieved directly + */ + +void traj_sbbox(const gdf_column& coor_x,const gdf_column& coor_y, + const gdf_column& len,const gdf_column& pos, + gdf_column& bbox_x1,gdf_column& bbox_y1,gdf_column& bbox_x2,gdf_column& bbox_y2 + /* ,cudaStream_t stream = 0 */); +} // namespace cuSpatial diff --git a/cpp/include/cuspatial/traj_thrust.h b/cpp/include/cuspatial/traj_thrust.h new file mode 100644 index 000000000..0e030b716 --- /dev/null +++ b/cpp/include/cuspatial/traj_thrust.h @@ -0,0 +1,72 @@ +#pragma once + +#include +#include + +/** + *@brief Thrust functor for comparing two Time variables; used in sorting based on timestamp + */ +__host__ __device__ +inline bool operator<(const Time & t1,const Time & t2) +{ + //cout<<"in operator<"< TBBox; + +/** + *@brief Thrust functor for lifting Time to an interval (1D box) + */ +struct TBBox_transformation : public thrust::unary_function +{ + __host__ __device__ + TBBox operator()(Time time) + { + return TBBox(time, time); + } +}; + +/** + *@brief Thrust functor for generating an 1D bounding box from two 1D box + */ +struct TBBox_reduction : public thrust::binary_function +{ + __host__ __device__ + TBBox operator()(TBBox a, TBBox b) + { + // lower left corner + Time minT=(a.first +{ + Location origin; + __host__ __device__ + coord_transformation(Location _origin): origin(_origin){} + + __host__ __device__ + Coord operator()(Location pt) + { + Coord c; + c.x = ((origin.lon - pt.lon) * 40000.0 * + cos((origin.lat + pt.lat) * M_PI / 360) / 360); + c.y = (origin.lat - pt.lat) * 40000.0 / 360; + return c; + } +}; \ No newline at end of file diff --git a/cpp/include/cuspatial/traj_util.h b/cpp/include/cuspatial/traj_util.h new file mode 100644 index 000000000..a3ebb41b2 --- /dev/null +++ b/cpp/include/cuspatial/traj_util.h @@ -0,0 +1,30 @@ +#pragma once + +#include + +/** + * @Brief retrive camera origin at a particular intersection from a configuration file + + * @param[in] df_fn: name of configuration file that contains paramters for all cameras + + * @param[in] inter_name: the unique name of an intersection + + * @param[out] camera_origin Location (lon/lat/alt) of the retrieved camera origin + * + * @return negative error code; 0 for success + */ +int get_camera_origin(const char *df_fn, const char * inter_name, Location & camera_origin); + + +/** + * @Brief retrive id, timestamp and location fields of a "raw" trajectory dataset, + * i.e., a set of coordiantes (lon/lat/alt) with a timestamp and an object (e.g., vehicle) identifier. + * + * @param[in] root_fn: the root of the three files stored in columnar format, + * with .objectid (uint type),.time (Time type) and .location(Location type) extensions, respectively. + * @param[out] objid: out array for ID + * @param[out] time: out array for ID + * @param[out] location: out array for ID + * @return the number of records (should be the same for all the three data files) + */ +size_t read_traj_soa(char *root_fn,int *& objid, Time *& time, Location*& location); diff --git a/cpp/src/io/soa/ply_soa_reader.cu b/cpp/src/io/soa/ply_soa_reader.cu new file mode 100644 index 000000000..2ced6331d --- /dev/null +++ b/cpp/src/io/soa/ply_soa_reader.cu @@ -0,0 +1,87 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cuSpatial +{ + /** + * @Brief read poygon data from file in SoA format; data type of vertices is fixed to double (GDF_FLOAT64) + + * @param[in] ply_fn: polygon data file name + + * @param[out] ply_fpos: pointer/array to index polygons, i.e., prefix-sum of #of rings of all polygons + + * @param[out] ply_rpos: pointer/array to index rings, i.e., prefix-sum of #of vertices of all rings + + * @param[out] ply_x: pointer/array of x coordiantes of concatenated polygons + + * @param[out] ply_y: pointer/array of x coordiantes of concatenated polygons + * + */ + + void read_ply_soa(const char *poly_fn,gdf_column& ply_fpos, gdf_column& ply_rpos, + gdf_column& ply_x,gdf_column& ply_y) + { + struct PolyMeta pm; + int num_p=read_polygon_soa(poly_fn,pm); + if(num_p<=0) return; + + ply_fpos.dtype=GDF_INT32; + ply_fpos.col_name=(char *)malloc(strlen("f_pos")+ 1); + strcpy(ply_fpos.col_name,"f_pos"); + ply_fpos.data=NULL; + RMM_TRY( RMM_ALLOC(&ply_fpos.data, pm.num_f * sizeof(uint), 0) ); + cudaMemcpy(ply_fpos.data, pm.p_f_len,pm.num_f * sizeof(uint) , cudaMemcpyHostToDevice); + thrust::device_ptr d_pfp_ptr=thrust::device_pointer_cast((uint *)ply_fpos.data); + //prefix-sum: len to pos + thrust::inclusive_scan(d_pfp_ptr,d_pfp_ptr+pm.num_f,d_pfp_ptr); + ply_fpos.size=pm.num_f; + ply_fpos.valid=nullptr; + ply_fpos.null_count=0; + delete[] pm.p_f_len; + + ply_rpos.dtype=GDF_INT32; + ply_rpos.col_name=(char *)malloc(strlen("r_pos")+ 1); + strcpy(ply_rpos.col_name,"r_pos"); + ply_rpos.data=NULL; + RMM_TRY( RMM_ALLOC(&ply_rpos.data, pm.num_r * sizeof(uint), 0) ); + cudaMemcpy(ply_rpos.data, pm.p_r_len,pm.num_r * sizeof(uint) , cudaMemcpyHostToDevice); + thrust::device_ptr d_prp_ptr=thrust::device_pointer_cast((uint *)ply_rpos.data); + //prefix-sum: len to pos + thrust::inclusive_scan(d_prp_ptr,d_prp_ptr+pm.num_r,d_prp_ptr); + ply_rpos.size=pm.num_r; + ply_rpos.valid=nullptr; + ply_rpos.null_count=0; + delete[] pm.p_r_len; + + ply_x.dtype= GDF_FLOAT64; + ply_x.col_name=(char *)malloc(strlen("x")+ 1); + strcpy(ply_x.col_name,"x"); + RMM_TRY( RMM_ALLOC(&ply_x.data, pm.num_v * sizeof(double), 0) ); + cudaMemcpy(ply_x.data, pm.p_x,pm.num_v * sizeof(double) , cudaMemcpyHostToDevice); + ply_x.size=pm.num_v; + ply_x.valid=nullptr; + ply_x.null_count=0; + delete[] pm.p_x; + + ply_y.dtype= GDF_FLOAT64; + ply_y.col_name=(char *)malloc(strlen("y")+ 1); + strcpy(ply_y.col_name,"y"); + ply_y.data=NULL; + RMM_TRY( RMM_ALLOC(&ply_y.data, pm.num_v * sizeof(double), 0) ); + cudaMemcpy(ply_y.data, pm.p_y,pm.num_v * sizeof(double) , cudaMemcpyHostToDevice); + ply_y.size=pm.num_v; + ply_y.valid=nullptr; + ply_y.null_count=0; + delete[] pm.p_y; + + delete[] pm.p_g_len; + } +} \ No newline at end of file diff --git a/cpp/src/io/soa/pnt_soa_reader.cu b/cpp/src/io/soa/pnt_soa_reader.cu new file mode 100644 index 000000000..c508a2293 --- /dev/null +++ b/cpp/src/io/soa/pnt_soa_reader.cu @@ -0,0 +1,77 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cuSpatial +{ + /** + * @Brief read lon/lat from file into two columns; data type is fixed to double (GDF_FLOAT64) + * + */ + void read_pnt_lonlat_soa(const char *pnt_fn,gdf_column& pnt_lon,gdf_column& pnt_lat) + { + + double * p_lon=NULL, *p_lat=NULL; + int num_p=read_point_ll(pnt_fn,p_lon,p_lat); + + pnt_lon.dtype= GDF_FLOAT64; + pnt_lon.col_name=(char *)malloc(strlen("lon")+ 1); + strcpy(pnt_lon.col_name,"lon"); + RMM_TRY( RMM_ALLOC(&pnt_lon.data, num_p * sizeof(double), 0) ); + cudaMemcpy(pnt_lon.data, p_lon,num_p * sizeof(double) , cudaMemcpyHostToDevice); + pnt_lon.size=num_p; + pnt_lon.valid=nullptr; + pnt_lon.null_count=0; + delete[] p_lon; + + pnt_lat.dtype= GDF_FLOAT64; + pnt_lat.col_name=(char *)malloc(strlen("lat")+ 1); + strcpy(pnt_lat.col_name,"lat"); + pnt_lat.data=NULL; + RMM_TRY( RMM_ALLOC(&pnt_lat.data, num_p * sizeof(double), 0) ); + cudaMemcpy(pnt_lat.data, p_lat,num_p * sizeof(double) , cudaMemcpyHostToDevice); + pnt_lat.size=num_p; + pnt_lat.valid=nullptr; + pnt_lat.null_count=0; + delete[] p_lat; + } + + /** + * @Brief read x/y from file into two columns; data type is fixed to double (GDF_FLOAT64) + * + */ + void read_pnt_xy_soa(const char *pnt_fn,gdf_column& pnt_x,gdf_column& pnt_y) + { + + double * p_x=NULL, *p_y=NULL; + int num_p=read_point_xy(pnt_fn,p_x,p_y); + + pnt_x.dtype= GDF_FLOAT64; + pnt_x.col_name=(char *)malloc(strlen("x")+ 1); + strcpy(pnt_x.col_name,"x"); + RMM_TRY( RMM_ALLOC(&pnt_x.data, num_p * sizeof(double), 0) ); + cudaMemcpy(pnt_x.data, p_x,num_p * sizeof(double) , cudaMemcpyHostToDevice); + pnt_x.size=num_p; + pnt_x.valid=nullptr; + pnt_x.null_count=0; + delete[] p_x; + + pnt_y.dtype= GDF_FLOAT64; + pnt_y.col_name=(char *)malloc(strlen("y")+ 1); + strcpy(pnt_y.col_name,"y"); + pnt_y.data=NULL; + RMM_TRY( RMM_ALLOC(&pnt_y.data, num_p * sizeof(double), 0) ); + cudaMemcpy(pnt_y.data, p_y,num_p * sizeof(double) , cudaMemcpyHostToDevice); + pnt_y.size=num_p; + pnt_y.valid=nullptr; + pnt_y.null_count=0; + delete[] p_y; + } +} \ No newline at end of file diff --git a/cpp/src/io/soa/ts_soa_reader.cu b/cpp/src/io/soa/ts_soa_reader.cu new file mode 100644 index 000000000..51d2ff658 --- /dev/null +++ b/cpp/src/io/soa/ts_soa_reader.cu @@ -0,0 +1,39 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cuSpatial +{ + /** + * @Brief read timestamp (ts: Time type) data from file as column + * + */ + void read_ts_soa(const char *ts_fn, gdf_column& ts) + { + Time * time=NULL; + size_t num_t=read_field