Point Cloud Library (PCL)  1.14.1-dev
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  __device__ __host__ __forceinline__
92  static uint3 decomposeCode(code_t code)
93  {
94  return make_uint3 (compactBits(code, 0), compactBits(code, 1), compactBits(code, 2));
95  }
96 
97  __host__ __device__ __forceinline__
98  static code_t extractLevelCode(code_t code, int level)
99  {
100  return (code >> (nbits - 3 * (level + 1) )) & 7;
101  }
102 
103  __host__ __device__ __forceinline__
104  static code_t shiftLevelCode(code_t level_code, int level)
105  {
106  return level_code << (nbits - 3 * (level + 1));
107  }
108  };
109 
110  struct CalcMorton
111  {
112  const static int depth_mult = 1 << Morton::levels;
113 
114  float3 minp_;
115  float3 dims_;
116 
117  __device__ __host__ __forceinline__ CalcMorton(float3 minp, float3 maxp) : minp_(minp)
118  {
119  dims_.x = maxp.x - minp.x;
120  dims_.y = maxp.y - minp.y;
121  dims_.z = maxp.z - minp.z;
122  }
123 
124  __device__ __host__ __forceinline__ Morton::code_t operator()(const float3& p) const
125  {
126  //Using floorf due to changes to MSVC 16.9. See details here: https://devtalk.blender.org/t/cuda-compile-error-windows-10/17886/4
127  //floorf is without std:: see why here: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=79700
128  const int cellx = min((int)floorf(depth_mult * min(1.f, max(0.f, (p.x - minp_.x)/dims_.x))), depth_mult - 1);
129  const int celly = min((int)floorf(depth_mult * min(1.f, max(0.f, (p.y - minp_.y)/dims_.y))), depth_mult - 1);
130  const int cellz = min((int)floorf(depth_mult * min(1.f, max(0.f, (p.z - minp_.z)/dims_.z))), depth_mult - 1);
131 
132  return Morton::createCode(cellx, celly, cellz);
133  }
134  __device__ __host__ __forceinline__ Morton::code_t operator()(const float4& p) const
135  {
136  return (*this)(make_float3(p.x, p.y, p.z));
137  }
138  };
139 
141  {
142  int level;
143 
144  __device__ __host__ __forceinline__
145  CompareByLevelCode(int level_arg) : level(level_arg) {}
146 
147  __device__ __host__ __forceinline__
148  bool operator()(Morton::code_t code1, Morton::code_t code2) const
149  {
151  }
152  };
153  }
154 }
155 
156 #endif /* PCL_GPU_OCTREE_MORTON_HPP */
Definition: inftrees.h:24
__device__ __host__ __forceinline__ Morton::code_t operator()(const float4 &p) const
Definition: morton.hpp:134
static const int depth_mult
Definition: morton.hpp:112
__device__ __host__ __forceinline__ CalcMorton(float3 minp, float3 maxp)
Definition: morton.hpp:117
__device__ __host__ __forceinline__ Morton::code_t operator()(const float3 &p) const
Definition: morton.hpp:124
__device__ __host__ __forceinline__ CompareByLevelCode(int level_arg)
Definition: morton.hpp:145
__device__ __host__ __forceinline__ bool operator()(Morton::code_t code1, Morton::code_t code2) const
Definition: morton.hpp:148
__device__ __host__ static __forceinline__ void decomposeCode(code_t code, int &cell_x, int &cell_y, int &cell_z)
Definition: morton.hpp:84
static const int nbits
Definition: morton.hpp:49
__host__ __device__ static __forceinline__ code_t extractLevelCode(code_t code, int level)
Definition: morton.hpp:98
__device__ __host__ static __forceinline__ int spreadBits(int x, int offset)
Definition: morton.hpp:54
__device__ __host__ static __forceinline__ code_t createCode(int cell_x, int cell_y, int cell_z)
Definition: morton.hpp:78
static const int levels
Definition: morton.hpp:47
__device__ __host__ static __forceinline__ int compactBits(int x, int offset)
Definition: morton.hpp:66
static const int bits_per_level
Definition: morton.hpp:48
__device__ __host__ static __forceinline__ uint3 decomposeCode(code_t code)
Definition: morton.hpp:92
__host__ __device__ static __forceinline__ code_t shiftLevelCode(code_t level_code, int level)
Definition: morton.hpp:104