[REVIEW] Code review for initial check-in

cpp/CMakeLists.txt

@@ -270,6 +270,7 @@ add_library(cuspatial SHARED
             src/trajectory/trajectory_spatial_bound.cu
             src/trajectory/trajectory_distance_speed.cu
             src/trajectory/derive_trajectories.cu
+            src/trajectory/subset_trajectories.cu
             src/utility/utility.cpp)
 
 #Override RPATH for cuspatial

cpp/include/cuspatial/trajectory.hpp

@@ -22,20 +22,20 @@ namespace cuspatial {
 
 /**
  * @brief derive trajectories from points, timestamps and object ids
- * 
+ *
  * Points are x/y coordinates relative to an origin. First sorts by object id
  * and timestamp and then groups by id.
- * 
+ *
  * @param[in/out] x: x coordinates relative to a camera origin
  *                  (before/after sorting)
  * @param[in/out] y: y coordinates relative to a camera origin
  *                   (before/after sorting)
- * @param[in/out] object_id: object (e.g., vehicle) id column (before/after 
+ * @param[in/out] object_id: object (e.g., vehicle) id column (before/after
  *                sorting); upon completion, unique ids become trajectory ids
  * @param[in/out] timestamp: timestamp column (before/after sorting)
  * @param[out] trajectory_id: trajectory id column (see comments on oid)
  * @param[out] length: #of points in the derived trajectories
- * @param[out] offset: position offsets of trajectories used to index x, y, 
+ * @param[out] offset: position offsets of trajectories used to index x, y,
  *                  object_id and timestamp
  * @return number of derived trajectories
  */
@@ -86,4 +86,27 @@ void trajectory_spatial_bounds(const gdf_column& x, const gdf_column& y,
                                gdf_column& bbox_x1, gdf_column& bbox_y1,
                                gdf_column& bbox_x2, gdf_column& bbox_y2);
 
+/**
+ * @brief subset a set of trajectories based on provided trajectory IDs
+
+ * @param[in] ids: ids of trajectories whose x/y/len/pos data will be kept
+ * @param[in] in_x: input x coordinates
+ * @param[in] in_y: input y coordinates
+ * @param[in] in_id: input ids of points
+ * @param[in] in_ts: input timestamps of points
+ * @param[out] out_x: output x coordinates ordered by (in_id,in_ts)
+ * @param[out] out_y: output y coordinates ordered by (in_id,in_ts)
+ * @param[out] out_id: output ids ordered by (in_id,in_ts)
+ * @param[out] out_ts: output timestamp ordered by (in_id,in_ts)
+
+ * Note: the API is useful for integraitng with cudf and serial python APIs,
+ * e.g., query based on trajectory level informaiton using serial python APIs or cudf APIs
+ * and identify a subset of trajectory IDs.
+ * These IDs can then be used to reterive x/y/len/pos data for futher processing.
+ */
+
+uint32_t subset_trajectory_id(const gdf_column& ids,
+		const gdf_column& in_x, const gdf_column& in_y, const gdf_column& in_id, const gdf_column& in_ts,
+		gdf_column& out_x, gdf_column& out_y,gdf_column& out_id, gdf_column& out_ts);
+
 }  // namespace cuspatial

cpp/src/trajectory/derive_trajectories.cu

@@ -17,8 +17,6 @@
 #include <cudf/utilities/legacy/type_dispatcher.hpp>
 #include <utilities/cuda_utils.hpp>
 #include <type_traits>
-#include <thrust/device_vector.h>
-#include <thrust/iterator/discard_iterator.h>
 #include <sys/time.h>
 #include <time.h>
 
@@ -38,22 +36,30 @@ struct derive_trajectories_functor {
                    gdf_column& timestamp, gdf_column& trajectory_id,
                    gdf_column& length, gdf_column& offset)
     {
-        T* x_ptr = static_cast<T*>(x.data);
-        T* y_ptr = static_cast<T*>(y.data);
-        uint32_t* id_ptr = static_cast<uint32_t*>(object_id.data);
-        cuspatial::its_timestamp * time_ptr = 
-            static_cast<cuspatial::its_timestamp*>(timestamp.data);
-
+        //T* x_ptr = static_cast<T*>(x.data);
+        thrust::device_ptr<T> x_ptr =
+            thrust::device_pointer_cast(static_cast<T*>(x.data));
+        //T* y_ptr = static_cast<T*>(y.data);
+        thrust::device_ptr<T> y_ptr =
+            thrust::device_pointer_cast(static_cast<T*>(y.data));
+        //uint32_t* id_ptr = static_cast<uint32_t*>(object_id.data);
+        thrust::device_ptr<uint32_t> id_ptr =
+            thrust::device_pointer_cast(static_cast<uint32_t*>(object_id.data));
+        //its_timestamp * time_ptr = static_cast<its_timestamp*>(timestamp.data);
+        thrust::device_ptr<cuspatial::its_timestamp> time_ptr =
+            thrust::device_pointer_cast(
+                static_cast<cuspatial::its_timestamp *>(timestamp.data));
+
 #ifdef DEBUG
         int num_print = (object_id.size < 10) ? object_id.size : 10;
         std::cout<<"showing the first "<< num_print<<" input records before sort"<<std::endl;
 
         std::cout<<"x"<<std::endl;
-        thrust::copy(x_ptr,x_ptr+num_print,std::ostream_iterator<col_type>(std::cout, " "));std::cout<<std::endl;  
+        thrust::copy(x_ptr,x_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
         std::cout<<"y"<<std::endl;
-        thrust::copy(y_ptr,y_ptr+num_print,std::ostream_iterator<col_type>(std::cout, " "));std::cout<<std::endl;  
+        thrust::copy(y_ptr,y_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
 
-        std::cout<<"oid"<<std::endl;
+        std::cout<<"object id"<<std::endl;
         thrust::copy(id_ptr,id_ptr+num_print,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;  
         std::cout<<"timestamp"<<std::endl;
         thrust::copy(time_ptr,time_ptr+num_print,std::ostream_iterator<its_timestamp>(std::cout, " "));std::cout<<std::endl;    
@@ -69,37 +75,41 @@ struct derive_trajectories_functor {
             thrust::make_zip_iterator(thrust::make_tuple(time_ptr, x_ptr, y_ptr)));
 
         //allocate sufficient memory to hold id,cnt and pos before reduce_by_key
-        uint32_t *objcnt = nullptr, *objpos = nullptr, *objid = nullptr;
+        uint32_t *objcnt{nullptr};
+        uint32_t *objid{nullptr};
         RMM_TRY( RMM_ALLOC((void**)&objcnt,num_rec* sizeof(uint32_t),0) ) ; 
-        RMM_TRY( RMM_ALLOC((void**)&objpos,num_rec* sizeof(uint32_t),0) ) ; 
         RMM_TRY( RMM_ALLOC((void**)&objid,num_rec* sizeof(uint32_t),0) ) ; 
+        thrust::device_ptr<uint32_t> objcnt_ptr=thrust::device_pointer_cast(objcnt);
+        thrust::device_ptr<uint32_t> objid_ptr=thrust::device_pointer_cast(objid);        
 
         int num_traj =
             thrust::reduce_by_key(thrust::device, id_ptr, id_ptr+num_rec,
                                   thrust::constant_iterator<int>(1),
-                                  objid, objcnt).second - objcnt;
+                                  objid_ptr, objcnt_ptr).second - objcnt_ptr;
+
 
         //allocate just enough memory (num_traj), copy over and then free large (num_rec) arrays         
-        uint32_t *trajid=nullptr,*trajcnt=nullptr,*trajpos=nullptr;
+        uint32_t *trajid{nullptr};
+        uint32_t *trajcnt{nullptr};
+        uint32_t *trajpos{nullptr};
         RMM_TRY( RMM_ALLOC((void**)&trajid,  num_traj * sizeof(uint32_t),0) ) ; 
         RMM_TRY( RMM_ALLOC((void**)&trajcnt, num_traj * sizeof(uint32_t),0) ) ; 
         RMM_TRY( RMM_ALLOC((void**)&trajpos, num_traj * sizeof(uint32_t),0) ) ; 
 
+        thrust::device_ptr<uint32_t> trajid_ptr=thrust::device_pointer_cast(trajid);
         thrust::device_ptr<uint32_t> trajcnt_ptr=thrust::device_pointer_cast(trajcnt);
         thrust::device_ptr<uint32_t> trajpos_ptr=thrust::device_pointer_cast(trajpos);        
 
-        thrust::copy(objid, objid + num_traj, trajid);
-        thrust::copy(objcnt, objcnt + num_traj, trajcnt);
-        thrust::copy(objpos, objpos + num_traj, trajpos);
+        thrust::copy(objid_ptr, objid_ptr + num_traj, trajid_ptr);
+        thrust::copy(objcnt_ptr, objcnt_ptr + num_traj, trajcnt_ptr);
+        thrust::inclusive_scan(trajcnt_ptr, trajcnt_ptr + num_traj, trajpos_ptr);
 
         RMM_TRY( RMM_FREE(objid, 0) );
         RMM_TRY( RMM_FREE(objcnt, 0) );
-        RMM_TRY( RMM_FREE(objpos, 0) );
-
+
         //to avoid lost memory problem when tid/cnt/pos gdf columns are associated with dvice memory
         gdf_column_view(&trajectory_id, trajid, nullptr, num_traj, GDF_INT32);
         gdf_column_view(&length, trajcnt, nullptr, num_traj, GDF_INT32);
-        thrust::inclusive_scan(thrust::device, trajcnt_ptr, trajcnt_ptr+num_traj, trajpos_ptr);
         gdf_column_view(&offset, trajpos, nullptr, num_traj, GDF_INT32);
 
         gettimeofday(&t1, nullptr);
@@ -110,9 +120,9 @@ struct derive_trajectories_functor {
         std::cout<<"#traj="<<num_traj<<std::endl;
         std::cout<<"showing the first "<< num_print<<" records aftr sort"<<std::endl;
         std::cout<<"x"<<std::endl;
-        thrust::copy(x_ptr,x_ptr+num_print,std::ostream_iterator<col_type>(std::cout, " "));std::cout<<std::endl;  
+        thrust::copy(x_ptr,x_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
         std::cout<<"y"<<std::endl;
-        thrust::copy(y_ptr,y_ptr+num_print,std::ostream_iterator<col_type>(std::cout, " "));std::cout<<std::endl;  
+        thrust::copy(y_ptr,y_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
 
         std::cout<<"oid"<<std::endl;
         thrust::copy(id_ptr,id_ptr+num_print,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;  
@@ -122,7 +132,6 @@ struct derive_trajectories_functor {
         num_print=(num_traj<10)?num_traj:10;
         std::cout<<"showing the first "<< num_print<<" trajectory records"<<std::endl;
         std::cout<<"trajectory id"<<std::endl;
-        thrust::device_ptr<uint32_t> trajid_ptr=thrust::device_pointer_cast(trajid);
         thrust::copy(trajid_ptr,trajid_ptr+num_print,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;
         std::cout<<"trajectory #of points"<<std::endl;
         thrust::copy(trajcnt_ptr,trajcnt_ptr+num_print,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;

cpp/src/trajectory/subset_trajectories.cu

@@ -0,0 +1,164 @@
+/*
+ * Copyright (c) 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.
+ */
+
+#include <cudf/utilities/legacy/type_dispatcher.hpp>
+#include <utilities/cuda_utils.hpp>
+#include <type_traits>
+#include <thrust/device_vector.h>
+#include <thrust/binary_search.h>
+#include <thrust/iterator/discard_iterator.h>
+#include <sys/time.h>
+#include <time.h>
+
+#include <utility/utility.hpp>
+#include <utility/trajectory_thrust.cuh>
+#include <cuspatial/trajectory.hpp>
+
+using namespace std; 
+using namespace cudf;
+using namespace cuspatial;
+
+struct is_true
+{
+	__device__
+	bool operator()(const bool t)
+	{
+		return(t);
+	}
+};
+
+struct subset_functor {
+    template <typename T>
+    static constexpr bool is_supported()
+    {
+        return std::is_floating_point<T>::value;
+    }
+
+    template <typename T, std::enable_if_t< is_supported<T>() >* = nullptr>
+    uint32_t operator()(const gdf_column& ids,
+    	const gdf_column& in_x,const gdf_column& in_y, const gdf_column& in_id,const gdf_column& in_ts,
+ 	gdf_column& out_x, gdf_column& out_y, gdf_column& out_id,gdf_column& out_ts)                
+    {
+
+        //struct timeval t0,t1;
+        //gettimeofday(&t0, nullptr);
+
+	thrust::device_ptr<T> in_x_ptr = thrust::device_pointer_cast(static_cast<T*>(in_x.data));
+        thrust::device_ptr<T> in_y_ptr = thrust::device_pointer_cast(static_cast<T*>(in_y.data));
+        thrust::device_ptr<uint32_t> in_id_ptr =thrust::device_pointer_cast(static_cast<uint32_t*>(in_id.data));
+        thrust::device_ptr<its_timestamp> in_ts_ptr = thrust::device_pointer_cast(static_cast<its_timestamp*>(in_ts.data));
+
+#ifdef DEBUG
+        int num_print = (in_id.size < 10) ? in_id.size : 10;
+        std::cout<<"showing the first "<< num_print<<" input records before sort"<<std::endl;
+
+        std::cout<<"input x"<<std::endl;
+        thrust::copy(in_x_ptr,in_x_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
+        std::cout<<"input y"<<std::endl;
+        thrust::copy(in_y_ptr,in_y_ptr+num_print,std::ostream_iterator<T>(std::cout, " "));std::cout<<std::endl;  
+
+        std::cout<<"input id"<<std::endl;
+        thrust::copy(in_id_ptr,in_id_ptr+num_print,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;  
+        std::cout<<"input timestamp"<<std::endl;
+        thrust::copy(in_ts_ptr,in_ts_ptr+num_print,std::ostream_iterator<its_timestamp>(std::cout, " "));std::cout<<std::endl;    
+#endif
+
+	uint32_t num_id=ids.size;
+	uint32_t num_rec=in_id.size;
+
+	uint32_t *temp_id{nullptr};
+        RMM_TRY( RMM_ALLOC((void**)&temp_id,num_id* sizeof(uint32_t),0) ) ; 
+        thrust::device_ptr<uint32_t> temp_id_ptr=thrust::device_pointer_cast(temp_id);
+
+ 	thrust::device_ptr<uint32_t> ids_ptr=thrust::device_pointer_cast(static_cast<uint32_t*>(ids.data));
+ 	thrust::copy(ids_ptr, ids_ptr+num_id,temp_id_ptr);
+        thrust::sort(temp_id_ptr,temp_id_ptr+num_id);
+        thrust::copy(temp_id_ptr,temp_id_ptr+num_id,std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;
+
+	thrust::device_vector<bool> hit_vec(num_rec);
+	std::cout<<"beginning binary_search .............."<<std::endl;
+	thrust::binary_search(temp_id_ptr, temp_id_ptr+num_id,in_id_ptr,in_id_ptr+num_rec,hit_vec.begin());
+	std::cout<<"binary_search results.............."<<std::endl;
+	thrust::copy(hit_vec.begin(),hit_vec.end(),std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;
+	thrust::device_vector<bool> temp_vec=hit_vec;
+	thrust::copy(temp_vec.begin(),temp_vec.end(),std::ostream_iterator<uint32_t>(std::cout, " "));std::cout<<std::endl;
+
+	uint32_t num_hit=thrust::count_if(temp_vec.begin(),temp_vec.end(),is_true());
+	std::cout<<"num_hit="<<num_hit<<std::endl;
+
+	RMM_TRY( RMM_ALLOC((void**)(&(out_x.data)),num_hit* sizeof(double),0) ) ; 
+	RMM_TRY( RMM_ALLOC((void**)(&(out_y.data)),num_hit* sizeof(double),0) ) ;
+	RMM_TRY( RMM_ALLOC((void**)(&(out_id.data)),num_hit* sizeof(uint32_t),0) ) ;
+	RMM_TRY( RMM_ALLOC((void**)(&(out_ts.data)),num_hit* sizeof(its_timestamp),0) ) ;
+
+	thrust::device_ptr<T> out_x_ptr = thrust::device_pointer_cast(static_cast<T*>(out_x.data));
+        thrust::device_ptr<T> out_y_ptr = thrust::device_pointer_cast(static_cast<T*>(out_y.data));
+        thrust::device_ptr<uint32_t> out_id_ptr =thrust::device_pointer_cast(static_cast<uint32_t*>(out_id.data));
+        thrust::device_ptr<its_timestamp> out_ts_ptr = thrust::device_pointer_cast(static_cast<its_timestamp*>(out_ts.data));
+
+	auto in_itr=thrust::make_zip_iterator(thrust::make_tuple(in_x_ptr, in_y_ptr, in_id_ptr,in_ts_ptr));
+	auto out_itr=thrust::make_zip_iterator(thrust::make_tuple(out_x_ptr, out_y_ptr, out_id_ptr,out_ts_ptr));
+	uint32_t num_keep=thrust::copy_if(in_itr,in_itr+num_rec,hit_vec.begin(),out_itr,is_true())-out_itr;
+	CUDF_EXPECTS(num_hit==num_keep,"expecting count_if and copy_if return the same size");
+
+	return(num_hit);
+   }
+
+    template <typename col_type, std::enable_if_t< !is_supported<col_type>() >* = nullptr>
+    uint32_t operator()(const gdf_column& ids,
+    	const gdf_column& in_x,const gdf_column& in_y, const gdf_column& in_id,const gdf_column& in_ts,
+ 	gdf_column& out_x,gdf_column& out_y, gdf_column& out_id,gdf_column& out_ts)                
+    {
+        CUDF_FAIL("Non-floating point operation is not supported");
+    }
+};
+
+
+namespace cuspatial {
+
+/*
+
+*/
+uint32_t subset_trajectory_id(const gdf_column& ids,
+		const gdf_column& in_x, const gdf_column& in_y, const gdf_column& in_id, const gdf_column& in_ts,
+		gdf_column& out_x, gdf_column& out_y,gdf_column& out_id, gdf_column& out_ts)
+{       
+    struct timeval t0,t1;
+    gettimeofday(&t0, nullptr);
+
+    CUDF_EXPECTS(in_x.data != nullptr && in_x.data != nullptr &&
+                 in_id.data != nullptr && in_ts.data != nullptr,
+                 "x/y/in_id/in_ts data cannot be null");
+    CUDF_EXPECTS(in_x.size == in_y.size && in_x.size == in_id.size &&
+                 in_x.size == in_ts.size ,
+                 "x/y/in_id/timestamp must have equal size");
+
+    // future versions might allow x/y/in_id/timestamp to have null_count > 0,
+    // which might be useful for taking query results as inputs 
+    CUDF_EXPECTS(in_x.null_count == 0 && in_y.null_count == 0 &&
+                 in_id.null_count==0 && in_ts.null_count==0, 
+                 "NULL support unimplemented");
+
+    uint32_t num_hit=cudf::type_dispatcher(in_x.dtype, subset_functor(),ids,
+                                                 in_x, in_y, in_id, in_ts,
+                                                 out_x, out_y,out_id,out_ts);
+    std::cout<<"number of resulting points:"<<num_hit<<std::endl;
+    gettimeofday(&t1, nullptr);
+    float subset_end2end_time=calc_time("trajectory subset end-to-end time in ms=",t0,t1);
+    return num_hit;
+}
+
+}// namespace cuspatial