Cycles: HIP-RT for AMD hardware ray-tracing #105538
Merged
Brecht Van Lommel
merged 3 commits from 2023-05-15 13:46:13 +02:00
salipour/AMD_HIPRT:hip-raytracing
into main
Reviewers
Request review
No reviewers
Labels
Clear labels
This issue affects/is about backward or forward compatibility
Issues relating to security: https://wiki.blender.org/wiki/Process/Vulnerability_Reports
Apply labels
Interest
Alembic
Interest
Animation & Rigging
Interest
Asset Browser
Interest
Asset Browser Project Overview
Interest
Audio
Interest
Automated Testing
Interest
Blender Asset Bundle
Interest
BlendFile
Interest
Collada
Interest
Compatibility
This issue affects/is about backward or forward compatibility
Interest
Compositing
Interest
Core
Interest
Cycles
Interest
Dependency Graph
Interest
Development Management
Interest
EEVEE & Viewport
Interest
Freestyle
Interest
Geometry Nodes
Interest
Grease Pencil
Interest
ID Management
Interest
Images & Movies
Interest
Import Export
Interest
Line Art
Interest
Masking
Interest
Metal
Interest
Modeling
Interest
Modifiers
Interest
Motion Tracking
Interest
Nodes & Physics
Interest
OpenGL
Interest
Overrides
Interest
Performance
Interest
Physics
Interest
Pipeline, Assets & IO
Interest
Platforms, Builds & Tests
Interest
Python API
Interest
Render & Cycles
Interest
Render Pipeline
Interest
Sculpt, Paint & Texture
Interest
Text Editor
Interest
Translations
Interest
Triaging
Interest
Undo
Interest
USD
Interest
User Interface
Interest
UV Editing
Interest
VFX & Video
Interest
Video Sequencer
Interest
Virtual Reality
Interest
Vulkan
Interest: Wayland
Legacy
Blender 2.8 Project
Legacy
Milestone 1: Basic, Local Asset Browser
Legacy
OpenGL Error
Meta
Good First Issue
Meta
Papercut
Meta
Retrospective
Meta
Security
Issues relating to security: https://wiki.blender.org/wiki/Process/Vulnerability_Reports
Module
Animation & Rigging
Module
Core
Module
Development Management
Module
EEVEE & Viewport
Module
Grease Pencil
Module
Modeling
Module
Nodes & Physics
Module
Pipeline, Assets & IO
Module
Platforms, Builds & Tests
Module
Python API
Module
Render & Cycles
Module
Sculpt, Paint & Texture
Module
Triaging
Module
User Interface
Module
VFX & Video
Platform
FreeBSD
Platform
Linux
Platform
macOS
Platform
Windows
Priority
High
Priority
Low
Priority
Normal
Priority
Unbreak Now!
Status
Archived
Status
Confirmed
Status
Duplicate
Status
Needs Info from Developers
Status
Needs Information from User
Status
Needs Triage
Status
Resolved
Type
Bug
Type
Design
Type
Known Issue
Type
Patch
Type
Report
Type
To Do
No Label
Interest
Alembic
Interest
Animation & Rigging
Interest
Asset Browser
Interest
Asset Browser Project Overview
Interest
Audio
Interest
Automated Testing
Interest
Blender Asset Bundle
Interest
BlendFile
Interest
Collada
Interest
Compatibility
Interest
Compositing
Interest
Core
Interest
Cycles
Interest
Dependency Graph
Interest
Development Management
Interest
EEVEE & Viewport
Interest
Freestyle
Interest
Geometry Nodes
Interest
Grease Pencil
Interest
ID Management
Interest
Images & Movies
Interest
Import Export
Interest
Line Art
Interest
Masking
Interest
Metal
Interest
Modeling
Interest
Modifiers
Interest
Motion Tracking
Interest
Nodes & Physics
Interest
OpenGL
Interest
Overrides
Interest
Performance
Interest
Physics
Interest
Pipeline, Assets & IO
Interest
Platforms, Builds & Tests
Interest
Python API
Interest
Render & Cycles
Interest
Render Pipeline
Interest
Sculpt, Paint & Texture
Interest
Text Editor
Interest
Translations
Interest
Triaging
Interest
Undo
Interest
USD
Interest
User Interface
Interest
UV Editing
Interest
VFX & Video
Interest
Video Sequencer
Interest
Virtual Reality
Interest
Vulkan
Interest: Wayland
Legacy
Blender 2.8 Project
Legacy
Milestone 1: Basic, Local Asset Browser
Legacy
OpenGL Error
Meta
Good First Issue
Meta
Papercut
Meta
Retrospective
Meta
Security
Module
Animation & Rigging
Module
Core
Module
Development Management
Module
EEVEE & Viewport
Module
Grease Pencil
Module
Modeling
Module
Nodes & Physics
Module
Pipeline, Assets & IO
Module
Platforms, Builds & Tests
Module
Python API
Module
Render & Cycles
Module
Sculpt, Paint & Texture
Module
Triaging
Module
User Interface
Module
VFX & Video
Platform
FreeBSD
Platform
Linux
Platform
macOS
Platform
Windows
Priority
High
Priority
Low
Priority
Normal
Priority
Unbreak Now!
Status
Archived
Status
Confirmed
Status
Duplicate
Status
Needs Info from Developers
Status
Needs Information from User
Status
Needs Triage
Status
Resolved
Type
Bug
Type
Design
Type
Known Issue
Type
Patch
Type
Report
Type
To Do
Milestone
Set milestone
Clear milestone
No items
No Milestone
Projects
Set Project
Clear projects
No project
Assignees
Assign users
Clear assignees
No Assignees
12 Participants
Notifications
Due Date
The due date is invalid or out of range. Please use the format 'yyyy-mm-dd'.
No due date set.
Dependencies
No dependencies set.
Reference: blender/blender#105538
Reference in New Issue
There is no content yet.
Delete Branch "salipour/AMD_HIPRT:hip-raytracing"
Deleting a branch is permanent. Although the deleted branch may exist for a short time before cleaning up, in most cases it CANNOT be undone. Continue?
HIPRT enables AMD hardware ray tracing on RDNA2 and above and fallbacks to shader implementation for graphic cards that support HIP but not hardware ray tracing.
The ray tracing feature functions are accessed through HIPRT SDK (available on GPUOpen). HIPRT SDK allows the developer to build the bvh on the GPU with different methods and trade-offs (fast, high quality, balanced). It also allows for user defined primitives and custom intersection functions. The device side of the SDK provides traversal functions. HIPRT traversal functionality is pre-compiled in bitcode format and is shipped with HIPRT SDK. Blender kernels are compiled with hiprt headers and then linked with hiprt bitcode that has the implementation of traversal and intersection functions.
HIPRTDevice and HIPRTDeviceQueue, derived from HIP, implement new functions or override existing functions to enable HIP ray tracing on the host side.
HIPRT offers an average improvement of 25% in sample rendering rate.
Benchmark test results (sample/second) and total rendering time are measured on W6800.
There's a known issue with motion blur (visual artifact) at the moment. All the test in Cycles regression tests that pass on HIP pass on HIP RT as well.
One of the buffers allocated for HIPRT (used in the event of stack overflow during traversal) has a static size. Ideally, the buffer has to be allocated dynamically.BTW This requires Adrenalin 23.3.1 on windows or newer https://www.amd.com/en/support/kb/release-notes/rn-rad-win-23-3-1
WIP-hip-raytracingto WIP: Cycles HIP-RT for hardware ray-tracingHaven't gotten through all the code yet, but some initial comments.
@ -252,0 +272,4 @@
/* Expected in C:/Windows/System32 or similar, no path needed. */
const char* hip_paths[] = {"amdhip64.dll", NULL};
const char *hiprtc_paths[] = {"hiprtc0504.dll",
"hiprtc0503.dll",
What do these numbers in the name mean? Are they a version number that will need to be updated for new drivers?
@ -255,2 +281,3 @@
#else
const char *hip_paths[] = {"libamdhip64.so", "/opt/rocm/hip/lib/libamdhip64.so", NULL};
const char *hip_paths[] = {"/opt/rocm/hip/lib/libamdhip64.so", NULL};
const char* hiprtc_paths[] = { "/opt/rocm/hip/lib/libhiprtc.so", NULL };
Can you leave in the libraries without absolute path?
This helps for Linux distribution that may package these libraries themselves and put them in a location in the library path.
@ -412,0 +459,4 @@
HIPRTC_LIBRARY_FIND_CHECKED( hiprtcLinkAddFile );
HIPRTC_LIBRARY_FIND_CHECKED( hiprtcLinkAddData );
HIPRTC_LIBRARY_FIND_CHECKED( hiprtcLinkComplete );
HIPRTC_LIBRARY_FIND_CHECKED( hiprtcLinkDestroy );
Various changes in this file seems to have inconsistent indentation and space around ().
@ -1545,1 +1545,4 @@
use_hiprt: BoolProperty(
name="HIPRT (Experimental)",
description="HIPRT enables AMD hardware ray tracing on RDNA2 and above. However this support is experimental and some scenes may render incorrectly",
Here it's called "HIP RT" rather than "HIPRT", so maybe that's better for a name in the user interface?
https://gpuopen.com/learn/introducing-hiprt/
This option also affects graphics cards before RDNA2, so maybe the text can be tweaked to clarify that.
@ -16,2 +16,4 @@
endif()
if(WITH_CYCLES_DEVICE_HIPRT)
set (HIPRTSDK $ENV{HIPRT_SDK_DIR})
We should not rely on environment variables, anything that needs to be configured needs to be a CMake variable.
@ -17,1 +17,4 @@
if(WITH_CYCLES_DEVICE_HIPRT)
set (HIPRTSDK $ENV{HIPRT_SDK_DIR})
list(APPEND INC
Indentation is inconsistent here.
@ -0,0 +98,4 @@
bool HIPRTDevice::set_function_table(hiprtFuncNameSet *func_name_set)
{
Code style: no empty line at the start of functions here and elsewhere.
@ -0,0 +111,4 @@
"motion_triangle_custom_intersect",
"point_custom_intersect"};
//"motion_triangle_custom_local_intersect", "motion_triangle_custom_volume_intersect"
For commented out code, please leave a comment explaining why it is so.
@ -0,0 +122,4 @@
func_name_set[table_index].intersectFuncName = intersect_function[prim];
}
else if (prim == Triangle)
// triangle primitives dont need a custom intersection function
Code style for comments, write as full sentence (here and elsewhere):
@ -151,2 +151,4 @@
};
LOG(FATAL) << "Unhandled kernel " << static_cast<int>(kernel) << ", should never happen.";
Leave out whitspace changes.
@ -1145,2 +1145,4 @@
case DEVICE_HIP:
return "HIP";
case DEVICE_HIPRT:
return "HIP-RT";
It's
"HIPRT"
in other places, let's pick one.@ -1003,2 +1099,4 @@
add_dependencies(cycles_kernel cycles_kernel_oneapi)
endif()
if(WITH_CYCLES_DEVICE_HIPRT)
# -D __HIPRT__
Remove code if it's not needed.
@ -48,6 +48,7 @@ ccl_device_inline int3 make_int3(int x, int y, int z);
#endif /* __KERNEL_NATIVE_VECTOR_TYPES__ */
ccl_device_inline int3 make_int3(int i);
Leave out whitespace changes.
@ -1255,0 +1256,4 @@
set (HIPRTSDK $ENV{HIPRT_SDK_DIR})
#Do not use environment variables for this, these kinds of paths should be specified through cmake options.
target_link_libraries(blender ${HIPRTSDK}/hiprt${HIPRT_VERSION}64.lib)
endif()
Can we leave this out now and only dynamically load libraries through hipew?
@blender-bot build
I think I got through all the code now.
For reference the previous review was here before we moved to projects.blender.org:
https://archive.blender.org/developer/D17098
I've repeated some of the comments that were not addressed yet.
@ -0,0 +22,4 @@
class Object;
class BVHHIPRT;
static void get_hiprt_transform(float matrix[][4], Transform &tfm);
This doesn't need to be declared in the header. I think you can just do a static definition inside the .cpp file without any declaration?
@ -0,0 +432,4 @@
type = segment.type;
prim = segment.prim;
/*if (u == 0.0f || u == 1.0f) {
Why was this commented out? Can you leave a comment explaining this if it's intentional?
@ -0,0 +602,4 @@
case 3:
case 7:
case 11:
case 15:
Can we avoid using magic numbers here somehow?
@ -0,0 +11,4 @@
#define HIPRT_SHARED_STACK
#define HIPRT_GLOBAL_STACK_SIZE 512 * 1024 * 1024
It's unclear to me where these numbers come from. At a minimum it should get a comment, but it sounds like something to dynamically compute based on device specs?
@ -0,0 +17,4 @@
64 // global stack allocation per thread (memory reserved for each thread in
// global_stack_buffer)
#define HIPRT_THREAD_GROUP_SIZE \
256 // total locaal stack size would be number of threads * HIPRT_SHARED_STACK_SIZE
locaal -> local
@ -0,0 +18,4 @@
// global_stack_buffer)
#define HIPRT_THREAD_GROUP_SIZE \
256 // total locaal stack size would be number of threads * HIPRT_SHARED_STACK_SIZE
Also for these other numbers, why are they these values? It's not clear to me if they would need to be updated if we make some change in the kernel.
@ -0,0 +1,271 @@
#pragma once
Can you add a comment explaining the purpose of this file?
It seems to define a bunch of common types and utility functions, it's not really clear to me why rather than using existing Cycles equivalents or things that come with the HIP headers. Are these meant to stay in sync with the HIP RT API somehow, and if so how exactly do we do that?
@ -0,0 +1,271 @@
#pragma once
This and a few following files are missing SPDX license headers.
@ -0,0 +1,271 @@
#pragma once
#if (defined(__CUDACC__) || defined(__HIPCC__))
# define __KERNELCC__
Can this get a better name? It's not really clear to me what the purpose is, why there is CUDA in this header at all.
@ -0,0 +1,384 @@
#pragma once
Missing license header.
@ -0,0 +1,384 @@
#pragma once
The contents of this file feels like it should be in hipew more than Cycles itself? If it's meant to match the HIP API somehow.
The code in this branch is based on a somewhat older version of main which does not work on the buildbot. It will need to be merged with main so we can verify if it builds correctly, though there is no rush for that as we can do it after review comments are addressed also.
649f2be265
to053c90af19
From what I can tell some comments were addressed but others not yet. I collapsed all the ones that seem to be done.
@ -0,0 +18,4 @@
// defined in HIPRT_THREAD_STACK_SIZE Currently allocating this buffer dynamically causes
// instability, so HIPRT_GLOBAL_STACK_SIZE and maximum number of threads is hard coded The total
// size will be (max numer of threads) x (thread stack size) x sizeof(int) x (headroom factor) A
// headroom factor of 2 is emperical
Thanks for the explanation. Is the plan to merge it with these hardcoded numbers, or is the instability from dynamic allocation planned to be solved still?
It feels like this is likely to break on newer hardware or compilers, or give excessive memory usage on older hardware. So I'm not sure that we should ship it like this.
@ -0,0 +20,4 @@
// size will be (max numer of threads) x (thread stack size) x sizeof(int) x (headroom factor) A
// headroom factor of 2 is emperical
#define HIPRT_THREAD_STACK_SIZE \
64 // The size of global stack availavle to eacj thread (memory reserved for each thread in
spelling: availavle, eacj, exectuion, emperical
@ -252,0 +278,4 @@
// Currently, kernels are compiled through hipcc, but if runtime compiler is used
// the hiprtc version should match the comgr version
const char *hiprtc_paths[] = {"hiprtc0504.dll",
"hiprtc0503.dll",
Thanks for the explanation.
But if these are associated with a particular compiler version, this is not future proof? A Blender release made now would not work with future driver versions that have a newer compiler version?
@ -252,0 +272,4 @@
/* Expected in C:/Windows/System32 or similar, no path needed. */
const char* hip_paths[] = {"amdhip64.dll", NULL};
const char *hiprtc_paths[] = {"hiprtc0504.dll",
"hiprtc0503.dll",
The numbers refer to compiler versions and hiprtc is only used if run-time compiler is used. Current implementation only uses hipcc.
@ -252,0 +278,4 @@
// Currently, kernels are compiled through hipcc, but if runtime compiler is used
// the hiprtc version should match the comgr version
const char *hiprtc_paths[] = {"hiprtc0504.dll",
"hiprtc0503.dll",
It won't be a problem if blender uses an older version, and the driver has a newer compiler. There is no reliance on the driver to pick the compiler version. The compiler is shipped with HIP SDK. If blender starts using a newer version of the SDK, then hiprtc version should be updated to match the compiler version. Also, right now hiprtc is not used for any kernel because all the compilation is done through hipcc (which is independent of hiprtc).
@ -0,0 +18,4 @@
// defined in HIPRT_THREAD_STACK_SIZE Currently allocating this buffer dynamically causes
// instability, so HIPRT_GLOBAL_STACK_SIZE and maximum number of threads is hard coded The total
// size will be (max numer of threads) x (thread stack size) x sizeof(int) x (headroom factor) A
// headroom factor of 2 is emperical
The instability issue will be eventually addressed but for the time being our plan is to ship it like this. Your point is absolutely valid, and we are working to resolve the issue.
@ -0,0 +1,271 @@
#pragma once
hiprt_common.h, hiprt_device.h, hiprt_types.h, and hiprt_vec.h are all headers from HIP RT SDK. In the latest code changes, these headers were removed and instead the include path passed to the compiler is updated to point to the location of these files.
@ -0,0 +18,4 @@
// defined in HIPRT_THREAD_STACK_SIZE Currently allocating this buffer dynamically causes
// instability, so HIPRT_GLOBAL_STACK_SIZE and maximum number of threads is hard coded The total
// size will be (max number of threads) x (thread stack size) x sizeof(int) x (headroom factor) A
// headroom factor of 2 is empirical
The instability issue will be addressed eventually. In the meantime the plan is to leave the buffer allocation as is.
@ -252,0 +278,4 @@
// Currently, kernels are compiled through hipcc, but if runtime compiler is used
// the hiprtc version should match the comgr version
const char *hiprtc_paths[] = {"hiprtc0504.dll",
"hiprtc0503.dll",
From what I can tell, no
hiprtc
functions are used in this patch, and so this code doesn't actually do anything besides loading a library that is not used.I don't think we should load any library we don't need, because it always carries a risk of breaking something. I think it would make more sense to have a separate hipew API function or a parameter to make loading of this optional, and only do it when needed.
If we would ship with runtime compilation on user machines, then having it tied to a specific version number still seems wrong. If we use this for offline compilation, it's also wrong to look for these dlls in arbitrary system paths. In that case it should be finding these in a path to an SDK provided through a CMake variable.
I installed the hiprt SDK on the buildbot now. The following code needs to be added to pick it up:
For the performance, it would be good to have a breakdown for the benchmarking scenes, since 25% is a bit vague. We have a benchmarking script that can help.
Are the regression tests passing besides the motion blur issue? Assuming you have the test files downloaded, the configuration and running would be something like this:
I have not checked recently if the regular HIP tests are passing, but the CUDA ones are passing here.
There are some compile errors on Linux.
https://builder.blender.org/admin/#/builders/132/builds/956/steps/3/logs/stdio
The code there seems to be Windows specific currently.
WIP: Cycles HIP-RT for hardware ray-tracingto Cycles: HIP-RT for AMD hardware ray-tracingThis warning I can only fix in the SDK header files, would be good if it got fixed in the SDK. For now I've edited the header file locally.
It built successfully now, the build is available for testing here:
https://builder.blender.org/download/patch/PR105538/
There are these warnings that seem like they might be a real problem:
HIP-RT seems to be always OFF in this build even when the option is selected in the systems settings? At least that's what I think because I'm getting identical perfs with it ON or Off (also same perfs as Blender 3.5). Radeon 6700XT Driver 23.4.1
I'm seeing the same here, no performance difference with HIPRT on or off. Tried driver versions 23.3.2 and 23.4.1, identical performance with HIPRT on/off on both drivers.
OS: Windows 11
GPU: Radeon RX 6800XT
Performance has become worse compared to 3.5
Hello all. Before reporting performance we need to upload an updated driver. I might upload a package you can install to test next week thought.
@BrianSavery, @salipour, there was an error in the previous build where it was not using HIP-RT. I fixed that now, but unfortunately it's now not giving correct renders for me with an 7900 XT. It shows black or some kind of corrupted geometry, hard to make anything out. I tried to default cube and a few benchmark scenes.
I suspected it might be my changes or a change in main, but it does not seem so. In #107224 and
08ebdb102d
I made just the minimal changes to get things to build for me, and it has the same problem.I see
removing function: +dot1-insts is not supported on the current target
when compiling for gfx1100 for the 7900 XT. So maybe that warning is related? Not sure if you tested on another generation card that might work?I seem to be getting a different issue entirely on my 6600 XT (Navi 23). It could be an issue with my setup (I'm using Adrenalin 23.4.2 on Windows 11 22624.1470).
It can't create the HIPRT context regardless of scene. I have tried on default cube and also more complex scenes.
Yup. Same error.
Thanks for all the testing. However, every user out there, you do NOT have the correct driver yet, so please hold off reporting anything else until we give an update of which driver to use (Will be coming in a driver update next week). Thanks.
@brecht I'll reply in blender.chat
I intend to commit the code tomorrow, while still leaving it disabled in daily builds. Then it's easier to do further changes and fixes without having to keep this big PR up to date.
The code landed now, and I merged with main. All that remains in this PR now is the change to actually enable it in daily/release builds.
Marking as request changes, regarding the render errors.
@blender-bot package windows
Package build started. Download here when ready.
For users wanting to test this, Please use the driver here: https://www.amd.com/en/support/kb/release-notes/rn-rad-win-22-40-51-06-blender-3-6-beta
Which Blender 3.6 Beta build should be used ?
Because this doesn't seem to work: https://builder.blender.org/download/patch/PR105538/
It either crashes when attempting to render with HIP-RT enabled or renders nothing/nonsense.
6700XT with the 22.40.51.drivers posted above.
9da482491a
toe0ec48535b
@BrianSavery I'm still getting the same issue as before with the new SDK and driver, testing on a 7900 XT. Also still seeing the same build warnings.
https://builder.blender.org/admin/#/builders/133/builds/1034/steps/4/logs/stdio
@Pliou the build was not updated yet when you tried it. It is now, but not super likely it will work for you if it didn't for me, though it's possible since you have a different GPU model.
@BrianSavery
@brecht Indeed. Just tried the latest build & I'm getting exactly the same crashes & messed rendering unfortunately.
With the beta driver and the new build on an RX6800, it's creating the HIPRT structures, but then freezes on the sampling.
In the sky demo this eventually causes a driver timeout.
In the 3.5 splash demo I'm getting a memory error:
Out of memory in hipStreamCreateWithFlags(&hip_stream_, hipStreamNonBlocking) (C:\Users\blender\git\blender-vexp\blender.git\intern\cycles\device\hip\queue.cpp:20)
Neither offline nor windowed rendering renders results
There's a new build now at the same link, but I have not tested it myself yet.
It now appears to render correctly, though it performs worse than HIP in the two renders I tried on my 6800XT.
My 6800xt renders standard classroom files (300) in 30s why does it take you 2 minutes? That's slower than 2.93 opencl.
@Bence-Koczogh Something's wrong on your side. Your 6800XT shouldn't be that slow even using HIP.
Classroom renders in 49 seconds on my 6700XT using HIP (on Blender 3.5) & 43 secs using HIPRT on 3.6.
Pavillon renders in 1:37 Blender 3.5 & 1:11 with HIPRT in 3.6
Scanlands : 4:01 minutes in Blender 3.5 & 2:34 with HIPRT in 3.6
@brecht
Some scenes only see a 5%-10% boost while others are up to +35% faster ...That's when it doesn't either crash Blender or the display driver when building the BVH. Scenes with geometry nodes seems to crash also.
@salipour @brecht Lighting/shading doesn't match in the Laundromat scene (see screenshots below)..
Looks like a regression. I checked the visual results from earlier this year, and the lighting is correct. Will track it down.
I 6800 xt rendering classroom 28s, classroom scene enhancements seem to be only 10 percent, (Note: 3.6 than the stable 3.5 version slow 1s,)
@salipour some scenes crash when building the BHV (try this one : https://cloud.blender.org/p/gallery/6220ae43b4a486f53171c89e) also when Geometry nodes are present.
@2905710881 try rendering Scanlands in 3.5 & then 3.6 with HIPRT https://www.blender.org/download/demo-files/
Crash is probably related to running out of GPU memory. Can you monitor task manager and see if the crash happens when GPU memory is all used up?
I did notice it was very slow, but I don't do these benchmark scenes often so I don't know what numbers to expect. I rendered those two scenes using a clean install of the driver posted above, but I'll have to check my setup before doing any more HIPRT testing.
I just ran the same renders in 3.5.1 on Linux and I got numbers much closer to yours (38s Classroom, 1:07 Pavillon).
Nope plenty of memory left. Had to take a photo of my screen because my PC was unresponsive (mouse pointer works but that's all..had to reset).
All scenes with geometry node simply crash/close Blender when attempting to build the BVH.
@salipour Also this scene crashes Blender when trying to build BVH using HIP or HIPRT (doesn't render in Blender 3.5 HIP either..)
https://cloud.blender.org/p/gallery/60337d495677e942564cce76
@Pliou if a file causes issues in the Main or release branches of Blender, then please make a bug report about it.
As you pointed out yourself, the issue you're experiencing isn't caused by HIP-RT, it's caused by something else. So it doesn't need to be reported on the HIP-RT pull request.
@salipour @BrianSavery can you clarify what the status is here?
What the bugs are, if they are important enough to block enabling this in main now, what is expected to be fixed still, etc. Any remaining issues and todo's should be listed in #104110 so we can keep track of them.
We'll also need some text for the release notes, can you provide that? It should then also includes any known limitations.
I don't think there is anything substantial blocking the enablement. There are a couple of regressions that we are tracking down.
Can you:
This is now in main, thanks @salipour and @BrianSavery for the contribution.
Please see the release notes for details and known issues, and report any bugs not listed there to the tracker.
https://wiki.blender.org/wiki/Reference/Release_Notes/3.6/Cycles
The latest release of the HIP-RT SDK (2.0.3, from a week ago) finally ships with Linux libraries, so I guess this could also be enabled on Linux now?
We are looking to enable it for Blender 4.0 in Linux so users don't have to install the HIP RT SDK themselves.
I can't get this HIP SDK to build on linux no matter what, any hint on how to enable hip rt for linux? Already have latest ROcm / Hip installed.
It's not meant to work on Linux yet, there are build instructions that can be given at this moment.
Reviewers