Fermat
ray_queues.h
1 /*
2  * Fermat
3  *
4  * Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved.
5  *
6  * Redistribution and use in source and binary forms, with or without
7  * modification, are permitted provided that the following conditions are met:
8  * * Redistributions of source code must retain the above copyright
9  * notice, this list of conditions and the following disclaimer.
10  * * Redistributions in binary form must reproduce the above copyright
11  * notice, this list of conditions and the following disclaimer in the
12  * documentation and/or other materials provided with the distribution.
13  * * Neither the name of the NVIDIA CORPORATION nor the
14  * names of its contributors may be used to endorse or promote products
15  * derived from this software without specific prior written permission.
16  *
17  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
18  * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
19  * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20  * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
21  * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
22  * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
23  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
24  * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
25  * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
26  * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27  */
28 
29 #pragma once
30 
31 // ------------------------------------------------------------------------- //
32 //
33 // Declaration of utility ray queues
34 //
35 // ------------------------------------------------------------------------- //
36 
37 #include <types.h>
38 #include <ray.h>
39 #include <cugar/basic/atomics.h>
40 #include <cugar/basic/cuda/warp_atomics.h>
41 #include <cugar/linalg/vector.h>
42 
43 union PixelInfo
44 {
45  FERMAT_HOST_DEVICE PixelInfo() {}
46  FERMAT_HOST_DEVICE PixelInfo(const uint32 _packed) : packed(_packed) {}
47  FERMAT_HOST_DEVICE PixelInfo(const uint32 _pixel, const uint32 _channel) : pixel(_pixel), channel(_channel) {}
48 
49  uint32 packed;
50  struct
51  {
52  uint32 pixel : 28;
53  uint32 channel : 4;
54  };
55 };
56 
57 struct RayQueue
58 {
59  Ray* rays;
60  Hit* hits;
61  float4* weights;
62  union {
63  float* probs;
64  uint32* light_path_id;
65  };
66  uint32* pixels;
67  float4* path_weights;
68  uint32* size;
69 
70  // construct a copy of this queue with all addressed shifted by constant
71  //
72  RayQueue offset(const uint32 _count, uint32* _size) const
73  {
74  RayQueue r;
75  r.rays = rays ? rays + _count : NULL;
76  r.hits = hits ? hits + _count : NULL;
77  r.weights = weights ? weights + _count : NULL;
78  r.probs = probs ? probs + _count : NULL;
79  r.pixels = pixels ? pixels + _count : NULL;
80  r.path_weights = path_weights ? path_weights + _count : NULL;
81  r.size = _size;
82  return r;
83  }
84 
85  FERMAT_HOST_DEVICE
86  uint32 append_slot() const { return cugar::atomic_add(size, 1u); }
87 
88  FERMAT_DEVICE
89  uint32 warp_append_slot() const { return cugar::cuda::warp_increment(size); }
90 
91  FERMAT_HOST_DEVICE
92  void append(const PixelInfo pixel, const Ray& ray, const float4 weight, const float p)
93  {
94  uint32 slot = append_slot();
95 
96  rays[slot] = ray;
97  weights[slot] = weight;
98  probs[slot] = p;
99  pixels[slot] = pixel.packed;
100  }
101  FERMAT_HOST_DEVICE
102  void append(const PixelInfo pixel, const Ray& ray, const float4 weight, const float p, float4 path_w)
103  {
104  uint32 slot = append_slot();
105 
106  rays[slot] = ray;
107  weights[slot] = weight;
108  probs[slot] = p;
109  pixels[slot] = pixel.packed;
110  path_weights[slot] = path_w;
111  }
112  FERMAT_HOST_DEVICE
113  void append(const PixelInfo pixel, const Ray& ray, const float4 weight, const uint32 path_id, float4 path_w)
114  {
115  uint32 slot = append_slot();
116 
117  rays[slot] = ray;
118  weights[slot] = weight;
119  light_path_id[slot] = path_id;
120  pixels[slot] = pixel.packed;
121  path_weights[slot] = path_w;
122  }
123 
124  FERMAT_DEVICE
125  void warp_append(const PixelInfo pixel, const Ray& ray, const float4 weight, const float p, float4 path_w)
126  {
127  const uint32 slot = cugar::cuda::warp_increment(size);
128 
129  rays[slot] = ray;
130  weights[slot] = weight;
131  probs[slot] = p;
132  pixels[slot] = pixel.packed;
133  path_weights[slot] = path_w;
134  }
135  FERMAT_DEVICE
136  void warp_append(const PixelInfo pixel, const Ray& ray, const float4 weight, const uint32 path_id, float4 path_w)
137  {
138  const uint32 slot = cugar::cuda::warp_increment(size);
139 
140  rays[slot] = ray;
141  weights[slot] = weight;
142  light_path_id[slot] = path_id;
143  pixels[slot] = pixel.packed;
144  path_weights[slot] = path_w;
145  }
146  FERMAT_DEVICE
147  void warp_append(const PixelInfo pixel, const Ray& ray, const float4 weight, const float p)
148  {
149  const uint32 slot = cugar::cuda::warp_increment(size);
150 
151  rays[slot] = ray;
152  weights[slot] = weight;
153  probs[slot] = p;
154  pixels[slot] = pixel.packed;
155  }
156  FERMAT_DEVICE
157  void warp_append(const PixelInfo pixel, const Ray& ray, const float4 weight, const uint32 path_id)
158  {
159  const uint32 slot = cugar::cuda::warp_increment(size);
160 
161  rays[slot] = ray;
162  weights[slot] = weight;
163  light_path_id[slot] = path_id;
164  pixels[slot] = pixel.packed;
165  }
166 };
cugar::cuda::warp_increment
__device__ __forceinline__ unsigned int warp_increment(unsigned int *ptr)
Definition: warp_atomics.h:56
RayQueue
Definition: ray_queues.h:57
cugar::atomic_add
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE float atomic_add(float *value, const float op)
Definition: atomics.h:100
Ray
Definition: ray.h:42
Hit
Definition: ray.h:68
PixelInfo
Definition: pathtracer_core.h:527
warp_atomics.h
Define CUDA based warp adders.
Generated by   doxygen 1.8.13