Skip to content

Commit 6441305

Browse files
RichardGeRichardGe
authored
Merge pull request #46 from GPUOpen-LibrariesAndSDKs/next-release-9
Update to 3.0.4fea77f
2 parents 8c9df52 + e4db2d7 commit 6441305

File tree

10 files changed

+88
-59
lines changed

10 files changed

+88
-59
lines changed

‎CMakeLists.txt‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -412,12 +412,15 @@ if(PRECOMPILE)
412412
${CMAKE_SOURCE_DIR}/hiprt/impl/Instance.h
413413
${CMAKE_SOURCE_DIR}/hiprt/impl/InstanceList.h
414414
${CMAKE_SOURCE_DIR}/hiprt/impl/MortonCode.h
415+
${CMAKE_SOURCE_DIR}/hiprt/impl/Header.h
415416
${CMAKE_SOURCE_DIR}/hiprt/impl/TriangleMesh.h
416417
${CMAKE_SOURCE_DIR}/hiprt/impl/Triangle.h
417418
${CMAKE_SOURCE_DIR}/hiprt/impl/BvhBuilderUtil.h
418419
${CMAKE_SOURCE_DIR}/hiprt/impl/SbvhCommon.h
419420
${CMAKE_SOURCE_DIR}/hiprt/impl/BvhConfig.h
421+
${CMAKE_SOURCE_DIR}/hiprt/impl/NodeList.h
420422
${CMAKE_SOURCE_DIR}/hiprt/impl/MemoryArena.h
423+
${CMAKE_SOURCE_DIR}/hiprt/impl/Obb.h
421424
${CMAKE_SOURCE_DIR}/hiprt/hiprt_types.h
422425
${CMAKE_SOURCE_DIR}/hiprt/hiprt_common.h
423426
)

‎hiprt/impl/BvhBuilderKernels.h‎

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -805,6 +805,7 @@ extern "C" __global__ void FitBounds_InstanceList_MatrixFrame_HwInstanceNode(
805805
FitBounds( header, primitives, boxNodes, primNodes );
806806
}
807807

808+
// assuming that OBBs are AMD specific, there are no warp syncs
808809
template <typename PrimitiveContainer, typename PrimitiveNode, typename Header>
809810
__device__ void FitOrientedBounds(
810811
Header* header,
@@ -878,6 +879,22 @@ __device__ void FitOrientedBounds(
878879
boxNodes[index].initBox(
879880
laneIndex, childCount, childIndex, childBox, nodeBox, childRange, MatrixIndexToId[matrixIndex] );
880881
}
882+
883+
// revert aabb if obb is not better
884+
if ( laneIndex == 0 )
885+
{
886+
// reconstructed quantized boxes
887+
float aabbArea = 0.0f;
888+
float obbArea = 0.0f;
889+
for ( uint32_t j = 0; j < node.getChildCount(); ++j )
890+
{
891+
aabbArea += node.getChildBox( j ).area();
892+
obbArea += boxNodes[index].getChildBox( j ).area();
893+
}
894+
895+
// compare to aabb surface area
896+
if ( aabbArea < ObbSurfaceAreaAlpha * obbArea ) boxNodes[index] = node;
897+
}
881898
}
882899

883900
index = node.getParentAddr();

‎hiprt/impl/BvhConfig.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,8 @@ static constexpr uint32_t BvhBuilderCompactionBlockSize = 1024;
3535
static constexpr uint32_t BatchBuilderMaxBlockSize = MaxBatchBuildMaxPrimCount;
3636
static constexpr uint32_t MaxFatLeafSize = 4u;
3737
static constexpr uint32_t LanesPerLeafPacketTask = 4u;
38+
static constexpr float ObbEnlargeEpsilon = 0.05f;
39+
static constexpr float ObbSurfaceAreaAlpha = 1.1f;
3840
// LBVH
3941
static constexpr uint32_t LbvhEmitBlockSize = 512u;
4042
// PLOC

‎hiprt/impl/BvhNode.h‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,6 @@
2626
#include <hiprt/hiprt_types.h>
2727
#include <hiprt/impl/Aabb.h>
2828
#include <hiprt/impl/Triangle.h>
29-
#include <hiprt/impl/Transform.h>
3029

3130
namespace hiprt
3231
{

‎hiprt/impl/InstanceList.h‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,6 @@
2525
#pragma once
2626
#include <hiprt/hiprt_types.h>
2727
#include <hiprt/impl/Header.h>
28-
#include <hiprt/impl/Transform.h>
2928

3029
namespace hiprt
3130
{
@@ -45,9 +44,9 @@ class InstanceList
4544
#if defined( __KERNELCC__ )
4645
HIPRT_DEVICE InstanceNode fetchPrimNode( const uint32_t index ) const
4746
{
48-
const Frame frame = fetchFrame( index );
4947
const hiprtInstance instance = fetchInstance( index );
5048
const hiprtTransformHeader transform = fetchTransformHeader( index );
49+
const Frame frame = fetchFrame( transform.frameIndex );
5150
const uint32_t mask = fetchMask( index );
5251

5352
InstanceNode instanceNode{};

‎hiprt/impl/Triangle.h‎

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@
2424

2525
#pragma once
2626
#include <hiprt/hiprt_types.h>
27-
#include <hiprt/impl/Aabb.h>
27+
#include <hiprt/impl/BvhConfig.h>
2828
#include <hiprt/impl/Obb.h>
2929

3030
namespace hiprt
@@ -113,17 +113,12 @@ class alignas( alignof( float3 ) ) Triangle
113113
HIPRT_HOST_DEVICE void
114114
crop( const uint32_t axis, const float position, const Aabb& box, BoundingVolume& boundingVolume ) const
115115
{
116-
constexpr float EnlargeEpsilon = 0.05f;
117-
constexpr float EdgeEnlargeRatio = 1.0f + 2.0f * EnlargeEpsilon;
118-
constexpr float AreaEnlargeDelta = EdgeEnlargeRatio * EdgeEnlargeRatio - 1.0f;
119-
120-
Aabb croppedBox;
121-
122116
// use enlarged box to make sure that the split points are inside
123117
Aabb enlargedBox = box;
124-
enlargedBox.m_min -= EnlargeEpsilon * box.extent();
125-
enlargedBox.m_max += EnlargeEpsilon * box.extent();
118+
enlargedBox.m_min -= ObbEnlargeEpsilon * box.extent();
119+
enlargedBox.m_max += ObbEnlargeEpsilon * box.extent();
126120

121+
Aabb croppedBox;
127122
const float3* vertices = &m_v0;
128123
const float3* v1 = &vertices[2];
129124
for ( uint32_t i = 0; i < 3; i++ )
@@ -151,7 +146,7 @@ class alignas( alignof( float3 ) ) Triangle
151146
}
152147

153148
// fallback when the cropped box is invalid
154-
if ( !croppedBox.valid() || box.area() - croppedBox.area() > AreaEnlargeDelta ) boundingVolume.grow( box );
149+
if ( !croppedBox.valid() ) boundingVolume.grow( box );
155150
}
156151

157152
HIPRT_HOST_DEVICE Aabb aabb() const

‎hiprt/impl/hiprt_device_impl.h‎

Lines changed: 36 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -277,17 +277,17 @@ class TraversalBase
277277
hiprtHit& hit );
278278

279279
HIPRT_DEVICE bool testTriangle(
280-
const hiprtRay& ray, const float3& invD, TriangleNode* nodes, uint32_t leafAddr, uint32_t leafType, hiprtHit& hit );
280+
const hiprtRay& ray, const float3& invD, TrianglePairNode* nodes, uint32_t leafAddr, uint32_t leafType, hiprtHit& hit );
281281

282282
HIPRT_DEVICE uint32_t testTrianglePair(
283-
const hiprtRay& ray,
284-
TriangleNode* nodes,
285-
uint32_t leafAddr,
286-
uint32_t triPairIndex,
287-
hiprtHit& hit0,
288-
hiprtHit& hit1,
289-
bool& nodeEnd,
290-
bool& rangeEnd );
283+
const hiprtRay& ray,
284+
TrianglePacketNode* nodes,
285+
uint32_t leafAddr,
286+
uint32_t triPairIndex,
287+
hiprtHit& hit0,
288+
hiprtHit& hit1,
289+
bool& nodeEnd,
290+
bool& rangeEnd );
291291

292292
protected:
293293
hiprtRay m_ray;
@@ -421,19 +421,20 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangleNode(
421421
bool hasHit = false;
422422
uint32_t leafAddr = getNodeAddr( leafIndex );
423423

424-
if constexpr ( Rtip >= 31 )
424+
if constexpr ( is_same<TriangleNode, TrianglePacketNode>::value ) // RTIP 3.1
425425
{
426+
TrianglePacketNode* packetNodes = reinterpret_cast<TrianglePacketNode*>( nodes );
427+
426428
hiprtHit secondHit;
427429
uint32_t triPairIndex = typeToTriPairIndex( getNodeType( leafIndex ) );
428-
429430
if constexpr ( TraversalType == hiprtTraversalTerminateAtAnyHit )
430431
{
431432
while ( true )
432433
{
433434
bool nodeEnd = false;
434435
bool rangeEnd = false;
435436
uint32_t hitMask =
436-
this->testTrianglePair( ray, nodes, leafAddr, triPairIndex, hit, secondHit, nodeEnd, rangeEnd );
437+
this->testTrianglePair( ray, packetNodes, leafAddr, triPairIndex, hit, secondHit, nodeEnd, rangeEnd );
437438

438439
bool firstHasHit = hitMask & 1;
439440
bool secondHasHit = hitMask & 2;
@@ -456,12 +457,6 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangleNode(
456457
triangleMask |= Triangle1Processed;
457458
}
458459

459-
if ( rangeEnd )
460-
{
461-
triangleMask = InvalidValue; // indicate range end by 'invalid value'
462-
break;
463-
}
464-
465460
triPairIndex++;
466461
triangleMask = 0;
467462
if ( nodeEnd )
@@ -470,9 +465,12 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangleNode(
470465
leafAddr++;
471466
}
472467

473-
if ( hasHit )
468+
if ( hasHit || rangeEnd )
474469
{
475-
leafIndex = encodeNodeIndex( leafAddr, triPairIndexToType( triPairIndex ) );
470+
if ( rangeEnd && !( firstHasHit && secondHasHit && ( triangleMask & Triangle1Processed ) == 0 ) )
471+
triangleMask = InvalidValue; // indicate range end by 'invalid value'
472+
else
473+
leafIndex = encodeNodeIndex( leafAddr, triPairIndexToType( triPairIndex ) );
476474
break;
477475
}
478476
}
@@ -486,7 +484,7 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangleNode(
486484
bool nodeEnd = false;
487485
bool rangeEnd = false;
488486
uint32_t hitMask =
489-
this->testTrianglePair( ray, nodes, leafAddr, triPairIndex, firstHit, secondHit, nodeEnd, rangeEnd );
487+
this->testTrianglePair( ray, packetNodes, leafAddr, triPairIndex, firstHit, secondHit, nodeEnd, rangeEnd );
490488

491489
bool firstHasHit = hitMask & 1;
492490
bool secondHasHit = hitMask & 2;
@@ -521,28 +519,30 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangleNode(
521519
}
522520
else
523521
{
522+
TrianglePairNode* pairNodes = reinterpret_cast<TrianglePairNode*>( nodes );
523+
524524
if constexpr ( TraversalType == hiprtTraversalTerminateAtAnyHit )
525525
{
526526
if ( ( triangleMask & Triangle0Processed ) == 0 )
527527
{
528-
hasHit = this->testTriangle( ray, invD, nodes, leafAddr, TriangleType, hit );
528+
hasHit = this->testTriangle( ray, invD, pairNodes, leafAddr, TriangleType, hit );
529529
triangleMask |= Triangle0Processed;
530530
}
531531

532532
if ( !hasHit )
533533
{
534-
hasHit = this->testTriangle( ray, invD, nodes, leafAddr, TriangleType + 1, hit );
534+
hasHit = this->testTriangle( ray, invD, pairNodes, leafAddr, TriangleType + 1, hit );
535535
triangleMask |= Triangle1Processed;
536536
}
537537

538538
if ( triangleMask & Triangle1Processed ) triangleMask = InvalidValue;
539539
}
540540
else
541541
{
542-
hasHit = this->testTriangle( ray, invD, nodes, leafAddr, TriangleType, hit );
542+
hasHit = this->testTriangle( ray, invD, pairNodes, leafAddr, TriangleType, hit );
543543

544544
hiprtHit secondHit;
545-
bool secondHasHit = this->testTriangle( ray, invD, nodes, leafAddr, TriangleType + 1, secondHit );
545+
bool secondHasHit = this->testTriangle( ray, invD, pairNodes, leafAddr, TriangleType + 1, secondHit );
546546

547547
if ( secondHasHit && ( !hasHit || hit.t > secondHit.t ) )
548548
{
@@ -561,12 +561,12 @@ template <typename Stack, hiprtTraversalType TraversalType>
561561
HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangle(
562562
const hiprtRay& ray,
563563
[[maybe_unused]] const float3& invD,
564-
TriangleNode* nodes,
564+
TrianglePairNode* nodes,
565565
uint32_t leafAddr,
566566
uint32_t leafType,
567567
hiprtHit& hit )
568568
{
569-
const TriangleNode& node = nodes[leafAddr];
569+
const TrianglePairNode& node = nodes[leafAddr];
570570
if ( node.getPrimIndex( 0 ) == node.getPrimIndex( 1 ) && leafType == TriangleType + 1 ) return false;
571571
bool hasHit = false;
572572
#if HIPRT_RTIP < 20
@@ -602,17 +602,17 @@ HIPRT_DEVICE bool TraversalBase<Stack, TraversalType>::testTriangle(
602602

603603
template <typename Stack, hiprtTraversalType TraversalType>
604604
HIPRT_DEVICE uint32_t TraversalBase<Stack, TraversalType>::testTrianglePair(
605-
const hiprtRay& ray,
606-
TriangleNode* nodes,
607-
uint32_t leafAddr,
608-
uint32_t triPairIndex,
609-
hiprtHit& hit0,
610-
hiprtHit& hit1,
611-
bool& nodeEnd,
612-
bool& rangeEnd )
605+
const hiprtRay& ray,
606+
TrianglePacketNode* nodes,
607+
uint32_t leafAddr,
608+
uint32_t triPairIndex,
609+
hiprtHit& hit0,
610+
hiprtHit& hit1,
611+
bool& nodeEnd,
612+
bool& rangeEnd )
613613
{
614614
#if HIPRT_RTIP >= 31
615-
const TriangleNode& node = nodes[leafAddr];
615+
const TrianglePacketNode& node = nodes[leafAddr];
616616

617617
hip_float3 dummy0, dummy1;
618618
auto result = __builtin_amdgcn_image_bvh8_intersect_ray(

‎test/common/allocator.h‎

Lines changed: 2 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -40,7 +40,6 @@ class PoolAllocator
4040
{
4141
T items[CHUNK_ITEMS];
4242
uint32_t index = 0;
43-
Chunk* ptr = nullptr;
4443
std::atomic<uint32_t> head = 0;
4544
};
4645

@@ -53,13 +52,8 @@ class PoolAllocator
5352

5453
~PoolAllocator()
5554
{
56-
Chunk* chunk = m_ptr;
57-
while ( chunk )
58-
{
59-
Chunk* p = chunk;
60-
chunk = chunk->ptr;
61-
delete p;
62-
}
55+
for ( auto chunk : m_chunks )
56+
delete chunk;
6357
}
6458

6559
PoolAllocator( const PoolAllocator& ) = delete;
@@ -88,7 +82,6 @@ class PoolAllocator
8882
{
8983
Chunk* newChunk = new Chunk();
9084
newChunk->index = chunk->index + 1;
91-
newChunk->ptr = chunk;
9285
newChunk->head = 0;
9386
m_chunks.push_back( newChunk );
9487
m_ptr = newChunk;

0 commit comments

Comments
 (0)