Point Cloud Library (PCL)  1.11.1
morton.hpp
1 /*
2 * Software License Agreement (BSD License)
3 *
4 * Copyright (c) 2011, Willow Garage, Inc.
5 * All rights reserved.
6 *
7 * Redistribution and use in source and binary forms, with or without
8 * modification, are permitted provided that the following conditions
9 * are met:
10 *
11 * * Redistributions of source code must retain the above copyright
12 * notice, this list of conditions and the following disclaimer.
13 * * Redistributions in binary form must reproduce the above
14 * copyright notice, this list of conditions and the following
15 * disclaimer in the documentation and/or other materials provided
16 * with the distribution.
17 * * Neither the name of Willow Garage, Inc. nor the names of its
18 * contributors may be used to endorse or promote products derived
19 * from this software without specific prior written permission.
20 *
21 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
22 * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
23 * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
24 * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
25 * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
26 * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
27 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
28 * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
29 * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
30 * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
31 * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
32 * POSSIBILITY OF SUCH DAMAGE.
33 *
34 * Author: Anatoly Baskeheev, Itseez Ltd, (myname.mysurname@mycompany.com)
35 */
36 
37 #ifndef PCL_GPU_OCTREE_MORTON_HPP
38 #define PCL_GPU_OCTREE_MORTON_HPP
39 
40 
41 namespace pcl
42 {
43  namespace device
44  {
45  struct Morton
46  {
47  const static int levels = 10;
48  const static int bits_per_level = 3;
49  const static int nbits = levels * bits_per_level;
50 
51  using code_t = int;
52 
53  __device__ __host__ __forceinline__
54  static int spreadBits(int x, int offset)
55  {
56  //......................9876543210
57  x = (x | (x << 10)) & 0x000f801f; //............98765..........43210
58  x = (x | (x << 4)) & 0x00e181c3; //........987....56......432....10
59  x = (x | (x << 2)) & 0x03248649; //......98..7..5..6....43..2..1..0
60  x = (x | (x << 2)) & 0x09249249; //....9..8..7..5..6..4..3..2..1..0
61 
62  return x << offset;
63  }
64 
65  __device__ __host__ __forceinline__
66  static int compactBits(int x, int offset)
67  {
68  x = ( x >> offset ) & 0x09249249; //....9..8..7..5..6..4..3..2..1..0
69  x = (x | (x >> 2)) & 0x03248649; //......98..7..5..6....43..2..1..0
70  x = (x | (x >> 2)) & 0x00e181c3; //........987....56......432....10
71  x = (x | (x >> 4)) & 0x000f801f; //............98765..........43210
72  x = (x | (x >> 10)) & 0x000003FF; //......................9876543210
73 
74  return x;
75  }
76 
77  __device__ __host__ __forceinline__
78  static code_t createCode(int cell_x, int cell_y, int cell_z)
79  {
80  return spreadBits(cell_x, 0) | spreadBits(cell_y, 1) | spreadBits(cell_z, 2);
81  }
82 
83  __device__ __host__ __forceinline__
84  static void decomposeCode(code_t code, int& cell_x, int& cell_y, int& cell_z)
85  {
86  cell_x = compactBits(code, 0);
87  cell_y = compactBits(code, 1);
88  cell_z = compactBits(code, 2);
89  }
90 
91  __host__ __device__ __forceinline__
92  static code_t extractLevelCode(code_t code, int level)
93  {
94  return (code >> (nbits - 3 * (level + 1) )) & 7;
95  }
96 
97  __host__ __device__ __forceinline__
98  static code_t shiftLevelCode(code_t level_code, int level)
99  {
100  return level_code << (nbits - 3 * (level + 1));
101  }
102  };
103 
104  struct CalcMorton
105  {
106  const static int depth_mult = 1 << Morton::levels;
107 
108  float3 minp_;
109  float3 dims_;
110 
111  __device__ __host__ __forceinline__ CalcMorton(float3 minp, float3 maxp) : minp_(minp)
112  {
113  dims_.x = maxp.x - minp.x;
114  dims_.y = maxp.y - minp.y;
115  dims_.z = maxp.z - minp.z;
116  }
117 
118  __device__ __host__ __forceinline__ Morton::code_t operator()(const float3& p) const
119  {
120  int cellx = min((int)std::floor(depth_mult * (p.x - minp_.x)/dims_.x), depth_mult - 1);
121  int celly = min((int)std::floor(depth_mult * (p.y - minp_.y)/dims_.y), depth_mult - 1);
122  int cellz = min((int)std::floor(depth_mult * (p.z - minp_.z)/dims_.z), depth_mult - 1);
123 
124  return Morton::createCode(cellx, celly, cellz);
125  }
126  __device__ __host__ __forceinline__ Morton::code_t operator()(const float4& p) const
127  {
128  return (*this)(make_float3(p.x, p.y, p.z));
129  }
130  };
131 
133  {
134  int level;
135 
136  __device__ __host__ __forceinline__
137  CompareByLevelCode(int level_arg) : level(level_arg) {}
138 
139  __device__ __host__ __forceinline__
140  bool operator()(Morton::code_t code1, Morton::code_t code2) const
141  {
143  }
144  };
145  }
146 }
147 
148 #endif /* PCL_GPU_OCTREE_MORTON_HPP */
pcl
Definition: convolution.h:46
pcl::device::Morton::decomposeCode
__device__ __host__ static __forceinline__ void decomposeCode(code_t code, int &cell_x, int &cell_y, int &cell_z)
Definition: morton.hpp:84
pcl::device::CompareByLevelCode::operator()
__device__ __host__ __forceinline__ bool operator()(Morton::code_t code1, Morton::code_t code2) const
Definition: morton.hpp:140
pcl::device::Morton::levels
static const int levels
Definition: morton.hpp:47
pcl::device::CalcMorton::CalcMorton
__device__ __host__ __forceinline__ CalcMorton(float3 minp, float3 maxp)
Definition: morton.hpp:111
pcl::device::Morton::nbits
static const int nbits
Definition: morton.hpp:49
pcl::device::Morton::createCode
__device__ __host__ static __forceinline__ code_t createCode(int cell_x, int cell_y, int cell_z)
Definition: morton.hpp:78
pcl::device::Morton::compactBits
__device__ __host__ static __forceinline__ int compactBits(int x, int offset)
Definition: morton.hpp:66
pcl::device::CompareByLevelCode::CompareByLevelCode
__device__ __host__ __forceinline__ CompareByLevelCode(int level_arg)
Definition: morton.hpp:137
pcl::device::Morton::bits_per_level
static const int bits_per_level
Definition: morton.hpp:48
pcl::device::CalcMorton::minp_
float3 minp_
Definition: morton.hpp:108
pcl::device::Morton::code_t
int code_t
Definition: morton.hpp:51
pcl::device::Morton::shiftLevelCode
__host__ __device__ static __forceinline__ code_t shiftLevelCode(code_t level_code, int level)
Definition: morton.hpp:98
pcl::device::Morton::spreadBits
__device__ __host__ static __forceinline__ int spreadBits(int x, int offset)
Definition: morton.hpp:54
pcl::device::CompareByLevelCode::level
int level
Definition: morton.hpp:134
pcl::device::CompareByLevelCode
Definition: morton.hpp:133
pcl::device::Morton
Definition: morton.hpp:46
pcl::device::CalcMorton::operator()
__device__ __host__ __forceinline__ Morton::code_t operator()(const float3 &p) const
Definition: morton.hpp:118
pcl::device::Morton::extractLevelCode
__host__ __device__ static __forceinline__ code_t extractLevelCode(code_t code, int level)
Definition: morton.hpp:92
pcl::device::CalcMorton::operator()
__device__ __host__ __forceinline__ Morton::code_t operator()(const float4 &p) const
Definition: morton.hpp:126
code
Definition: inftrees.h:24
pcl::device::CalcMorton
Definition: morton.hpp:105
pcl::device::CalcMorton::depth_mult
static const int depth_mult
Definition: morton.hpp:106
pcl::device::CalcMorton::dims_
float3 dims_
Definition: morton.hpp:109