Point Cloud Library (PCL) 1.12.0
Loading...
Searching...
No Matches
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
41namespace 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
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 */
__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