Fermat
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Modules Pages
radixtree_inline.h
1 /*
2  * Copyright (c) 2010-2016, NVIDIA Corporation
3  * All rights reserved.
4  *
5  * Redistribution and use in source and binary forms, with or without
6  * modification, are permitted provided that the following conditions are met:
7  * * Redistributions of source code must retain the above copyright
8  * notice, this list of conditions and the following disclaimer.
9  * * Redistributions in binary form must reproduce the above copyright
10  * notice, this list of conditions and the following disclaimer in the
11  * documentation and/or other materials provided with the distribution.
12  * * Neither the name of NVIDIA Corporation nor the
13  * names of its contributors may be used to endorse or promote products
14  * derived from this software without specific prior written permission.
15  *
16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19  * DISCLAIMED. IN NO EVENT SHALL <COPYRIGHT HOLDER> BE LIABLE FOR ANY
20  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
26  */
27 
28 #include <cugar/basic/functors.h>
29 #include <cugar/basic/algorithms.h>
30 #include <cugar/basic/utils.h>
31 #include <stack>
32 
33 namespace cugar {
34 namespace bintree {
35 
36  struct Split_task
37  {
38  CUGAR_HOST_DEVICE Split_task() {}
39  CUGAR_HOST_DEVICE Split_task(const uint32 id, const uint32 begin, const uint32 end, const uint32 level, const uint32 parent)
40  : m_node( id ), m_begin( begin ), m_end( end ), m_parent( parent ), m_level( level ) {}
41 
42  uint32 m_node;
43  uint32 m_begin;
44  uint32 m_end;
45  uint32 m_parent : 26;
46  uint32 m_level : 6;
47  };
48 
49  // find the most significant bit smaller than start by which code0 and code1 differ
50  template <typename Integer>
51  CUGAR_FORCEINLINE CUGAR_HOST_DEVICE int32 find_leading_bit_difference(
52  const int32 start_level,
53  const Integer code0,
54  const Integer code1)
55  {
56  int32 level = start_level;
57 
58  while (level >= 0)
59  {
60  const Integer mask = Integer(1u) << level;
61 
62  if ((code0 & mask) !=
63  (code1 & mask))
64  break;
65 
66  --level;
67  }
68  return level;
69  }
70 
71 } // namespace bintree
72 
73 template <typename Tree, typename Integer>
75  const uint32 n_codes,
76  const Integer* codes,
77  const uint32 bits,
78  const uint32 max_leaf_size,
79  const bool keep_singletons,
80  const bool middle_splits,
81  Tree& tree)
82 {
84 
85  tree.reserve_nodes( n_codes * 2 );
86  tree.reserve_leaves( n_codes );
87 
88  std::vector<Split_task> task_queues[2];
89  std::vector<uint32> skip_queues[2];
90 
91  task_queues[0].push_back( Split_task(0,0,n_codes,bits-1,uint32(-1)) );
92  skip_queues[0].push_back( uint32(-1) );
93 
94  uint32 node_count = 1;
95  uint32 leaf_count = 0;
96 
97  typename Tree::context_type context = tree.get_context();
98 
99  uint32 in_queue = 0;
100  uint32 out_queue = 1;
101 
102  while (task_queues[ in_queue ].size())
103  {
104  task_queues[ out_queue ].erase(
105  task_queues[ out_queue ].begin(),
106  task_queues[ out_queue ].end() );
107 
108  for (uint32 task_id = 0; task_id < task_queues[ in_queue ].size(); ++task_id)
109  {
110  const Split_task task = task_queues[ in_queue ][ task_id ];
111  const uint32 skip_node = skip_queues[ in_queue ][ task_id ];
112 
113  const uint32 node = task.m_node;
114  const uint32 begin = task.m_begin;
115  const uint32 end = task.m_end;
116  uint32 level = task.m_level;
117  uint32 parent = task.m_parent;
118 
119  if (!keep_singletons)
120  {
121  level = bintree::find_leading_bit_difference(
122  level,
123  codes[ begin ],
124  codes[ end-1 ] );
125  }
126 
127  uint32 output_count = 0;
128  uint32 split_index;
129 
130  // check whether the input node really needs to be split
131  if (end - begin > max_leaf_size && level != uint32(-1))
132  {
133  // find the "partitioning pivot" using a binary search
134  split_index = find_pivot(
135  codes + begin,
136  end - begin,
137  mask_and<Integer>( Integer(1u) << level ) ) - codes;
138 
139  output_count = (split_index == begin || split_index == end) ? 1u : 2u;
140  }
141 
142  const uint32 node_offset = node_count;
143  const uint32 first_end = (output_count == 1) ? end : split_index;
144  const uint32 first_skip = (output_count == 1) ? skip_node : node_offset+1;
145 
146  if (output_count >= 1) { task_queues[ out_queue ].push_back( Split_task( node_offset+0, begin, first_end, level-1, node ) ); skip_queues[ out_queue ].push_back( first_skip ); }
147  if (output_count == 2) { task_queues[ out_queue ].push_back( Split_task( node_offset+1, split_index, end, level-1, node ) ); skip_queues[ out_queue ].push_back( skip_node ); }
148 
149  const bool generate_leaf = output_count == 0;
150 
151  // count how many leaves we need to generate
152  const uint32 leaf_index = leaf_count;
153 
154  node_count += output_count;
155  leaf_count += generate_leaf ? 1 : 0;
156 
157  // write the parent node
158  context.write_node(
159  node,
160  parent,
161  output_count ? split_index != begin : false,
162  output_count ? split_index != end : false,
163  output_count ? node_offset : leaf_index,
164  skip_node,
165  level,
166  begin,
167  end,
168  output_count ? split_index : uint32(-1) );
169 
170  // make a leaf if necessary
171  if (generate_leaf)
172  context.write_leaf( leaf_index, node, begin, end );
173  }
174 
175  std::swap( in_queue, out_queue );
176  }
177 }
178 
179 } // namespace cugar
thrust::device_vector< T >::iterator begin(thrust::device_vector< T > &vec)
Definition: thrust_view.h:89
Defines some general purpose functors.
Define CUDA utilities.
Definition: radixtree_inline.h:36
Defines some general purpose algorithms.
Definition: functors.h:481
Define a vector_view POD type and plain_view() for std::vector.
Definition: diff.h:38
void generate_radix_tree(const uint32 n_codes, const Integer *codes, const uint32 bits, const uint32 max_leaf_size, const bool keep_singletons, const bool middle_splits, Tree_writer &tree)
Definition: radixtree_inline.h:381
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE Iterator find_pivot(Iterator begin, const uint32 n, const Predicate predicate)
Definition: algorithms.h:63