Loading...
Searching...
No Matches
octree_iterator.hpp
/*
* Software License Agreement (BSD License)
*
* Copyright (c) 2011, Willow Garage, Inc.
* All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
*
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above
* copyright notice, this list of conditions and the following
* disclaimer in the documentation and/or other materials provided
* with the distribution.
* * Neither the name of Willow Garage, Inc. nor the names of its
* contributors may be used to endorse or promote products derived
* from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
* "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
* LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
* FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
* COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
* INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
* LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
* ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
*
* Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com)
*/
#ifndef PCL_GPU_OCTREE_ITERATOR
#define PCL_GPU_OCTREE_ITERATOR
namespace pcl
{
namespace device
{
template<int CTA_SIZE, int STACK_DEPTH>
struct OctreeIteratorDevice
{
using Storage = int[STACK_DEPTH][CTA_SIZE];
int level;
Storage& storage;
__device__ __forceinline__ OctreeIteratorDevice(Storage& storage_arg) : storage(storage_arg)
{
level = 0; // root level
storage[level][threadIdx.x] = (0 << 8) + 1;
}
__device__ __forceinline__ void gotoNextLevel(int first, int len)
{
++level;
storage[level][threadIdx.x] = (first << 8) + len;
}
__device__ __forceinline__ int operator*() const
{
return storage[level][threadIdx.x] >> 8;
}
__device__ __forceinline__ void operator++()
{
while(level >= 0)
{
int data = storage[level][threadIdx.x];
if ((data & 0xFF) > 1) // there are another siblings, can goto there
{
data += (1 << 8) - 1; // +1 to first and -1 from len
storage[level][threadIdx.x] = data;
break;
}
else
--level; //goto parent;
}
}
};
struct OctreeIteratorDeviceNS
{
int level;
int node_idx;
int length;
const OctreeGlobalWithBox& octree;
__device__ __forceinline__ OctreeIteratorDeviceNS(const OctreeGlobalWithBox& octree_arg) : octree(octree_arg)
{
node_idx = 0;
level = 0;
length = 1;
}
__device__ __forceinline__ void gotoNextLevel(int first, int len)
{
node_idx = first;
length = len;
++level;
}
__device__ __forceinline__ int operator*() const
{
return node_idx;
}
__device__ __forceinline__ void operator++()
{
#if 1
while(level >= 0)
{
if (length > 1)
{
length--;
node_idx++;
break;
}
if (node_idx == 0)
{
level = -1;
return;
}
node_idx = octree.parent[node_idx];
--level;
if (node_idx == 0)
{
level = -1;
return;
}
int parent = octree.nodes[octree.parent[node_idx]];
int parent_first = parent >> 8;
int parent_len = __popc(parent & 0xFF);
int pos = node_idx - parent_first;
length = parent_len - pos;
}
#endif
}
};
}
}
#endif /* PCL_GPU_OCTREE_ITERATOR */
pcl
Definition convolution.h:46
pcl::device::OctreeGlobal::parent
int * parent
Definition internal.hpp:57
pcl::device::OctreeGlobal::nodes
int * nodes
Definition internal.hpp:50
pcl::device::OctreeGlobalWithBox
Definition internal.hpp:63
pcl::device::OctreeIteratorDevice
Definition octree_iterator.hpp:47
pcl::device::OctreeIteratorDevice::operator++
__device__ __forceinline__ void operator++()
Definition octree_iterator.hpp:70
pcl::device::OctreeIteratorDevice::storage
Storage & storage
Definition octree_iterator.hpp:51
pcl::device::OctreeIteratorDevice::operator*
__device__ __forceinline__ int operator*() const
Definition octree_iterator.hpp:65
pcl::device::OctreeIteratorDevice::level
int level
Definition octree_iterator.hpp:50
pcl::device::OctreeIteratorDevice::gotoNextLevel
__device__ __forceinline__ void gotoNextLevel(int first, int len)
Definition octree_iterator.hpp:59
pcl::device::OctreeIteratorDevice::Storage
int[STACK_DEPTH][CTA_SIZE] Storage
Definition octree_iterator.hpp:48
pcl::device::OctreeIteratorDevice::OctreeIteratorDevice
__device__ __forceinline__ OctreeIteratorDevice(Storage &storage_arg)
Definition octree_iterator.hpp:53
pcl::device::OctreeIteratorDeviceNS
Definition octree_iterator.hpp:89
pcl::device::OctreeIteratorDeviceNS::gotoNextLevel
__device__ __forceinline__ void gotoNextLevel(int first, int len)
Definition octree_iterator.hpp:102
pcl::device::OctreeIteratorDeviceNS::OctreeIteratorDeviceNS
__device__ __forceinline__ OctreeIteratorDeviceNS(const OctreeGlobalWithBox &octree_arg)
Definition octree_iterator.hpp:95
pcl::device::OctreeIteratorDeviceNS::operator++
__device__ __forceinline__ void operator++()
Definition octree_iterator.hpp:114
pcl::device::OctreeIteratorDeviceNS::octree
const OctreeGlobalWithBox & octree
Definition octree_iterator.hpp:93
pcl::device::OctreeIteratorDeviceNS::level
int level
Definition octree_iterator.hpp:90
pcl::device::OctreeIteratorDeviceNS::node_idx
int node_idx
Definition octree_iterator.hpp:91
pcl::device::OctreeIteratorDeviceNS::length
int length
Definition octree_iterator.hpp:92
pcl::device::OctreeIteratorDeviceNS::operator*
__device__ __forceinline__ int operator*() const
Definition octree_iterator.hpp:109