blender/build_files/build_environment/patches/hiprt_rocm7_support.diff
Jonas Holzman cd754c5ccc Deps: Library changes for Blender 5.3
New libraries:
* KTX
* OpenTimelineIO
* Jolt
* hwloc
* RapidJSON

Updated libraries:
* LLVM
* OSL
* OpenImageIO
* OpenUSD
* DPC++
* minizip-ng
* flac

Build changes for existing libraries:
* OpenEXR (build Python bindings)
* FFmpeg (add additional hardware acceleration backends)
* Enable CUDA/OptiX support for Windows ARM64 (LLVM, OSL, OIDN)

Removed libraries:
* fastjsonschema

For a full list of upgraded dependencies, specific versions and additional notes, see the main library change tracking issue below.

Pull Request: https://projects.blender.org/blender/blender/pulls/162972

Ref: #158444

Co-authored-by: Thomas Dinges <thomas@blender.org>
Co-authored-by: Anthony Roberts <anthony.roberts@linaro.org>
Co-authored-by: Ray Molenkamp <github@lazydodo.com>
Co-authored-by: Brecht Van Lommel <brecht@blender.org>
Co-authored-by: Campbell Barton <campbell@blender.org>
Co-authored-by: Sebastian Parborg <sebastian@blender.org>
Co-authored-by: Patrick Mours <pmours@nvidia.com>
Co-authored-by: Nikita Sirgienko <nikita.sirgienko@intel.com>
Co-authored-by: John Kiril Swenson <john@blender.org>
2026-10-01 18:21:55 +02:00

90 lines
3.8 KiB
Diff

diff --git a/hiprt/hiprt_libpath.h b/hiprt/hiprt_libpath.h
index 58daa96..88cbcf2 100644
--- a/hiprt/hiprt_libpath.h
+++ b/hiprt/hiprt_libpath.h
@@ -34,7 +34,7 @@
#ifdef HIPRT_PREFER_HIP_5
const char* g_hip_paths[] = { "amdhip64.dll", "amdhip64_6.dll", NULL };
#else
-const char* g_hip_paths[] = { "amdhip64_6.dll", "amdhip64.dll", NULL };
+const char* g_hip_paths[] = { "amdhip64_7.dll", "amdhip64_6.dll", "amdhip64.dll", NULL };
#endif
const char* g_hiprtc_paths[] = {
@@ -72,9 +72,13 @@ const char* g_hip_paths[] = {
"/opt/rocm/lib/libamdhip64.so",
"/opt/rocm/hip/lib/libamdhip64.so",
- "libamdhip64.so.5",
- "/opt/rocm/lib/libamdhip64.so.5",
- "/opt/rocm/hip/lib/libamdhip64.so.5",
+ "libamdhip64.so.7",
+ "/opt/rocm/lib/libamdhip64.so.7",
+ "/opt/rocm/hip/lib/libamdhip64.so.7",
+
+ "libamdhip64.so.6",
+ "/opt/rocm/lib/libamdhip64.so.6",
+ "/opt/rocm/hip/lib/libamdhip64.so.6",
NULL };
#endif
diff --git a/hiprt/impl/hiprt_device_impl.h b/hiprt/impl/hiprt_device_impl.h
index 0687234..272594d 100644
--- a/hiprt/impl/hiprt_device_impl.h
+++ b/hiprt/impl/hiprt_device_impl.h
@@ -37,7 +37,8 @@
#include <hiprt/hiprt_device.h>
#if __gfx1030__ || __gfx1031__ || __gfx1032__ || __gfx1033__ || __gfx1034__ || __gfx1035__ || __gfx1036__ || __gfx1100__ || \
- __gfx1101__ || __gfx1102__ || __gfx1103__ || __gfx1150__ || __gfx1151__ || __gfx1152__ || __gfx1200__ || __gfx1201__
+ __gfx1101__ || __gfx1102__ || __gfx1103__ || __gfx1150__ || __gfx1151__ || __gfx1152__ || __gfx1153__ || \
+ __gfx1200__ || __gfx1201__
#ifndef __USE_HWI__
#define __USE_HWI__
#endif
@@ -362,10 +363,10 @@ TraversalBase<Stack>::testInternalNode( const hiprtRay& ray, const float3& invD,
auto result = __builtin_amdgcn_image_bvh_intersect_ray_l(
encodeBaseAddr( nodes, nodeIndex ),
ray.maxT,
- float4{ ray.origin.x, ray.origin.y, ray.origin.z, 0.0f }.data,
- float4{ ray.direction.x, ray.direction.y, ray.direction.z, 0.0f }.data,
- float4{ invD.x, invD.y, invD.z, 0.0f }.data,
- m_descriptor.data );
+ { ray.origin.x, ray.origin.y, ray.origin.z, 0.0f },
+ { ray.direction.x, ray.direction.y, ray.direction.z, 0.0f },
+ { invD.x, invD.y, invD.z, 0.0f },
+ { m_descriptor.x, m_descriptor.y, m_descriptor.z, m_descriptor.w } );
#endif
if ( m_stack.vacancy() < 3 )
{
@@ -397,11 +398,13 @@ HIPRT_DEVICE bool TraversalBase<Stack>::testTriangleNode(
hit.normal = node.m_triPair.fetchTriangle( leafIndex & 1 ).normal( node.m_flags >> ( ( leafIndex & 1 ) * 8 ) );
}
#else
- const float4 origin = float4{ ray.origin.x, ray.origin.y, ray.origin.z, 0.0f };
- const float4 direction = float4{ ray.direction.x, ray.direction.y, ray.direction.z, 0.0f };
- const float4 invDir = float4{ invD.x, invD.y, invD.z, 0.0f };
- auto result = __builtin_amdgcn_image_bvh_intersect_ray_l(
- encodeBaseAddr( nodes, leafIndex ), ray.maxT, origin.data, direction.data, invDir.data, m_descriptor.data );
+ auto result = __builtin_amdgcn_image_bvh_intersect_ray_l(
+ encodeBaseAddr( nodes, leafIndex ),
+ ray.maxT,
+ { ray.origin.x, ray.origin.y, ray.origin.z, 0.0f },
+ { ray.direction.x, ray.direction.y, ray.direction.z, 0.0f },
+ { invD.x, invD.y, invD.z, 0.0f },
+ { m_descriptor.x, m_descriptor.y, m_descriptor.z, m_descriptor.w } );
float invDenom = __ocml_native_recip_f32( __int_as_float( result[1] ) );
float t = __int_as_float( result[0] ) * invDenom;
hasHit = ray.minT <= t && t <= ray.maxT;
diff --git a/scripts/bitcodes/common_tools.py b/scripts/bitcodes/common_tools.py
index 10604e5..bd46b5f 100644
--- a/scripts/bitcodes/common_tools.py
+++ b/scripts/bitcodes/common_tools.py
@@ -58,6 +58,7 @@ def getAMDGPUArchs(hip_sdk_version_num):
if hip_sdk_version_num >= 63:
gpus_archs.append('gfx1152')
+ gpus_archs.append('gfx1153')
if hip_sdk_version_num >= 62: # Navi4 supported from 6.2
gpus_archs.append('gfx1200')