ArborX::RangeAdaptor

ArborX

ArborX::RangeAdaptor

Defined in header <ArborX_RangeAdaptor.hpp>

template <typename T>
struct RangeAdaptor;

The struct template ArborX::RangeAdaptor tells the ArborX::BVH constructor how much data to index and how to access that data, or ArborX::BVH::query() what predicates to query for.

ArborX provides the following partial specialization for Kokkos views

template <typename View>
RangeAdaptor<View, std::enable_if_t<Kokkos::is_view<View>{}>>;

Template parameters

T : User data type.

Member types

Member type	Description
memory_space	A valid Kokkos memory space
value_type	A type of a single indexed value

Static member functions

size	returns the number of elements
get	access the specified element

Example

#include <ArborX.hpp>

#include <Kokkos_Core.hpp>

#include <array>
#include <iostream>
#include <numeric>

struct PointCloud
{
  float *d_x;
  float *d_y;
  float *d_z;
  int N;
};

struct NearestToOrigin
{
  int k;
};

template <>
struct ArborX::RangeAdaptor<PointCloud>
{
  static KOKKOS_FUNCTION std::size_t size(PointCloud const &cloud)
  {
    return cloud.N;
  }
  static KOKKOS_FUNCTION static ArborX::Point get(PointCloud const &cloud,
                                                  std::size_t i)
  {
    return {{cloud.d_x[i], cloud.d_y[i], cloud.d_z[i]}};
  }
  using memory_space = Kokkos::CudaSpace;
  using value_type = ArborX::Point;
};

template <>
struct ArborX::RangeAdaptor<NearestToOrigin>
{
  static KOKKOS_FUNCTION std::size_t size(NearestToOrigin const &)
  {
    return 1;
  }
  static KOKKOS_FUNCTION auto get(NearestToOrigin const &d, std::size_t)
  {
    return ArborX::nearest(ArborX::Point{{0, 0, 0}}, d.k);
  }
  using memory_space = Kokkos::CudaSpace;
  using value_type = ArborX::Nearest<ArborX::Point>;
};

int main(int argc, char *argv[])
{
  Kokkos::ScopeGuard guard(argc, argv);

  constexpr std::size_t N = 1000;
  std::array<float, N> a;

  float *d_a;
  cudaMalloc(&d_a, sizeof(float) * N);

  std::iota(std::begin(a), std::end(a), 1.0);

  cudaMemcpy(d_a, a.data(), sizeof(float) * N, cudaMemcpyHostToDevice);

  using device_type = Kokkos::Cuda::device_type;
  ArborX::BVH<Kokkos::CudaSpace> bvh{Kokkos::Cuda{}, ArborX::AttachIndices{PointCloud{d_a, d_a, d_a, N}}};

  Kokkos::View<int *, device_type> indices("indices", 0);
  Kokkos::View<int *, device_type> offset("offset", 0);
  bvh.query(Kokkos::Cuda{}, NearestToOrigin{5}, indices, offset);

  Kokkos::parallel_for(1, KOKKOS_LAMBDA(int i) {
    for (int j = offset(i); j < offset(i + 1); ++j)
    {
      printf("%i %i\n", i, indices(j));
    }
  });

  return 0;
}

Output

0 0
0 1
0 2
0 3
0 4

ArborX::RandomAdaptor<T>::size

static KOKKOS_FUNCTION size_type size(T const& data);

Parameters

data : data argument passed to the ArborX::BVH<MemorySpace>::BVH constructor or ArborX::BVH<MemorySpace>::query()

Returns

The number of data values.

ArborX::RandomAdaptor<T>::get

static KOKKOS_FUNCTION context_dependent get(T const& data, size_type index);

The return value type must decay to value_type.

get() needs to be marked with the KOKKOS_FUNCTION or KOKKOS_INLINE_FUNCTION macro because it is called from the device.

Parameters

data : data argument passed to the ArborX::BVH<MemorySpace>::BVH constructor or ArborX::BVH<MemorySpace>::query() index : index of a data point

Returns

Returns a data point.

Notes