paleolimbot commented on code in PR #556:
URL: https://github.com/apache/sedona-db/pull/556#discussion_r2776423794
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/gpuspatial_c.h:
##########
@@ -14,60 +14,199 @@
// KIND, either express or implied. See the License for the
// specific language governing permissions and limitations
// under the License.
+#include <stdbool.h>
#include <stdint.h>
#ifdef __cplusplus
extern "C" {
#endif
-struct GpuSpatialJoinerConfig {
- uint32_t concurrency;
+struct ArrowSchema;
+struct ArrowArray;
+
+// Interfaces for ray-tracing engine (OptiX)
+struct GpuSpatialRuntimeConfig {
+ /** Path to PTX files */
const char* ptx_root;
+ /** Device ID to use, 0 is the first GPU */
+ int device_id;
+ /** Whether to use CUDA memory pool for allocations */
+ bool use_cuda_memory_pool;
+ /** Ratio of initial memory pool size to total GPU memory, between 0 and 100
*/
+ int cuda_memory_pool_init_precent;
};
-struct GpuSpatialJoinerContext {
- const char* last_error; // Pointer to std::string to store last error
message
- void* private_data; // GPUSpatial context
- void* build_indices; // Pointer to std::vector<uint32_t> to store results
- void* stream_indices;
+/** Opaque runtime for GPU spatial operations
+ * Each process should only has one instance of GpuSpatialRuntime exactly
+ */
+struct GpuSpatialRuntime {
+ /** Initialize the runtime (OptiX) with the given configuration
+ * @return 0 on success, non-zero on failure
+ */
+ int (*init)(struct GpuSpatialRuntime* self, struct GpuSpatialRuntimeConfig*
config);
+ void (*release)(struct GpuSpatialRuntime* self);
+ const char* (*get_last_error)(struct GpuSpatialRuntime* self);
+ void* private_data;
};
-enum GpuSpatialPredicate {
- GpuSpatialPredicateEquals = 0,
- GpuSpatialPredicateDisjoint,
- GpuSpatialPredicateTouches,
- GpuSpatialPredicateContains,
- GpuSpatialPredicateCovers,
- GpuSpatialPredicateIntersects,
- GpuSpatialPredicateWithin,
- GpuSpatialPredicateCoveredBy
+/** Create an instance of GpuSpatialRuntime */
+void GpuSpatialRuntimeCreate(struct GpuSpatialRuntime* runtime);
+
+struct GpuSpatialIndexConfig {
+ /** Pointer to an initialized GpuSpatialRuntime struct */
+ struct GpuSpatialRuntime* runtime;
+ /** How many threads will concurrently call Probe method */
+ uint32_t concurrency;
+};
+
+// An opaque context for concurrent probing
+struct SedonaSpatialIndexContext {
+ void* private_data;
+};
+
+struct SedonaFloatIndex2D {
+ /** Clear the spatial index, removing all built data */
+ int (*clear)(struct SedonaFloatIndex2D* self);
+ /** Create a new context for concurrent probing */
+ void (*create_context)(struct SedonaSpatialIndexContext* context);
+ /** Destroy a previously created context */
+ void (*destroy_context)(struct SedonaSpatialIndexContext* context);
+ /** Push rectangles for building the spatial index, each rectangle is
represented by 4
+ * floats: [min_x, min_y, max_x, max_y] Points can also be indexed by
providing [x, y,
+ * x, y] but points and rectangles cannot be mixed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*push_build)(struct SedonaFloatIndex2D* self, const float* buf,
uint32_t n_rects);
+ /**
+ * Finish building the spatial index after all rectangles have been pushed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*finish_building)(struct SedonaFloatIndex2D* self);
+ /**
+ * Probe the spatial index with the given rectangles, each rectangle is
represented by 4
+ * floats: [min_x, min_y, max_x, max_y] Points can also be probed by
providing [x, y, x,
+ * y] but points and rectangles cannot be mixed in one Probe call. The
results of the
+ * probe will be stored in the context.
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*probe)(struct SedonaFloatIndex2D* self, struct
SedonaSpatialIndexContext* context,
+ const float* buf, uint32_t n_rects);
+ /** Get the build indices buffer from the context
+ *
+ * @return A pointer to the buffer and its length
+ */
+ void (*get_build_indices_buffer)(struct SedonaSpatialIndexContext* context,
+ uint32_t** build_indices,
+ uint32_t* build_indices_length);
+ /** Get the probe indices buffer from the context
+ *
+ * @return A pointer to the buffer and its length
+ */
+ void (*get_probe_indices_buffer)(struct SedonaSpatialIndexContext* context,
+ uint32_t** probe_indices,
+ uint32_t* probe_indices_length);
+ /** Get the last error message from either the index
+ *
+ * @return A pointer to the error message string
+ */
+ const char* (*get_last_error)(struct SedonaFloatIndex2D* self);
+ /** Get the last error message from the context
+ *
+ * @return A pointer to the error message string
+ */
+ const char* (*context_get_last_error)(struct SedonaSpatialIndexContext*
context);
+ /** Release the spatial index and free all resources */
+ void (*release)(struct SedonaFloatIndex2D* self);
+ void* private_data;
};
-struct GpuSpatialJoiner {
- int (*init)(struct GpuSpatialJoiner* self, struct GpuSpatialJoinerConfig*
config);
- void (*clear)(struct GpuSpatialJoiner* self);
- void (*create_context)(struct GpuSpatialJoiner* self,
- struct GpuSpatialJoinerContext* context);
- void (*destroy_context)(struct GpuSpatialJoinerContext* context);
- int (*push_build)(struct GpuSpatialJoiner* self, const struct ArrowSchema*
schema,
- const struct ArrowArray* array, int64_t offset, int64_t
length);
- int (*finish_building)(struct GpuSpatialJoiner* self);
- int (*push_stream)(struct GpuSpatialJoiner* self,
- struct GpuSpatialJoinerContext* context,
- const struct ArrowSchema* schema, const struct
ArrowArray* array,
- int64_t offset, int64_t length, enum GpuSpatialPredicate
predicate,
- int32_t array_index_offset);
- void (*get_build_indices_buffer)(struct GpuSpatialJoinerContext* context,
- void** build_indices, uint32_t*
build_indices_length);
- void (*get_stream_indices_buffer)(struct GpuSpatialJoinerContext* context,
- void** stream_indices,
- uint32_t* stream_indices_length);
- void (*release)(struct GpuSpatialJoiner* self);
+/** Create an instance of GpuSpatialIndex for 2D float rectangles/points
+ * @return 0 on success, non-zero on failure
+ */
+int GpuSpatialIndexFloat2DCreate(struct SedonaFloatIndex2D* index,
+ const struct GpuSpatialIndexConfig* config);
+
+struct GpuSpatialRefinerConfig {
+ /** Pointer to an initialized GpuSpatialRuntime struct */
+ struct GpuSpatialRuntime* runtime;
+ /** How many threads will concurrently call Probe method */
+ uint32_t concurrency;
+ /** Whether to compress the BVH structures to save memory */
+ bool compress_bvh;
+ /** Number of batches to pipeline for parsing and refinement; setting to 1
disables
+ * pipelining */
+ uint32_t pipeline_batches;
+};
+
+enum SedonaSpatialRelationPredicate {
+ SedonaSpatialPredicateEquals = 0,
+ SedonaSpatialPredicateDisjoint,
+ SedonaSpatialPredicateTouches,
+ SedonaSpatialPredicateContains,
+ SedonaSpatialPredicateCovers,
+ SedonaSpatialPredicateIntersects,
+ SedonaSpatialPredicateWithin,
+ SedonaSpatialPredicateCoveredBy
+};
+
+/** An opaque spatial refiner that can refine candidate pairs of geometries */
+struct SedonaSpatialRefiner {
+ /** Clear all built geometries from the refiner */
+ int (*clear)(struct SedonaSpatialRefiner* self);
+
+ /** Push geometries for building the spatial refiner
+ *
+ * @param build_schema The Arrow schema of the build geometries
+ * @param build_array The Arrow array of the build geometries
+ * @return 0 on success, non-zero on failure
+ */
+ int (*push_build)(struct SedonaSpatialRefiner* self,
+ const struct ArrowSchema* build_schema,
+ const struct ArrowArray* build_array);
+ /**
+ * Finish building the spatial refiner after all geometries have been pushed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*finish_building)(struct SedonaSpatialRefiner* self);
+
+ /**
+ * Refine candidate pairs of geometries
+ *
+ * @param probe_schema The Arrow schema of the probe geometries
+ * @param probe_array The Arrow array of the probe geometries
+ * @param predicate The spatial relation predicate to evaluate
+ * @param build_indices The build indices of candidate pairs
+ * @param probe_indices The probe indices of candidate pairs
+ * @param indices_size The number of candidate pairs
+ * @param new_indices_size Output parameter to store the number of refined
pairs
+ * @return 0 on success, non-zero on failure
+ */
+ int (*refine)(struct SedonaSpatialRefiner* self, const struct ArrowSchema*
probe_schema,
+ const struct ArrowArray* probe_array,
+ enum SedonaSpatialRelationPredicate predicate, uint32_t*
build_indices,
+ uint32_t* probe_indices, uint32_t indices_size,
+ uint32_t* new_indices_size);
Review Comment:
It is worth clarifying here that build indices are global and probe indices
are local, that build_indices and probe_indices and new_indices_size are output
arguments, and what happens when there are more indices than `indices_size`.
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/loader/parallel_wkb_loader.hpp:
##########
@@ -105,6 +66,7 @@ template <typename POINT_T, typename INDEX_T>
struct HostParsedGeometries {
constexpr static int n_dim = POINT_T::n_dim;
using mbr_t = Box<Point<float, n_dim>>;
+ GeometryType type; // A general type that can reprs
Review Comment:
```suggestion
GeometryType type;
```
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/gpuspatial_c.h:
##########
@@ -14,60 +14,199 @@
// KIND, either express or implied. See the License for the
// specific language governing permissions and limitations
// under the License.
+#include <stdbool.h>
#include <stdint.h>
#ifdef __cplusplus
extern "C" {
#endif
-struct GpuSpatialJoinerConfig {
- uint32_t concurrency;
+struct ArrowSchema;
+struct ArrowArray;
+
+// Interfaces for ray-tracing engine (OptiX)
+struct GpuSpatialRuntimeConfig {
+ /** Path to PTX files */
const char* ptx_root;
+ /** Device ID to use, 0 is the first GPU */
+ int device_id;
+ /** Whether to use CUDA memory pool for allocations */
+ bool use_cuda_memory_pool;
+ /** Ratio of initial memory pool size to total GPU memory, between 0 and 100
*/
+ int cuda_memory_pool_init_precent;
};
-struct GpuSpatialJoinerContext {
- const char* last_error; // Pointer to std::string to store last error
message
- void* private_data; // GPUSpatial context
- void* build_indices; // Pointer to std::vector<uint32_t> to store results
- void* stream_indices;
+/** Opaque runtime for GPU spatial operations
+ * Each process should only has one instance of GpuSpatialRuntime exactly
+ */
+struct GpuSpatialRuntime {
+ /** Initialize the runtime (OptiX) with the given configuration
+ * @return 0 on success, non-zero on failure
+ */
+ int (*init)(struct GpuSpatialRuntime* self, struct GpuSpatialRuntimeConfig*
config);
+ void (*release)(struct GpuSpatialRuntime* self);
+ const char* (*get_last_error)(struct GpuSpatialRuntime* self);
+ void* private_data;
};
-enum GpuSpatialPredicate {
- GpuSpatialPredicateEquals = 0,
- GpuSpatialPredicateDisjoint,
- GpuSpatialPredicateTouches,
- GpuSpatialPredicateContains,
- GpuSpatialPredicateCovers,
- GpuSpatialPredicateIntersects,
- GpuSpatialPredicateWithin,
- GpuSpatialPredicateCoveredBy
+/** Create an instance of GpuSpatialRuntime */
+void GpuSpatialRuntimeCreate(struct GpuSpatialRuntime* runtime);
+
+struct GpuSpatialIndexConfig {
+ /** Pointer to an initialized GpuSpatialRuntime struct */
+ struct GpuSpatialRuntime* runtime;
+ /** How many threads will concurrently call Probe method */
+ uint32_t concurrency;
+};
+
+// An opaque context for concurrent probing
+struct SedonaSpatialIndexContext {
+ void* private_data;
+};
+
+struct SedonaFloatIndex2D {
+ /** Clear the spatial index, removing all built data */
+ int (*clear)(struct SedonaFloatIndex2D* self);
+ /** Create a new context for concurrent probing */
+ void (*create_context)(struct SedonaSpatialIndexContext* context);
+ /** Destroy a previously created context */
+ void (*destroy_context)(struct SedonaSpatialIndexContext* context);
+ /** Push rectangles for building the spatial index, each rectangle is
represented by 4
+ * floats: [min_x, min_y, max_x, max_y] Points can also be indexed by
providing [x, y,
+ * x, y] but points and rectangles cannot be mixed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*push_build)(struct SedonaFloatIndex2D* self, const float* buf,
uint32_t n_rects);
Review Comment:
I am not sure I understand the comment about mixing points and rectangles.
Are points specified by only two `float`s? Or are degenerate rectangles like
(0, 1, 0, 1) not valid?
##########
c/sedona-libgpuspatial/libgpuspatial/test/refiner_test.cu:
##########
@@ -0,0 +1,356 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 "array_stream.hpp"
+#include "test_common.hpp"
+
+#include "gpuspatial/index/rt_spatial_index.hpp"
+#include "gpuspatial/loader/device_geometries.hpp"
+#include "gpuspatial/refine/rt_spatial_refiner.hpp"
+
+#include "nanoarrow/nanoarrow.hpp"
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+#include "gpuspatial/index/rt_spatial_index.cuh"
+#include "gpuspatial/refine/rt_spatial_refiner.cuh"
+
+namespace gpuspatial {
+
+void TestJoiner(ArrowSchema* build_schema, std::vector<ArrowArray*>&
build_arrays,
+ ArrowSchema* probe_schema, std::vector<ArrowArray*>&
probe_arrays,
+ Predicate predicate, bool pipelined = false) {
+ using namespace TestUtils;
+ using coord_t = double;
+ using fpoint_t = Point<coord_t, 2>;
+ using box_t = Box<fpoint_t>;
+
+ auto rt_engine = std::make_shared<RTEngine>();
+ {
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+ auto config = get_default_rt_config(ptx_root);
+ rt_engine->Init(config);
+ }
+
+ RTSpatialIndexConfig idx_config;
+ idx_config.rt_engine = rt_engine;
+ auto rt_index = CreateRTSpatialIndex<coord_t, 2>(idx_config);
+
+ RTSpatialRefinerConfig refiner_config;
+ refiner_config.rt_engine = rt_engine;
+ if (pipelined) {
+ refiner_config.pipeline_batches = 10;
+ }
+ auto rt_refiner = CreateRTSpatialRefiner(refiner_config);
+
+ // Initialize GEOS C++ components
+ auto geos_factory = geos::geom::GeometryFactory::create();
+ geos::io::WKBReader wkb_reader(*geos_factory);
+ geos::index::strtree::STRtree tree(10);
+
+ // Storage for GEOS geometries to ensure they outlive the tree
+ // The STRtree stores raw pointers, so we must own the objects
+ std::vector<std::unique_ptr<geos::geom::Geometry>> build_geoms_storage;
+ size_t total_build_length = 0;
+ for (auto& array : build_arrays) {
+ total_build_length += array->length;
+ }
+ build_geoms_storage.reserve(total_build_length);
+
+ size_t tail_build = 0;
+ ArrowError error;
+
+ // --- Build Phase ---
+ for (auto& array : build_arrays) {
+ nanoarrow::UniqueArrayView array_view;
+ ASSERT_EQ(ArrowArrayViewInitFromSchema(array_view.get(), build_schema,
&error),
+ NANOARROW_OK)
+ << error.message;
+ ASSERT_EQ(ArrowArrayViewSetArray(array_view.get(), array, &error),
NANOARROW_OK)
+ << error.message;
+
+ std::vector<box_t> rects;
+ rects.reserve(array->length);
+
+ for (int64_t i = 0; i < array->length; i++) {
+ // Parse WKB
+ ArrowStringView wkb_view =
ArrowArrayViewGetStringUnsafe(array_view.get(), i);
+ // Copy the view to a buffer because WKBReader reads from istream or
byte array
+ // We can cast directly if the underlying type allows
+ std::stringstream iss;
+ // WKBReader::readHEX parses hex strings, read parses raw bytes
+ // Arrow WKB is usually raw bytes (binary)
Review Comment:
```suggestion
```
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/gpuspatial_c.h:
##########
@@ -14,60 +14,199 @@
// KIND, either express or implied. See the License for the
// specific language governing permissions and limitations
// under the License.
+#include <stdbool.h>
#include <stdint.h>
#ifdef __cplusplus
extern "C" {
#endif
-struct GpuSpatialJoinerConfig {
- uint32_t concurrency;
+struct ArrowSchema;
+struct ArrowArray;
+
+// Interfaces for ray-tracing engine (OptiX)
+struct GpuSpatialRuntimeConfig {
+ /** Path to PTX files */
const char* ptx_root;
+ /** Device ID to use, 0 is the first GPU */
+ int device_id;
+ /** Whether to use CUDA memory pool for allocations */
+ bool use_cuda_memory_pool;
+ /** Ratio of initial memory pool size to total GPU memory, between 0 and 100
*/
+ int cuda_memory_pool_init_precent;
};
-struct GpuSpatialJoinerContext {
- const char* last_error; // Pointer to std::string to store last error
message
- void* private_data; // GPUSpatial context
- void* build_indices; // Pointer to std::vector<uint32_t> to store results
- void* stream_indices;
+/** Opaque runtime for GPU spatial operations
+ * Each process should only has one instance of GpuSpatialRuntime exactly
Review Comment:
```suggestion
* Each process should have exactly one instance of GpuSpatialRuntime
```
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/refine/spatial_refiner.hpp:
##########
@@ -0,0 +1,45 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 "gpuspatial/relate/predicate.hpp"
+
+#include "nanoarrow/nanoarrow.h"
+
+namespace gpuspatial {
+/** This class refines candidate pairs of geometries based on a spatial
predicate.
+ *
+ * The SpatialRefiner is initialized by pushing build-side geometries via
PushBuild(),
+ * followed by a call to FinishBuilding(). After that, the Refine() method can
be called
+ * multiple times with probe-side geometries and candidate index pairs to
filter out
+ * non-matching pairs based on the specified spatial predicate.
+ */
+class SpatialRefiner {
+ public:
+ virtual ~SpatialRefiner() = default;
+
+ virtual void Clear() = 0;
+
+ virtual void PushBuild(const ArrowArrayView* build_array) = 0;
+
+ virtual void FinishBuilding() = 0;
+
+ virtual uint32_t Refine(const ArrowArrayView* probe_array, Predicate
predicate,
+ uint32_t* build_indices, uint32_t* probe_indices,
+ uint32_t len) = 0;
Review Comment:
Can you also add docstrings to these functions since this is one of the main
APIs? In particular that `build_indices` and `probe_indices` are output
parameters and what `len` is referring to (is `probe_array->length`
sufficient?).
I think one of my earlier comments was not correct...here, `build_indices`
is global and not local, correct? This implies the build side can never be
larger than 4 billion rows (which is probably fine given the current memory
limitations of GPUs). All of this is fine, but these parameters do need to be
documented.
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/gpuspatial_c.h:
##########
@@ -14,60 +14,199 @@
// KIND, either express or implied. See the License for the
// specific language governing permissions and limitations
// under the License.
+#include <stdbool.h>
#include <stdint.h>
#ifdef __cplusplus
extern "C" {
#endif
-struct GpuSpatialJoinerConfig {
- uint32_t concurrency;
+struct ArrowSchema;
+struct ArrowArray;
+
+// Interfaces for ray-tracing engine (OptiX)
+struct GpuSpatialRuntimeConfig {
+ /** Path to PTX files */
const char* ptx_root;
+ /** Device ID to use, 0 is the first GPU */
+ int device_id;
+ /** Whether to use CUDA memory pool for allocations */
+ bool use_cuda_memory_pool;
+ /** Ratio of initial memory pool size to total GPU memory, between 0 and 100
*/
+ int cuda_memory_pool_init_precent;
};
-struct GpuSpatialJoinerContext {
- const char* last_error; // Pointer to std::string to store last error
message
- void* private_data; // GPUSpatial context
- void* build_indices; // Pointer to std::vector<uint32_t> to store results
- void* stream_indices;
+/** Opaque runtime for GPU spatial operations
+ * Each process should only has one instance of GpuSpatialRuntime exactly
+ */
+struct GpuSpatialRuntime {
+ /** Initialize the runtime (OptiX) with the given configuration
+ * @return 0 on success, non-zero on failure
+ */
+ int (*init)(struct GpuSpatialRuntime* self, struct GpuSpatialRuntimeConfig*
config);
+ void (*release)(struct GpuSpatialRuntime* self);
+ const char* (*get_last_error)(struct GpuSpatialRuntime* self);
+ void* private_data;
};
-enum GpuSpatialPredicate {
- GpuSpatialPredicateEquals = 0,
- GpuSpatialPredicateDisjoint,
- GpuSpatialPredicateTouches,
- GpuSpatialPredicateContains,
- GpuSpatialPredicateCovers,
- GpuSpatialPredicateIntersects,
- GpuSpatialPredicateWithin,
- GpuSpatialPredicateCoveredBy
+/** Create an instance of GpuSpatialRuntime */
+void GpuSpatialRuntimeCreate(struct GpuSpatialRuntime* runtime);
+
+struct GpuSpatialIndexConfig {
+ /** Pointer to an initialized GpuSpatialRuntime struct */
+ struct GpuSpatialRuntime* runtime;
+ /** How many threads will concurrently call Probe method */
+ uint32_t concurrency;
+};
+
+// An opaque context for concurrent probing
+struct SedonaSpatialIndexContext {
+ void* private_data;
+};
+
+struct SedonaFloatIndex2D {
+ /** Clear the spatial index, removing all built data */
+ int (*clear)(struct SedonaFloatIndex2D* self);
+ /** Create a new context for concurrent probing */
+ void (*create_context)(struct SedonaSpatialIndexContext* context);
+ /** Destroy a previously created context */
+ void (*destroy_context)(struct SedonaSpatialIndexContext* context);
+ /** Push rectangles for building the spatial index, each rectangle is
represented by 4
+ * floats: [min_x, min_y, max_x, max_y] Points can also be indexed by
providing [x, y,
+ * x, y] but points and rectangles cannot be mixed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*push_build)(struct SedonaFloatIndex2D* self, const float* buf,
uint32_t n_rects);
+ /**
+ * Finish building the spatial index after all rectangles have been pushed
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*finish_building)(struct SedonaFloatIndex2D* self);
+ /**
+ * Probe the spatial index with the given rectangles, each rectangle is
represented by 4
+ * floats: [min_x, min_y, max_x, max_y] Points can also be probed by
providing [x, y, x,
+ * y] but points and rectangles cannot be mixed in one Probe call. The
results of the
+ * probe will be stored in the context.
+ *
+ * @return 0 on success, non-zero on failure
+ */
+ int (*probe)(struct SedonaFloatIndex2D* self, struct
SedonaSpatialIndexContext* context,
+ const float* buf, uint32_t n_rects);
+ /** Get the build indices buffer from the context
+ *
+ * @return A pointer to the buffer and its length
+ */
+ void (*get_build_indices_buffer)(struct SedonaSpatialIndexContext* context,
+ uint32_t** build_indices,
+ uint32_t* build_indices_length);
+ /** Get the probe indices buffer from the context
+ *
+ * @return A pointer to the buffer and its length
+ */
+ void (*get_probe_indices_buffer)(struct SedonaSpatialIndexContext* context,
+ uint32_t** probe_indices,
+ uint32_t* probe_indices_length);
Review Comment:
In the `refine` callback the build and probe indices are returned directly
from the callback into caller-allocated memory. Is it possible to use the same
pattern for both of these interfaces?
(Also worth clarifying the local vs. globalness of the build indices here)
##########
c/sedona-libgpuspatial/libgpuspatial/test/test_common.hpp:
##########
@@ -91,6 +102,139 @@ std::string GetCanonicalPath(const std::string&
relative_path_str) {
}
}
+// Helper to evaluate predicates using GEOS C++ API
+static bool EvaluateGeosPredicate(gpuspatial::Predicate predicate,
+ const geos::geom::Geometry* geom1,
+ const geos::geom::Geometry* geom2) {
+ switch (predicate) {
+ case gpuspatial::Predicate::kContains:
+ return geom1->contains(geom2);
+ case gpuspatial::Predicate::kIntersects:
+ return geom1->intersects(geom2);
+ case gpuspatial::Predicate::kWithin:
+ return geom1->within(geom2);
+ case gpuspatial::Predicate::kEquals:
+ return geom1->equals(geom2);
+ case gpuspatial::Predicate::kTouches:
+ return geom1->touches(geom2);
+ default:
+ throw std::out_of_range("Unsupported GEOS predicate enumeration value.");
+ }
+}
+
+// Helper structure to keep visitor context
+struct JoinVisitorContext {
+ const geos::geom::Geometry* probe_geom;
+ std::vector<uint32_t>* build_indices;
+ std::vector<uint32_t>* probe_indices;
+ size_t current_probe_index;
+ gpuspatial::Predicate predicate;
+};
+
+// GEOS Visitor Implementation
+class JoinVisitor : public geos::index::ItemVisitor {
+ public:
+ JoinVisitorContext* ctx;
+ explicit JoinVisitor(JoinVisitorContext* c) : ctx(c) {}
+
+ void visitItem(void* item) override {
+ const auto* build_geom = static_cast<const geos::geom::Geometry*>(item);
+
+ // Use the existing predicate evaluator from your TestUtils
Review Comment:
```suggestion
// Use the existing predicate evaluator from TestUtils
```
##########
c/sedona-libgpuspatial/libgpuspatial/test/refiner_test.cu:
##########
@@ -0,0 +1,356 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 "array_stream.hpp"
+#include "test_common.hpp"
+
+#include "gpuspatial/index/rt_spatial_index.hpp"
+#include "gpuspatial/loader/device_geometries.hpp"
+#include "gpuspatial/refine/rt_spatial_refiner.hpp"
+
+#include "nanoarrow/nanoarrow.hpp"
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+#include "gpuspatial/index/rt_spatial_index.cuh"
+#include "gpuspatial/refine/rt_spatial_refiner.cuh"
+
+namespace gpuspatial {
+
+void TestJoiner(ArrowSchema* build_schema, std::vector<ArrowArray*>&
build_arrays,
+ ArrowSchema* probe_schema, std::vector<ArrowArray*>&
probe_arrays,
+ Predicate predicate, bool pipelined = false) {
+ using namespace TestUtils;
+ using coord_t = double;
+ using fpoint_t = Point<coord_t, 2>;
+ using box_t = Box<fpoint_t>;
+
+ auto rt_engine = std::make_shared<RTEngine>();
+ {
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+ auto config = get_default_rt_config(ptx_root);
+ rt_engine->Init(config);
+ }
+
+ RTSpatialIndexConfig idx_config;
+ idx_config.rt_engine = rt_engine;
+ auto rt_index = CreateRTSpatialIndex<coord_t, 2>(idx_config);
+
+ RTSpatialRefinerConfig refiner_config;
+ refiner_config.rt_engine = rt_engine;
+ if (pipelined) {
+ refiner_config.pipeline_batches = 10;
+ }
+ auto rt_refiner = CreateRTSpatialRefiner(refiner_config);
+
+ // Initialize GEOS C++ components
+ auto geos_factory = geos::geom::GeometryFactory::create();
+ geos::io::WKBReader wkb_reader(*geos_factory);
+ geos::index::strtree::STRtree tree(10);
+
+ // Storage for GEOS geometries to ensure they outlive the tree
+ // The STRtree stores raw pointers, so we must own the objects
+ std::vector<std::unique_ptr<geos::geom::Geometry>> build_geoms_storage;
+ size_t total_build_length = 0;
+ for (auto& array : build_arrays) {
+ total_build_length += array->length;
+ }
+ build_geoms_storage.reserve(total_build_length);
+
+ size_t tail_build = 0;
+ ArrowError error;
+
+ // --- Build Phase ---
+ for (auto& array : build_arrays) {
+ nanoarrow::UniqueArrayView array_view;
+ ASSERT_EQ(ArrowArrayViewInitFromSchema(array_view.get(), build_schema,
&error),
+ NANOARROW_OK)
+ << error.message;
+ ASSERT_EQ(ArrowArrayViewSetArray(array_view.get(), array, &error),
NANOARROW_OK)
+ << error.message;
+
+ std::vector<box_t> rects;
+ rects.reserve(array->length);
+
+ for (int64_t i = 0; i < array->length; i++) {
+ // Parse WKB
+ ArrowStringView wkb_view =
ArrowArrayViewGetStringUnsafe(array_view.get(), i);
+ // Copy the view to a buffer because WKBReader reads from istream or
byte array
+ // We can cast directly if the underlying type allows
+ std::stringstream iss;
+ // WKBReader::readHEX parses hex strings, read parses raw bytes
+ // Arrow WKB is usually raw bytes (binary)
+ auto geom = wkb_reader.read(reinterpret_cast<const unsigned
char*>(wkb_view.data),
+ wkb_view.size_bytes);
+
+ // Calculate Envelope for GPU Index
+ const geos::geom::Envelope* env = geom->getEnvelopeInternal();
+
+ double xmin = 0, ymin = 0, xmax = -1, ymax = -1;
+ if (!env->isNull()) {
+ xmin = env->getMinX();
+ ymin = env->getMinY();
+ xmax = env->getMaxX();
+ ymax = env->getMaxY();
+ }
+
+ box_t bbox(fpoint_t((float)xmin, (float)ymin), fpoint_t((float)xmax,
(float)ymax));
+ rects.push_back(bbox);
+
+ // Store User Data (global offset)
+ size_t global_offset = tail_build + i;
+ geom->setUserData((void*)global_offset);
+
+ // Insert into GEOS STRtree
+ tree.insert(env, geom.get());
+
+ // Transfer ownership to storage vector
+ build_geoms_storage.push_back(std::move(geom));
+ }
+
+ rt_index->PushBuild(rects.data(), rects.size());
+ tail_build += array->length;
+
+ rt_refiner->PushBuild(array_view.get());
+ }
+
+ rt_index->FinishBuilding();
+ rt_refiner->FinishBuilding();
+
+ // --- Probe Phase ---
+ for (auto& probe_array : probe_arrays) {
+ nanoarrow::UniqueArrayView probe_view;
+ ASSERT_EQ(ArrowArrayViewInitFromSchema(probe_view.get(), probe_schema,
&error),
+ NANOARROW_OK)
+ << error.message;
+ ASSERT_EQ(ArrowArrayViewSetArray(probe_view.get(), probe_array, &error),
NANOARROW_OK)
+ << error.message;
+
+ std::vector<box_t> queries;
+ std::vector<std::unique_ptr<geos::geom::Geometry>> probe_geoms;
+ probe_geoms.reserve(probe_array->length);
+
+ for (int64_t i = 0; i < probe_array->length; i++) {
+ ArrowStringView wkb_view =
ArrowArrayViewGetStringUnsafe(probe_view.get(), i);
Review Comment:
```suggestion
ArrowBufferView wkb_view =
ArrowArrayViewGetBytesUnsafe(probe_view.get(), i);
```
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/loader/parallel_wkb_loader.hpp:
##########
@@ -531,9 +548,7 @@ class ParallelWkbLoader {
public:
struct Config {
- // How many rows of WKBs to process in one chunk
- // This value affects the peak memory usage and overheads
- int chunk_size = 16 * 1024;
+ float memory_quota = 0.8f; // percentage of free memory to use
Review Comment:
Just making sure this value is updated from the input config somewhere (I
think I just missed where this happens)
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/loader/parallel_wkb_loader.hpp:
##########
@@ -543,82 +558,94 @@ class ParallelWkbLoader {
: thread_pool_(thread_pool) {}
void Init(const Config& config = Config()) {
- ArrowArrayViewInitFromType(&array_view_, NANOARROW_TYPE_BINARY);
config_ = config;
- geometry_type_ = GeometryType::kNull;
+ Clear(rmm::cuda_stream_default);
}
void Clear(rmm::cuda_stream_view stream) {
geometry_type_ = GeometryType::kNull;
geoms_.Clear(stream);
}
- void Parse(rmm::cuda_stream_view stream, const ArrowArray* array, int64_t
offset,
+ void Parse(rmm::cuda_stream_view stream, const ArrowArrayView* array,
int64_t offset,
int64_t length) {
+ auto begin = thrust::make_counting_iterator<int64_t>(offset);
+ auto end = begin + length;
+
+ Parse(stream, array, begin, end);
+ }
+
+ template <typename OFFSET_IT>
+ void Parse(rmm::cuda_stream_view stream, const ArrowArrayView* array_view,
+ OFFSET_IT begin, OFFSET_IT end) {
using host_geometries_t = detail::HostParsedGeometries<POINT_T, INDEX_T>;
- ArrowError arrow_error;
- if (ArrowArrayViewSetArray(&array_view_, array, &arrow_error) !=
NANOARROW_OK) {
- throw std::runtime_error("ArrowArrayViewSetArray error " +
- std::string(arrow_error.message));
- }
+
+ size_t num_offsets = std::distance(begin, end);
+ if (num_offsets == 0) return;
+
auto parallelism = thread_pool_->num_threads();
- auto est_bytes = estimateTotalBytes(array, offset, length);
- auto free_memory = detail::get_free_physical_memory_linux();
+ uint64_t est_bytes = estimateTotalBytes(array_view, begin, end);
+
+ uint64_t free_memory = MemoryManager::get_available_host_memory();
+ uint64_t memory_quota = free_memory * config_.memory_quota;
Review Comment:
Not for this PR but I think eventually we'll have to specify a host memory
quote in MB based on some accounting in the RuntimeEnv of DataFusion.
##########
c/sedona-libgpuspatial/libgpuspatial/include/gpuspatial/refine/rt_spatial_refiner.hpp:
##########
@@ -0,0 +1,54 @@
+// Licensed to the Apache Software Foundation (ASF) under one
+// or more contributor license agreements. See the NOTICE file
+// distributed with this work for additional information
+// regarding copyright ownership. The ASF licenses this file
+// to you 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 "gpuspatial/refine/spatial_refiner.hpp"
+#include "gpuspatial/rt/rt_engine.hpp"
+
+#include <memory>
+
+namespace gpuspatial {
+/** Configuration for RTSpatialRefiner for initializing a RTSpatialRefiner
instance
+ */
+struct RTSpatialRefinerConfig {
+ // The ray-tracing engine to use
+ std::shared_ptr<RTEngine> rt_engine;
+ // Prefer fast build the BVH
+ bool prefer_fast_build = false;
+ // Compress the BVH to save memory
+ bool compact = true;
+ // Loader configurations
+ // How many threads to use for parsing WKBs
+ uint32_t parsing_threads = std::thread::hardware_concurrency();
+ // How many threads are allowed to call PushStream concurrently
+ uint32_t concurrency = 1;
+ // Overlapping parsing and refinement by pipelining multiple batches; 1
means no
+ // pipelining
+ uint32_t pipeline_batches = 1;
+ // the host memory quota for WKB parser compared to the available memory
+ float wkb_parser_memory_quota = 0.8;
+ // the device memory quota for relate engine compared to the available memory
+ float relate_engine_memory_quota = 0.8;
Review Comment:
Optional, but these are probably more effective to tweak from SedonaDB as
memory quotas in MB (since SedonaDB has its own memory accounting).
##########
c/sedona-libgpuspatial/libgpuspatial/test/c_wrapper_test.cc:
##########
@@ -15,92 +15,304 @@
// specific language governing permissions and limitations
// under the License.
+#include "array_stream.hpp"
#include "test_common.hpp"
#include "gpuspatial/gpuspatial_c.h"
#include <gmock/gmock.h>
#include <gtest/gtest.h>
+
+#include "nanoarrow/nanoarrow.hpp"
+
#include <random>
#include <vector>
-#include "array_stream.hpp"
-#include "nanoarrow/nanoarrow.hpp"
-namespace TestUtils {
-std::string GetTestDataPath(const std::string& relative_path_to_file);
+TEST(RuntimeTest, InitializeRuntime) {
+ GpuSpatialRuntime runtime;
+ GpuSpatialRuntimeCreate(&runtime);
+ GpuSpatialRuntimeConfig config;
+
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+ config.ptx_root = ptx_root.c_str();
+ config.device_id = 0;
+ config.use_cuda_memory_pool = false;
+ ASSERT_EQ(runtime.init(&runtime, &config), 0);
+
+ runtime.release(&runtime);
+}
+
+TEST(RuntimeTest, ErrorTest) {
+ GpuSpatialRuntime runtime;
+ GpuSpatialRuntimeCreate(&runtime);
+ GpuSpatialRuntimeConfig runtime_config;
+
+ runtime_config.ptx_root = "/invalid/path/to/ptx";
+ runtime_config.device_id = 0;
+ runtime_config.use_cuda_memory_pool = false;
+
+ EXPECT_NE(runtime.init(&runtime, &runtime_config), 0);
+
+ const char* raw_error = runtime.get_last_error(&runtime);
+ printf("Error received: %s\n", raw_error);
+
+ std::string error_msg(raw_error);
+
+ EXPECT_NE(error_msg.find("No such file or directory"), std::string::npos)
+ << "Error message was corrupted or incorrect. Got: " << error_msg;
+
+ runtime.release(&runtime);
+}
+
+TEST(SpatialIndexTest, InitializeIndex) {
+ GpuSpatialRuntime runtime;
+ GpuSpatialRuntimeCreate(&runtime);
+ GpuSpatialRuntimeConfig runtime_config;
+
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+ runtime_config.ptx_root = ptx_root.c_str();
+ runtime_config.device_id = 0;
+ runtime_config.use_cuda_memory_pool = true;
+ runtime_config.cuda_memory_pool_init_precent = 10;
+ ASSERT_EQ(runtime.init(&runtime, &runtime_config), 0);
+
+ SedonaFloatIndex2D index;
+ GpuSpatialIndexConfig index_config;
+
+ index_config.runtime = &runtime;
+ index_config.concurrency = 1;
+
+ ASSERT_EQ(GpuSpatialIndexFloat2DCreate(&index, &index_config), 0);
+
+ index.release(&index);
+ runtime.release(&runtime);
+}
+
+TEST(RefinerTest, InitializeRefiner) {
+ GpuSpatialRuntime runtime;
+ GpuSpatialRuntimeCreate(&runtime);
+ GpuSpatialRuntimeConfig runtime_config;
+
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+ runtime_config.ptx_root = ptx_root.c_str();
+ runtime_config.device_id = 0;
+ runtime_config.use_cuda_memory_pool = true;
+ runtime_config.cuda_memory_pool_init_precent = 10;
+ ASSERT_EQ(runtime.init(&runtime, &runtime_config), 0);
+
+ SedonaSpatialRefiner refiner;
+ GpuSpatialRefinerConfig refiner_config;
+
+ refiner_config.runtime = &runtime;
+ refiner_config.concurrency = 1;
+
+ ASSERT_EQ(GpuSpatialRefinerCreate(&refiner, &refiner_config), 0);
+
+ refiner.release(&refiner);
+ runtime.release(&runtime);
}
class CWrapperTest : public ::testing::Test {
protected:
void SetUp() override {
- // Initialize the GpuSpatialJoiner
- GpuSpatialJoinerCreate(&joiner_);
- struct GpuSpatialJoinerConfig config_;
- std::string ptx_root = TestUtils::GetTestDataPath("shaders_ptx");
+ std::string ptx_root = TestUtils::GetTestShaderPath();
+
+ GpuSpatialRuntimeCreate(&runtime_);
+ GpuSpatialRuntimeConfig runtime_config;
- // Set up the configuration
- config_.concurrency = 2; // Example concurrency level
- config_.ptx_root = ptx_root.c_str();
+ runtime_config.ptx_root = ptx_root.c_str();
+ runtime_config.device_id = 0;
+ runtime_config.use_cuda_memory_pool = true;
+ runtime_config.cuda_memory_pool_init_precent = 10;
+ ASSERT_EQ(runtime_.init(&runtime_, &runtime_config), 0);
- ASSERT_EQ(joiner_.init(&joiner_, &config_), 0);
- // Initialize the context
+ GpuSpatialIndexConfig index_config;
+
+ index_config.runtime = &runtime_;
+ index_config.concurrency = 1;
+
+ ASSERT_EQ(GpuSpatialIndexFloat2DCreate(&index_, &index_config), 0);
+
+ GpuSpatialRefinerConfig refiner_config;
+
+ refiner_config.runtime = &runtime_;
+ refiner_config.concurrency = 1;
+ refiner_config.compress_bvh = false;
+ refiner_config.pipeline_batches = 1;
+
+ ASSERT_EQ(GpuSpatialRefinerCreate(&refiner_, &refiner_config), 0);
}
void TearDown() override {
- // Clean up
- joiner_.release(&joiner_);
+ refiner_.release(&refiner_);
+ index_.release(&index_);
+ runtime_.release(&runtime_);
}
-
- struct GpuSpatialJoiner joiner_;
+ GpuSpatialRuntime runtime_;
+ SedonaFloatIndex2D index_;
+ SedonaSpatialRefiner refiner_;
};
TEST_F(CWrapperTest, InitializeJoiner) {
- // Test if the joiner initializes correctly
- struct GpuSpatialJoinerContext context_;
- joiner_.create_context(&joiner_, &context_);
+ // Define types matching the CWrapper expectation (float for GPU)
+ using coord_t = float;
+ using fpoint_t = gpuspatial::Point<coord_t, 2>;
+ using box_t = gpuspatial::Box<fpoint_t>;
+
+ // 1. Load Data using ReadParquet
+ auto poly_path =
TestUtils::GetTestDataPath("synthetic_pip/polygons.parquet");
+ auto point_path = TestUtils::GetTestDataPath("synthetic_pip/points.parquet");
+
+ // ReadParquet loads the file into batches
(std::vector<std::shared_ptr<arrow::Array>>)
+ // Assuming batch_size=1000 or similar default inside ReadParquet
+ auto poly_arrays = gpuspatial::ReadParquet(poly_path);
+ auto point_arrays = gpuspatial::ReadParquet(point_path);
+
+ // 2. Setup GEOS C++ Objects
+ auto factory = geos::geom::GeometryFactory::create();
+ geos::io::WKBReader wkb_reader(*factory);
+
+ // Iterate over batches (replacing the stream loop)
+ size_t num_batches = std::min(poly_arrays.size(), point_arrays.size());
+
+ for (size_t i = 0; i < num_batches; i++) {
+ auto build_arrow_arr = poly_arrays[i];
+ auto probe_arrow_arr = point_arrays[i];
+
+ // Export to C-style ArrowArrays for the CWrapper (SUT)
+ nanoarrow::UniqueArray build_array, probe_array;
+ nanoarrow::UniqueSchema build_schema, probe_schema;
+
+ ARROW_THROW_NOT_OK(
+ arrow::ExportArray(*build_arrow_arr, build_array.get(),
build_schema.get()));
+ ARROW_THROW_NOT_OK(
+ arrow::ExportArray(*probe_arrow_arr, probe_array.get(),
probe_schema.get()));
+
+ // --- Build Phase ---
+
+ // Boxes for the GPU Index (SUT)
+ std::vector<box_t> rects;
+
+ // View for parsing WKB
+ nanoarrow::UniqueArrayView build_view;
+ ArrowError error;
- auto poly_path = TestUtils::GetTestDataPath("arrowipc/test_polygons.arrows");
- auto point_path = TestUtils::GetTestDataPath("arrowipc/test_points.arrows");
- nanoarrow::UniqueArrayStream poly_stream, point_stream;
+ ASSERT_EQ(ArrowArrayViewInitFromSchema(build_view.get(),
build_schema.get(), &error),
+ NANOARROW_OK)
+ << error.message;
+ ASSERT_EQ(ArrowArrayViewSetArray(build_view.get(), build_array.get(),
&error),
+ NANOARROW_OK)
+ << error.message;
- gpuspatial::ArrayStreamFromIpc(poly_path, "geometry", poly_stream.get());
- gpuspatial::ArrayStreamFromIpc(point_path, "geometry", point_stream.get());
+ for (int64_t j = 0; j < build_array->length; j++) {
+ // Parse WKB
+ ArrowStringView wkb = ArrowArrayViewGetStringUnsafe(build_view.get(), j);
+ auto geom = wkb_reader.read(reinterpret_cast<const unsigned
char*>(wkb.data),
+ wkb.size_bytes);
- nanoarrow::UniqueSchema build_schema, stream_schema;
- nanoarrow::UniqueArray build_array, stream_array;
- ArrowError error;
- ArrowErrorSet(&error, "");
+ // Extract Envelope for GPU
+ const auto* env = geom->getEnvelopeInternal();
+ double xmin = 0, ymin = 0, xmax = -1, ymax = -1;
+ if (!env->isNull()) {
+ xmin = env->getMinX();
+ ymin = env->getMinY();
+ xmax = env->getMaxX();
+ ymax = env->getMaxY();
+ }
- int n_row_groups = 100;
+ // Add to GPU Build Data
+ rects.emplace_back(fpoint_t((float)xmin, (float)ymin),
+ fpoint_t((float)xmax, (float)ymax));
+ }
- for (int i = 0; i < n_row_groups; i++) {
- ASSERT_EQ(ArrowArrayStreamGetNext(poly_stream.get(), build_array.get(),
&error),
- NANOARROW_OK);
- ASSERT_EQ(ArrowArrayStreamGetSchema(poly_stream.get(), build_schema.get(),
&error),
- NANOARROW_OK);
+ // Initialize SUT (CWrapper Index)
+ index_.clear(&index_);
+ ASSERT_EQ(index_.push_build(&index_, (float*)rects.data(), rects.size()),
0);
+ ASSERT_EQ(index_.finish_building(&index_), 0);
- ASSERT_EQ(ArrowArrayStreamGetNext(point_stream.get(), stream_array.get(),
&error),
- NANOARROW_OK);
- ASSERT_EQ(ArrowArrayStreamGetSchema(point_stream.get(),
stream_schema.get(), &error),
- NANOARROW_OK);
+ // --- Probe Phase ---
+ std::vector<box_t> queries;
- joiner_.push_build(&joiner_, build_schema.get(), build_array.get(), 0,
- build_array->length);
- joiner_.finish_building(&joiner_);
+ nanoarrow::UniqueArrayView probe_view;
- joiner_.push_stream(&joiner_, &context_, stream_schema.get(),
stream_array.get(), 0,
- stream_array->length, GpuSpatialPredicateContains, 0);
+ ASSERT_EQ(ArrowArrayViewInitFromSchema(probe_view.get(),
probe_schema.get(), &error),
+ NANOARROW_OK)
+ << error.message;
+ ASSERT_EQ(ArrowArrayViewSetArray(probe_view.get(), probe_array.get(),
&error),
+ NANOARROW_OK)
+ << error.message;
- void* build_indices_ptr;
- void* stream_indices_ptr;
+ for (int64_t j = 0; j < probe_array->length; j++) {
+ ArrowStringView wkb = ArrowArrayViewGetStringUnsafe(probe_view.get(), j);
Review Comment:
```suggestion
ArrowBufferView wkb = ArrowArrayViewGetBytesUnsafe(probe_view.get(),
j);
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
