`ArborX::AccessTraits`¶

Defined in header <ArborX_AccessTraits.hpp>

template <typename T>
struct AccessTraits;

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

ArborX provides constrained specializations for rank-1 and rank-2 Kokkos views. Rank-1 views expose their element at each position. Rank-2 views expose each row as a point. A statically known second extent determines the point dimension; for a dynamic second extent, the specialization constructs a three-dimensional point from the first three columns.

template <typename View>
  requires(Kokkos::is_view_v<View> && (View::rank() == 1 || View::rank() == 2))
struct AccessTraits<View>;

Template parameters¶

T: User data container.

Member types¶

Member type	Description
`memory_space`	A valid Kokkos memory space that the data resides in

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;
  int N;
};

struct NearestToOrigin
{
  int k;
};

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

template <>
struct ArborX::AccessTraits<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<2>{{0, 0}}, d.k);
  }
  using memory_space = Kokkos::CudaSpace;
};

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::BoundingVolumeHierarchy bvh{Kokkos::Cuda{}, PointCloud{d_a, d_a, N}};

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

  Kokkos::parallel_for(1, KOKKOS_LAMBDA(int i) {
    for (int j = offset(i); j < offset(i + 1); ++j)
      printf("[%i] (%.2f, %.2f)\n", i, values(j)[0], values(j)[1]);
  });

  return 0;
}

Output

[0] (1.00, 1.00)
[0] (2.00, 2.00)
[0] (3.00, 3.00)
[0] (4.00, 4.00)
[0] (5.00, 5.00)