31 #include <cugar/basic/memory_arena.h> 56 for (uint32 i = 0; i < 16; ++i)
61 uint32 size(
const uint32 i)
const 66 case UINT2:
return 8u;
67 case UINT4:
return 16u;
68 case FLOAT:
return 4u;
69 case FLOAT2:
return 8u;
70 case FLOAT4:
return 16u;
75 template <
typename TUserData>
79 serialize( *
this, 0u, user );
85 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i, uint v) { queue.desc[m] = QueueDescriptor::UINT; }
86 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i, uint2 v) { queue.desc[m] = QueueDescriptor::UINT2; }
87 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i, uint4 v) { queue.desc[m] = QueueDescriptor::UINT4; }
88 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i,
float v) { queue.desc[m] = QueueDescriptor::FLOAT; }
89 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i, float2 v) { queue.desc[m] = QueueDescriptor::FLOAT2; }
90 template <u
int32 m>
void serialize_member(
QueueDescriptor& queue, uint32 i, float4 v) { queue.desc[m] = QueueDescriptor::FLOAT4; }
106 capacity = _capacity;
109 for (uint32 i = 0; i < 16; ++i)
111 const uint32 el_size = _desc.size(i);
120 offset += el_size * capacity;
127 uint32 byte_size()
const {
return offsets[ 15 ]; }
132 void alloc(uint8* _ptr, uint32* _size) { ptr = _ptr; size = _size; }
136 template <u
int32 m,
typename T>
138 T*
member_base()
const {
return reinterpret_cast<T*
>(ptr + offsets[m]); }
142 template <u
int32 m,
typename T>
144 const T&
member(
const uint32 slot)
const {
return member_base<m>()[slot]; }
148 template <u
int32 m,
typename T>
150 T&
member(
const uint32 slot) {
return member_base<m>()[slot]; }
163 void set_size(
const uint32 _size) { *size = _size; }
171 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot, uint v) { queue.
member<m,uint>(slot) = v; }
172 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot, uint2 v) { queue.
member<m,uint2>(slot) = v; }
173 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot, uint4 v) { queue.
member<m,uint4>(slot) = v; }
174 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot,
float v) { queue.
member<m,
float>(slot) = v; }
175 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot, float2 v) { queue.
member<m,float2>(slot) = v; }
176 template <u
int32 m>
void serialize_member(
WavefrontQueue& queue,
const uint32 slot, float4 v) { queue.
member<m,float4>(slot) = v; }
178 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot, uint& v) { v = queue.
member<m,uint>(slot); }
179 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot, uint2& v) { v = queue.
member<m,uint2>(slot); }
180 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot, uint4& v) { v = queue.
member<m,uint4>(slot); }
181 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot,
float& v) { v = queue.
member<m,
float>(slot); }
182 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot, float2& v) { v = queue.
member<m,float2>(slot); }
183 template <u
int32 m>
void deserialize_member(
const WavefrontQueue& queue,
const uint32 slot, float4& v) { v = queue.
member<m,float4>(slot); }
192 void alloc(uint8* _ptr, uint32* _size) { ptr = _ptr +
sizeof(float4)*3*capacity; size = _size; ray_ptr = _ptr; }
195 Ray* rays()
const {
return (
Ray*)ray_ptr; }
198 Hit* hits()
const {
return (
Hit*)(ray_ptr +
sizeof(float4)*2*capacity); }
206 reinterpret_cast<float4*
>(queue.ray_ptr)[2*i + 0] = make_float4( ray.origin.x, ray.origin.y, ray.origin.x, ray.tmin );
207 reinterpret_cast<float4*
>(queue.ray_ptr)[2*i + 1] = make_float4( ray.dir.x, ray.dir.y, ray.dir.x, ray.tmax );
213 const float4 val =
reinterpret_cast<const float4*
>(queue.ray_ptr)[2*i + 0];
214 ray.origin.x = val.x;
215 ray.origin.y = val.y;
216 ray.origin.z = val.z;
220 const float4 val =
reinterpret_cast<const float4*
>(queue.ray_ptr)[2*i + 1];
227 const float4 val =
reinterpret_cast<const float4*
>(queue.ray_ptr +
sizeof(float4)*2*queue.capacity)[i];
236 struct ScatteringPayload
239 CacheInfo cache_info;
240 CacheInfo prev_cache_info;
245 template <
typename TQueue,
typename TQueueEntry>
247 void serialize(TQueue& queue,
const TQueueEntry& i,
const ScatteringPayload& payload)
249 serialize_member<0u>( queue, i, make_uint4(
250 payload.pixel_info.packed,
251 payload.cache_info.packed,
252 payload.prev_cache_info.packed,
253 cugar::binary_cast<uint32>(payload.roughness)) );
255 serialize_member<1u>( queue, i, payload.cone );
256 serialize_member<2u>( queue, i, payload.weight );
258 template <
typename TQueue,
typename TQueueEntry>
260 void deserialize(
const TQueue& queue,
const TQueueEntry& i, ScatteringPayload& payload)
263 deserialize_member<0u>( queue, i, u );
265 payload.cache_info = CacheInfo(u.y);
266 payload.prev_cache_info = CacheInfo(u.z);
269 deserialize_member<1u>( queue, i, payload.cone );
270 deserialize_member<2u>( queue, i, payload.weight );
FERMAT_HOST_DEVICE void setup(const QueueDescriptor &_desc, const uint32 _capacity)
Definition: wavefront_queues.h:104
__device__ __forceinline__ unsigned int warp_increment(unsigned int *ptr)
Definition: warp_atomics.h:56
Definition: wavefront_queues.h:92
CUGAR_HOST_DEVICE L round_i(const L x, const R y)
Definition: numbers.h:198
FERMAT_HOST_DEVICE const T & member(const uint32 slot) const
Definition: wavefront_queues.h:144
FERMAT_HOST_DEVICE T * member_base() const
Definition: wavefront_queues.h:138
FERMAT_HOST_DEVICE void alloc(uint8 *_ptr, uint32 *_size)
Definition: wavefront_queues.h:192
FERMAT_HOST_DEVICE void set_size(const uint32 _size)
Definition: wavefront_queues.h:163
Definition: wavefront_queues.h:187
FERMAT_HOST_DEVICE void alloc(uint8 *_ptr, uint32 *_size)
Definition: wavefront_queues.h:132
Definition: wavefront_queues.h:40
CUGAR_FORCEINLINE CUGAR_HOST_DEVICE Out binary_cast(const In in)
Definition: types.h:288
Definition: pathtracer_core.h:527
Define CUDA based warp adders.
FERMAT_HOST_DEVICE WavefrontQueue()
Definition: wavefront_queues.h:99
FERMAT_HOST_DEVICE T & member(const uint32 slot)
Definition: wavefront_queues.h:150