ArborX::AccessTraits - aprokop/ArborX GitHub Wiki

ArborX

Note

For ArborX 1.x, see ArborX::AccessTraits (v1))

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 the following partial specialization for Kokkos views

template <typename View>
AccessTraits<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

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::AccessTraits<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;
};

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{{0, 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::BVH<Kokkos::CudaSpace> bvh{Kokkos::Cuda{}, 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::AccessTraits<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 (primitives or queries).

ArborX::AccessTraits<T>::get

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

If provided to a constructor, the return type must decay to the Value template argument of the spatial index.

If provided in a query, the return value type must decay to a valid ArborX predicate, e.g. a predicate generated by ArborX::nearest() or ArborX::intersects().

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() pos : position of the primitive from which to return the centroid or the smallest bounding volume

Returns

Notes

⚠️ **GitHub.com Fallback** ⚠️