Skip to content

Commit e694335

Browse files
RichardGemeistdan
andcommitted
Update to 3.0.2.4242e39
Co-Authored-By: Daniel Meister <[email protected]>
1 parent 60e6460 commit e694335

File tree

12 files changed

+224
-153
lines changed

12 files changed

+224
-153
lines changed

.github/workflows/format.yaml

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4,24 +4,25 @@ on:
44
pull_request:
55
branches:
66
- master
7-
- HRT-0-format*
87

98

109
jobs:
1110
build:
1211
runs-on: ubuntu-latest
1312

1413
steps:
15-
- uses: actions/checkout@v3
14+
- uses: actions/checkout@v4
1615
with:
1716
ref: ${{ github.event.pull_request.head.ref }}
18-
- uses: DoozyX/[email protected]
17+
18+
- uses: DoozyX/[email protected]
1919
with:
2020
source: './hiprt ./test'
2121
exclude: './contrib ./tools ./scripts'
2222
extensions: 'h,cpp,c'
2323
clangFormatVersion: 15
24-
inplace: True
24+
inplace: 'True'
25+
2526
- uses: EndBug/[email protected]
2627
with:
2728
author_name: rprbuild

.gitignore

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88
*.hipfb
99
*.bc
1010
build/
11+
build-cmake/
1112
dist/
1213
hiprt/hiprt.h
1314
hiprt/hiprtew.h
@@ -16,3 +17,4 @@ hiprt/cache/KernelArgs.h
1617
PUBLIC_OUT/
1718
hiprt/impl/bvh_build_array.h
1819
scripts/bitcodes/__pycache__/
20+
test/common/meshes/lfs

CMakeLists.txt

Lines changed: 14 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -84,10 +84,12 @@ function(get_version file)
8484
list(GET lines 0 major)
8585
list(GET lines 1 minor)
8686
list(GET lines 2 patch)
87-
set(patch_hex "0x${patch}")
87+
list(GET lines 3 hash)
88+
set(patch_hex "0x${hash}")
8889
set(major ${major} PARENT_SCOPE)
8990
set(minor ${minor} PARENT_SCOPE)
90-
set(patch ${patch_hex} PARENT_SCOPE)
91+
set(patch ${patch} PARENT_SCOPE)
92+
set(hash ${patch_hex} PARENT_SCOPE)
9193
endfunction()
9294

9395

@@ -166,8 +168,13 @@ function(get_hip_sdk_version result result_path)
166168
if(useHipFromPATH)
167169
set(hipCommand "hipcc")
168170
else()
169-
# try classic path used by HIPRT developers
170-
set(hipCommand "hipSdk\\bin\\hipcc")
171+
# use special path used by HIPRT developers
172+
if(EXISTS "hipSdk\\bin\\hipcc.exe")
173+
set(exec_perl "")
174+
set(hipCommand "hipSdk\\bin\\hipcc.exe")
175+
else()
176+
set(hipCommand "hipSdk\\bin\\hipcc")
177+
endif()
171178
endif()
172179
else()
173180

@@ -246,7 +253,7 @@ function(write_version_info in_file header_file version_file version_str_out)
246253
get_version(${version_file})
247254

248255
# set(version "${major}${minor}")
249-
# set(version_str "${version}_${patch}")
256+
# set(version_str "${version}_${hash}")
250257

251258
# Read the content of the header template file
252259
file(READ ${in_file} header_content)
@@ -265,14 +272,15 @@ function(write_version_info in_file header_file version_file version_str_out)
265272
set(version_str "${HIPRT_VERSION}" )
266273
endif()
267274

268-
# message(STATUS "HIPRT_API_VERSION: ${version_str}_${patch}")
275+
# message(STATUS "HIPRT_API_VERSION: ${version_str}_${hash}")
269276

270277
set(HIPRT_API_VERSION ${HIPRT_VERSION})
271278

272279
# Replace placeholders with actual version values
273280
string(REPLACE "@HIPRT_MAJOR_VERSION@" "${major}" header_content "${header_content}")
274281
string(REPLACE "@HIPRT_MINOR_VERSION@" "${minor}" header_content "${header_content}")
275282
string(REPLACE "@HIPRT_PATCH_VERSION@" "${patch}" header_content "${header_content}")
283+
string(REPLACE "@HIPRT_HASH_VERSION@" "${hash}" header_content "${header_content}")
276284
string(REPLACE "@HIPRT_VERSION_STR@" "\"${version_str}\"" header_content "${header_content}")
277285
string(REPLACE "@HIPRT_API_VERSION@" "${HIPRT_API_VERSION}" header_content "${header_content}")
278286

hiprt/hiprt.h.in

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@
2727
#define HIPRT_MAJOR_VERSION @HIPRT_MAJOR_VERSION@
2828
#define HIPRT_MINOR_VERSION @HIPRT_MINOR_VERSION@
2929
#define HIPRT_PATCH_VERSION @HIPRT_PATCH_VERSION@
30+
#define HIPRT_HASH_VERSION @HIPRT_HASH_VERSION@
3031

3132
#define HIPRT_API_VERSION @HIPRT_API_VERSION@
3233
#define HIPRT_VERSION_STR @HIPRT_VERSION_STR@

hiprt/hiprtew.h.in

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@
2727
#define HIPRT_MAJOR_VERSION @HIPRT_MAJOR_VERSION@
2828
#define HIPRT_MINOR_VERSION @HIPRT_MINOR_VERSION@
2929
#define HIPRT_PATCH_VERSION @HIPRT_PATCH_VERSION@
30+
#define HIPRT_HASH_VERSION @HIPRT_HASH_VERSION@
3031

3132
#define HIPRT_API_VERSION @HIPRT_API_VERSION@
3233
#define HIPRT_VERSION_STR @HIPRT_VERSION_STR@

hiprt/impl/BvhBuilderKernels.h

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -709,7 +709,7 @@ __device__ void FitBounds( Header* header, PrimitiveContainer& primitives, BoxNo
709709
{
710710
node = boxNodes[index];
711711
childCount = node.getChildCount();
712-
internal = laneIndex < childCount && node.getChildType( laneIndex ) == BoxType;
712+
internal = sublaneIndex < childCount && node.getChildType( sublaneIndex ) == BoxType;
713713
}
714714

715715
uint32_t internalCount = __popcll( hiprt::ballot( internal ) & subwarpMask );
@@ -725,7 +725,7 @@ __device__ void FitBounds( Header* header, PrimitiveContainer& primitives, BoxNo
725725
if ( !done && sublaneIndex < childCount )
726726
{
727727
childIndex = node.getChildIndex( sublaneIndex );
728-
childRange = node.getChildRange( laneIndex );
728+
childRange = node.getChildRange( sublaneIndex );
729729
childBox = getNodeBox( childIndex, primitives, boxNodes, primNodes );
730730
}
731731

@@ -736,21 +736,18 @@ __device__ void FitBounds( Header* header, PrimitiveContainer& primitives, BoxNo
736736

737737
if ( !done )
738738
{
739-
boxNodes[index].initBox( sublaneIndex, childCount, childIndex, childBox, nodeBox, childRange );
740-
741739
if ( sublaneIndex < childCount )
742-
{
743-
index = node.getParentAddr();
744-
if ( index == InvalidValue ) done = true;
745-
}
740+
boxNodes[index].initBox( sublaneIndex, childCount, childIndex, childBox, nodeBox, childRange );
741+
index = node.getParentAddr();
742+
if ( index == InvalidValue ) done = true;
746743
}
747744

748745
internal = false;
749746
if ( !done )
750747
{
751748
node = boxNodes[index];
752749
childCount = node.getChildCount();
753-
internal = laneIndex < childCount && node.getChildType( laneIndex ) == BoxType;
750+
internal = sublaneIndex < childCount && node.getChildType( sublaneIndex ) == BoxType;
754751
}
755752

756753
internalCount = __popcll( hiprt::ballot( internal ) & subwarpMask );
@@ -1967,6 +1964,9 @@ __device__ void PackLeavesWarp(
19671964
TrianglePairData( pairIndices, triIndices0, triIndices1, rangeEnd );
19681965

19691966
rangeOffset++;
1967+
// not sure why but this fixes the issue on linux
1968+
// otherwise triIndices1 are not correcly written to the final packet
1969+
__threadfence_block();
19701970
}
19711971
sync_warp();
19721972
__threadfence_block();

hiprt/impl/BvhConfig.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -35,7 +35,7 @@ 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;
38+
static constexpr float ObbEpsilon = 1.0e-4f;
3939
static constexpr float ObbSurfaceAreaAlpha = 1.1f;
4040
// LBVH
4141
static constexpr uint32_t LbvhEmitBlockSize = 512u;

hiprt/impl/BvhNode.h

Lines changed: 2 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -998,22 +998,12 @@ struct alignas( DefaultAlignment ) TrianglePacketNode
998998
while ( true )
999999
{
10001000
const Triangle& tri0 = triPacketNode->fetchTriangle( triPairIndex, 0 );
1001-
tri0.crop( 0, box.m_min.x, box, obb );
1002-
tri0.crop( 0, box.m_max.x, box, obb );
1003-
tri0.crop( 1, box.m_min.y, box, obb );
1004-
tri0.crop( 1, box.m_max.y, box, obb );
1005-
tri0.crop( 2, box.m_min.z, box, obb );
1006-
tri0.crop( 2, box.m_max.z, box, obb );
1001+
obb.m_box.grow( tri0.obb( box, matrixIndex ).m_box );
10071002

10081003
if ( triPacketNode->getPrimIndex( triPairIndex, 0 ) != triPacketNode->getPrimIndex( triPairIndex, 1 ) )
10091004
{
10101005
const Triangle& tri1 = triPacketNode->fetchTriangle( triPairIndex, 1 );
1011-
tri1.crop( 0, box.m_min.x, box, obb );
1012-
tri1.crop( 0, box.m_max.x, box, obb );
1013-
tri1.crop( 1, box.m_min.y, box, obb );
1014-
tri1.crop( 1, box.m_max.y, box, obb );
1015-
tri1.crop( 2, box.m_min.z, box, obb );
1016-
tri1.crop( 2, box.m_max.z, box, obb );
1006+
obb.m_box.grow( tri1.obb( box, matrixIndex ).m_box );
10171007
}
10181008

10191009
bool nodeEnd = triPairIndex + 1 == triPacketNode->getTrianglePairCount();

hiprt/impl/Triangle.h

Lines changed: 73 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,58 @@ namespace hiprt
3131
{
3232
class alignas( alignof( float3 ) ) Triangle
3333
{
34+
private:
35+
template <uint32_t N>
36+
struct PolygonN
37+
{
38+
static constexpr uint32_t MaxVertexCount = N;
39+
float3 m_vertices[MaxVertexCount];
40+
uint32_t m_count;
41+
};
42+
using Polygon = PolygonN<16>;
43+
44+
HIPRT_HOST_DEVICE static bool inside( const float3& p, uint32_t axis, float pos, bool isMin, const float maxExtent )
45+
{
46+
if ( isMin )
47+
return ( &p.x )[axis] >= pos - ObbEpsilon * maxExtent;
48+
else
49+
return ( &p.x )[axis] <= pos + ObbEpsilon * maxExtent;
50+
}
51+
52+
HIPRT_HOST_DEVICE static float3 intersect( const float3& a, const float3& b, uint32_t axis, float pos )
53+
{
54+
float da = ( &a.x )[axis] - pos;
55+
float db = ( &b.x )[axis] - pos;
56+
float t = da / ( da - db );
57+
return a + t * ( b - a );
58+
}
59+
60+
HIPRT_HOST_DEVICE static void
61+
clip( const Polygon& inPoly, uint32_t axis, float pos, bool isMin, Polygon& outPoly, const float maxExtent )
62+
{
63+
outPoly.m_count = 0;
64+
if ( inPoly.m_count == 0 ) return;
65+
66+
for ( uint32_t i = 0; i < inPoly.m_count; i++ )
67+
{
68+
const float3& curr = inPoly.m_vertices[i];
69+
const float3& prev = inPoly.m_vertices[( i + inPoly.m_count - 1 ) % inPoly.m_count];
70+
bool currIn = inside( curr, axis, pos, isMin, maxExtent );
71+
bool prevIn = inside( prev, axis, pos, isMin, maxExtent );
72+
73+
if ( currIn )
74+
{
75+
if ( !prevIn ) outPoly.m_vertices[outPoly.m_count++] = intersect( prev, curr, axis, pos );
76+
if ( outPoly.m_count < Polygon::MaxVertexCount ) outPoly.m_vertices[outPoly.m_count++] = curr;
77+
}
78+
else if ( prevIn )
79+
{
80+
if ( outPoly.m_count < Polygon::MaxVertexCount )
81+
outPoly.m_vertices[outPoly.m_count++] = intersect( prev, curr, axis, pos );
82+
}
83+
}
84+
}
85+
3486
public:
3587
Triangle() = default;
3688

@@ -109,44 +161,28 @@ class alignas( alignof( float3 ) ) Triangle
109161
rightBox.intersect( box );
110162
}
111163

112-
template <typename BoundingVolume>
113-
HIPRT_HOST_DEVICE void
114-
crop( const uint32_t axis, const float position, const Aabb& box, BoundingVolume& boundingVolume ) const
164+
HIPRT_HOST_DEVICE Obb obb( const Aabb& box, uint32_t matrixIndex ) const
115165
{
116-
// use enlarged box to make sure that the split points are inside
117-
Aabb enlargedBox = box;
118-
enlargedBox.m_min -= ObbEnlargeEpsilon * box.extent();
119-
enlargedBox.m_max += ObbEnlargeEpsilon * box.extent();
120-
121-
Aabb croppedBox;
122-
const float3* vertices = &m_v0;
123-
const float3* v1 = &vertices[2];
124-
for ( uint32_t i = 0; i < 3; i++ )
125-
{
126-
const float3* v0 = v1;
127-
v1 = &vertices[i];
128-
const float v0p = ( &v0->x )[axis];
129-
const float v1p = ( &v1->x )[axis];
130-
131-
if ( enlargedBox.contains( *v0 ) )
132-
{
133-
boundingVolume.grow( *v0 );
134-
croppedBox.grow( *v0 );
135-
}
136-
137-
if ( ( v0p < position && v1p > position ) || ( v0p > position && v1p < position ) )
138-
{
139-
const float3 t = mix( *v0, *v1, clamp( ( position - v0p ) / ( v1p - v0p ), 0.0f, 1.0f ) );
140-
if ( enlargedBox.contains( t ) )
141-
{
142-
boundingVolume.grow( t );
143-
croppedBox.grow( t );
144-
}
145-
}
146-
}
147-
148-
// fallback when the cropped box is invalid
149-
if ( !croppedBox.valid() ) boundingVolume.grow( box );
166+
Polygon poly0, poly1;
167+
poly0.m_vertices[0] = m_v0;
168+
poly0.m_vertices[1] = m_v1;
169+
poly0.m_vertices[2] = m_v2;
170+
poly0.m_count = 3;
171+
poly1.m_count = 0;
172+
173+
const float maxExtent = fmaxf( fmaxf( box.extent().x, box.extent().y ), box.extent().z );
174+
clip( poly0, 0, box.m_min.x, true, poly1, maxExtent );
175+
clip( poly1, 0, box.m_max.x, false, poly0, maxExtent );
176+
clip( poly0, 1, box.m_min.y, true, poly1, maxExtent );
177+
clip( poly1, 1, box.m_max.y, false, poly0, maxExtent );
178+
clip( poly0, 2, box.m_min.z, true, poly1, maxExtent );
179+
clip( poly1, 2, box.m_max.z, false, poly0, maxExtent );
180+
181+
Obb obb( matrixIndex );
182+
for ( uint32_t i = 0; i < poly0.m_count; i++ )
183+
obb.grow( poly0.m_vertices[i] );
184+
185+
return obb;
150186
}
151187

152188
HIPRT_HOST_DEVICE Aabb aabb() const

0 commit comments

Comments
 (0)