Fermat
bvh.h
Go to the documentation of this file.
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 
32 #pragma once
33 
34 #include <cugar/bvh/bvh_node.h>
36 #include <cugar/linalg/vector.h>
37 #include <cugar/linalg/bbox.h>
38 #include <cugar/basic/cuda/pointers.h>
39 #include <vector>
40 #include <stack>
41 
42 namespace cugar {
43 
76 
79 
83 template <uint32 DIM>
84 struct Bvh
85 {
88 
89  typedef Bvh_node node_type;
90 
91  std::vector<node_type> m_nodes;
92  std::vector<bbox_type> m_bboxes;
93 };
94 
98 template <uint32 DIM>
100 {
101 public:
104 
107  Bvh_builder() : m_max_leaf_size( 4u ) {}
108 
112  void set_params(const uint32 max_leaf_size) { m_max_leaf_size = max_leaf_size; }
113 
121  template <typename Iterator>
122  void build(
123  const Iterator begin,
124  const Iterator end,
125  Bvh<DIM>* bvh);
126 
129  uint32 index(const uint32 i) { return m_points[i].m_index; }
130 
131 private:
132  struct Point
133  {
134  bbox_type m_bbox;
135  uint32 m_index;
136 
137  float center(const uint32 dim) const { return (m_bbox[0][dim] + m_bbox[1][dim])*0.5f; }
138  };
139 
140  struct Node
141  {
142  uint32 m_begin;
143  uint32 m_end;
144  uint32 m_node;
145  uint32 m_depth;
146  };
147  typedef std::stack<Node> Node_stack;
148 
149  void compute_bbox(
150  const uint32 begin,
151  const uint32 end,
152  bbox_type& bbox);
153 
154  struct Bvh_partitioner;
155 
156  uint32 m_max_leaf_size;
157  std::vector<Point> m_points;
158 };
159 
161 inline float compute_sah_cost(const Bvh<3>& bvh, uint32 node_index = 0);
162 
164 inline void build_skip_nodes(const Bvh_node* nodes, uint32* skip_nodes);
165 
170 {
171  struct Reference
172  {
173  CUGAR_HOST_DEVICE
174  Reference(cugar::Bvh_node_3d* _node) : node(_node) {}
175 
176  CUGAR_HOST_DEVICE
177  operator cugar::Bbox3f() const
178  {
179  #if 0
180  return ((volatile cugar::Bvh_node_3d*)node)->bbox;
181  #elif 0
182  // fetch the node's bbox as a float2 and a float4
183  const float2 f1 = *(reinterpret_cast<const volatile float2*>( node ) + 1);
184  const float4 f2 = *(reinterpret_cast<const volatile float4*>( node ) + 1);
185  return cugar::Bbox3f( cugar::Vector3f(f1.x,f1.y,f2.x), cugar::Vector3f(f2.y,f2.z,f2.w) );
186  #else
187  // fetch the node's bbox as a float2 and a float4
188  const float2 f1 = cuda::load<cuda::LOAD_CG>(reinterpret_cast<const float2*>( node ) + 1);
189  const float4 f2 = cuda::load<cuda::LOAD_CG>(reinterpret_cast<const float4*>( node ) + 1);
190  return cugar::Bbox3f( cugar::Vector3f(f1.x,f1.y,f2.x), cugar::Vector3f(f2.y,f2.z,f2.w) );
191  #endif
192  }
193 
194  CUGAR_HOST_DEVICE
195  void operator=(const cugar::Bbox3f& bbox)
196  {
197  #if 0
198  // overwrite the node's bbox only
199  ((volatile cugar::Bvh_node_3d*)node)->bbox[0] = bbox[0];
200  ((volatile cugar::Bvh_node_3d*)node)->bbox[1] = bbox[1];
201  #elif 0
202  // store the node's bbox as a float2 and a float4
203  const float2 f1 = make_float2( bbox[0].x, bbox[0].y );
204  const float4 f2 = make_float4( bbox[0].z, bbox[1].x, bbox[1].y, bbox[1].z );
205  *(reinterpret_cast<volatile float2*>( node ) + 1) = f1;
206  *(reinterpret_cast<volatile float4*>( node ) + 1) = f2;
207  #else
208  // store the node's bbox as a float2 and a float4
209  const float2 f1 = make_float2( bbox[0].x, bbox[0].y );
210  const float4 f2 = make_float4( bbox[0].z, bbox[1].x, bbox[1].y, bbox[1].z );
211  cuda::store<cuda::STORE_CG>( reinterpret_cast<float2*>( node ) + 1, f1 );
212  cuda::store<cuda::STORE_CG>( reinterpret_cast<float4*>( node ) + 1, f2 );
213  #endif
214  }
215 
216  cugar::Bvh_node_3d* node;
217  };
218 
219  CUGAR_HOST_DEVICE
220  Bvh_node_3d_bbox_iterator(cugar::Bvh_node_3d* _nodes) : nodes(_nodes) {}
221 
222  CUGAR_HOST_DEVICE
223  Reference operator[] (const uint32 i) const { return Reference(nodes + i); }
224 
225  CUGAR_HOST_DEVICE
226  Reference operator[] (const uint32 i) { return Reference(nodes + i); }
227 
228  cugar::Bvh_node_3d* nodes;
229 };
230 
232 
233 } // namespace cugar
234 
235 #include <cugar/bvh/bvh_inline.h>
void set_params(const uint32 max_leaf_size)
Definition: bvh.h:112
thrust::device_vector< T >::iterator begin(thrust::device_vector< T > &vec)
Definition: thrust_view.h:89
Definition: bvh_node.h:45
void build_skip_nodes(const Bvh_node *nodes, uint32 *skip_nodes)
build skip nodes for a tree
Definition: bvh_inline.h:53
Entry point to the generic Bounding Volume Hierarchy library.
Definition: bvh_node.h:80
uint32 index(const uint32 i)
Definition: bvh.h:129
Bvh_builder()
Definition: bvh.h:107
Definition: bvh.h:84
Define CUDA based scan primitives.
Defines an axis-aligned bounding box class.
Definition: vector.h:54
Define a vector_view POD type and plain_view() for std::vector.
Definition: diff.h:38
Definition: bvh_inline.h:31
Definition: bvh.h:99
float compute_sah_cost(const Bvh< 3 > &bvh, uint32 node_index=0)
compute SAH cost of a subtree
Definition: bvh_inline.h:184