Cycles: HIP-RT for AMD hardware ray-tracing #105538

Merged
Brecht Van Lommel merged 3 commits from salipour/AMD_HIPRT:hip-raytracing into main 2023-05-15 13:46:13 +02:00
Member

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.

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.~~
Brian Savery (AMD) requested review from Brecht Van Lommel 2023-03-07 18:20:22 +01:00

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

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
Brecht Van Lommel changed title from WIP-hip-raytracing to WIP: Cycles HIP-RT for hardware ray-tracing 2023-03-07 18:29:32 +01:00
Brecht Van Lommel added this to the Render & Cycles project 2023-03-07 18:29:37 +01:00
Brecht Van Lommel approved these changes 2023-03-09 13:36:30 +01:00
Dismissed
Brecht Van Lommel left a comment
Owner

Haven't gotten through all the code yet, but some initial comments.

Haven'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?

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?

  const char *hip_paths[] = {"libamdhip64.so", "/opt/rocm/hip/lib/libamdhip64.so", NULL};
  const char* hiprtc_paths[] = {"libhiprtc.so", "/opt/rocm/hip/lib/libhiprtc.so", NULL }; 

This helps for Linux distribution that may package these libraries themselves and put them in a location in the library path.

Can you leave in the libraries without absolute path? ``` const char *hip_paths[] = {"libamdhip64.so", "/opt/rocm/hip/lib/libamdhip64.so", NULL}; const char* hiprtc_paths[] = {"libhiprtc.so", "/opt/rocm/hip/lib/libhiprtc.so", NULL }; ``` This helps for Linux distribution that may package these libraries themselves and put them in a location in the library path.
salipour marked this conversation as resolved
@ -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 ().

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.

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.
brecht marked this conversation as resolved
@ -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.

We should not rely on environment variables, anything that needs to be configured needs to be a CMake variable.
brecht marked this conversation as resolved
@ -17,1 +17,4 @@
if(WITH_CYCLES_DEVICE_HIPRT)
set (HIPRTSDK $ENV{HIPRT_SDK_DIR})
list(APPEND INC

Indentation is inconsistent here.

Indentation is inconsistent here.
brecht marked this conversation as resolved
@ -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.

Code style: no empty line at the start of functions here and elsewhere.
brecht marked this conversation as resolved
@ -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.

For commented out code, please leave a comment explaining why it is so.
brecht marked this conversation as resolved
@ -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):

// Triangle primitives dont need a custom intersection function.
Code style for comments, write as full sentence (here and elsewhere): ``` // Triangle primitives dont need a custom intersection function. ```
brecht marked this conversation as resolved
@ -151,2 +151,4 @@
};
LOG(FATAL) << "Unhandled kernel " << static_cast<int>(kernel) << ", should never happen.";

Leave out whitspace changes.

Leave out whitspace changes.
brecht marked this conversation as resolved
@ -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.

It's `"HIPRT"` in other places, let's pick one.
brecht marked this conversation as resolved
@ -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.

Remove code if it's not needed.
brecht marked this conversation as resolved
@ -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.

Leave out whitespace changes.
brecht marked this conversation as resolved
@ -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?

Can we leave this out now and only dynamically load libraries through hipew?
brecht marked this conversation as resolved
Brecht Van Lommel dismissed brecht’s review 2023-03-09 13:37:17 +01:00
Reason:

Didn't intend to click approve yet

@blender-bot build

@blender-bot build
Brecht Van Lommel requested changes 2023-03-09 18:49:24 +01:00
Brecht Van Lommel left a comment
Owner

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.

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?

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?
brecht marked this conversation as resolved
@ -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?

Why was this commented out? Can you leave a comment explaining this if it's intentional?
brecht marked this conversation as resolved
@ -0,0 +602,4 @@
case 3:
case 7:
case 11:
case 15:

Can we avoid using magic numbers here somehow?

Can we avoid using magic numbers here somehow?
brecht marked this conversation as resolved
@ -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?

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?
brecht marked this conversation as resolved
@ -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

locaal -> local
brecht marked this conversation as resolved
@ -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.

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.
brecht marked this conversation as resolved
@ -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?

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?
brecht marked this conversation as resolved
@ -0,0 +1,271 @@
#pragma once

This and a few following files are missing SPDX license headers.

This and a few following files are missing SPDX license headers.
brecht marked this conversation as resolved
@ -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.

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.
brecht marked this conversation as resolved
@ -0,0 +1,384 @@
#pragma once

Missing license header.

Missing license header.
brecht marked this conversation as resolved
@ -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 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.
brecht marked this conversation as resolved

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.

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.
Sahar A. Kashi force-pushed hip-raytracing from 649f2be265 to 053c90af19 2023-03-17 03:42:02 +01:00 Compare
Brecht Van Lommel requested review from Brecht Van Lommel 2023-03-17 20:52:40 +01:00
Brecht Van Lommel requested changes 2023-03-20 16:18:19 +01:00
Brecht Van Lommel left a comment
Owner

From what I can tell some comments were addressed but others not yet. I collapsed all the ones that seem to be done.

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.

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

spelling: availavle, eacj, exectuion, emperical
brecht marked this conversation as resolved
Sahar A. Kashi added 2 commits 2023-03-21 19:38:51 +01:00
Brecht Van Lommel reviewed 2023-03-21 19:48:06 +01:00
@ -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?

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?
salipour marked this conversation as resolved
Sahar A. Kashi reviewed 2023-03-21 20:02:37 +01:00
@ -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",
Author
Member

The numbers refer to compiler versions and hiprtc is only used if run-time compiler is used. Current implementation only uses hipcc.

The numbers refer to compiler versions and hiprtc is only used if run-time compiler is used. Current implementation only uses hipcc.
salipour marked this conversation as resolved
@ -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",
Author
Member

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).

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
Author
Member

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.

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.
salipour marked this conversation as resolved
@ -0,0 +1,271 @@
#pragma once
Author
Member

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.

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.
salipour marked this conversation as resolved
Sahar A. Kashi reviewed 2023-03-21 20:12:33 +01:00
@ -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
Author
Member

The instability issue will be addressed eventually. In the meantime the plan is to leave the buffer allocation as is.

The instability issue will be addressed eventually. In the meantime the plan is to leave the buffer allocation as is.
Brecht Van Lommel requested review from Brecht Van Lommel 2023-03-22 15:39:25 +01:00
Brecht Van Lommel requested changes 2023-03-27 13:20:44 +02:00
@ -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.

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:

diff --git a/build_files/config/pipeline_config.yaml b/build_files/config/pipeline_config.yaml
index 11a095e..91f6df8 100644
--- a/build_files/config/pipeline_config.yaml
+++ b/build_files/config/pipeline_config.yaml
@@ -10,6 +10,8 @@ buildbot:
         version: '11.4.1'
     hip:
         version: '5.3.22480'
+    hiprt:
+        version: '2.0.0'
     optix:
         version: '7.3.0'
     ocloc:
I installed the hiprt SDK on the buildbot now. The following code needs to be added to pick it up: ``` diff --git a/build_files/config/pipeline_config.yaml b/build_files/config/pipeline_config.yaml index 11a095e..91f6df8 100644 --- a/build_files/config/pipeline_config.yaml +++ b/build_files/config/pipeline_config.yaml @@ -10,6 +10,8 @@ buildbot: version: '11.4.1' hip: version: '5.3.22480' + hiprt: + version: '2.0.0' optix: version: '7.3.0' ocloc: ```

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:

https://wiki.blender.org/wiki/Tools/Tests/Setup

cmake -DCYCLES_TEST_DEVICES=CPU;HIPRT .
ctest -C Release -R cycles

I have not checked recently if the regular HIP tests are passing, but the CUDA ones are passing here.

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](https://wiki.blender.org/wiki/Tools/Tests/Performance) that can help. Are the regression tests passing besides the motion blur issue? Assuming you have the [test files downloaded](https://wiki.blender.org/wiki/Tools/Tests/Setup), the configuration and running would be something like this: ``` https://wiki.blender.org/wiki/Tools/Tests/Setup cmake -DCYCLES_TEST_DEVICES=CPU;HIPRT . ctest -C Release -R cycles ``` I have not checked recently if the regular HIP tests are passing, but the CUDA ones are passing here.
Sahar A. Kashi added 1 commit 2023-03-28 02:04:41 +02:00
Brecht Van Lommel requested review from Brecht Van Lommel 2023-04-04 20:54:32 +02:00
Sahar A. Kashi added 3 commits 2023-04-11 00:32:48 +02:00

There are some compile errors on Linux.
https://builder.blender.org/admin/#/builders/132/builds/956/steps/3/logs/stdio

/home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:36:8: error: ‘HMODULE’ does not name a type
   36 | static HMODULE hiprt_lib;
      |        ^~~~~~~
/home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc: In function ‘void hipewHipRtExit()’:
/home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:54:7: error: ‘hiprt_lib’ was not declared in this scope; did you mean ‘hiprtHit’?
   54 |   if (hiprt_lib != NULL) {
      |       ^~~~~~~~~
      |       hiprtHit
/home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:56:5: error: ‘FreeLibrary’ was not declared in this scope
   56 |     FreeLibrary(hiprt_lib);
      |     ^~~~~~~~~~~

The code there seems to be Windows specific currently.

There are some compile errors on Linux. https://builder.blender.org/admin/#/builders/132/builds/956/steps/3/logs/stdio ``` /home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:36:8: error: ‘HMODULE’ does not name a type 36 | static HMODULE hiprt_lib; | ^~~~~~~ /home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc: In function ‘void hipewHipRtExit()’: /home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:54:7: error: ‘hiprt_lib’ was not declared in this scope; did you mean ‘hiprtHit’? 54 | if (hiprt_lib != NULL) { | ^~~~~~~~~ | hiprtHit /home/blender/git/blender-vexp/blender.git/extern/hipew/src/hiprtew.cc:56:5: error: ‘FreeLibrary’ was not declared in this scope 56 | FreeLibrary(hiprt_lib); | ^~~~~~~~~~~ ``` The code there seems to be Windows specific currently.
Brecht Van Lommel changed title from WIP: Cycles HIP-RT for hardware ray-tracing to Cycles: HIP-RT for AMD hardware ray-tracing 2023-04-20 18:40:23 +02:00

This 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.

In file included from hiprt2.0.0/hiprt/hiprt_types.h:8:
hiprt2.0.0/hiprt/hiprt_vec.h:12:9: warning: 'HIPRT_HOST_DEVICE' macro redefined [-Wmacro-redefined]
#define HIPRT_HOST_DEVICE
        ^
hiprt2.0.0/hiprt/hiprt_common.h:60:9: note: previous definition is here
#define HIPRT_HOST_DEVICE __host__ __device__
        ^
This 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. ``` In file included from hiprt2.0.0/hiprt/hiprt_types.h:8: hiprt2.0.0/hiprt/hiprt_vec.h:12:9: warning: 'HIPRT_HOST_DEVICE' macro redefined [-Wmacro-redefined] #define HIPRT_HOST_DEVICE ^ hiprt2.0.0/hiprt/hiprt_common.h:60:9: note: previous definition is here #define HIPRT_HOST_DEVICE __host__ __device__ ^ ```

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:

_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
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: ``` _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target ```
First-time contributor

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:

_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target
_ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target

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

> 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: > ``` > _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN35hiprtGeomTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN37hiprtSceneTraversalClosestCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > _ZN36hiprtSceneTraversalAnyHitCustomStackI16hiprtGlobalStackE10getNextHitEv: removing function: +dot1-insts is not supported on the current target > ``` 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
First-time contributor

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

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
First-time contributor

Performance has become worse compared to 3.5

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.

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 seeremoving 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?

@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 08ebdb102d1e6a84e4a9b41e018a72b034d8c514 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?
First-time contributor

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.

image

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. ![image](/attachments/0276c037-f536-4e06-92c2-8ff197ddbf74)
7.7 KiB
First-time contributor

Yup. Same error.

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

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.

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.

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.
Brecht Van Lommel requested changes 2023-04-25 21:14:03 +02:00
Brecht Van Lommel left a comment
Owner

Marking as request changes, regarding the render errors.

Marking as request changes, regarding the render errors.

@blender-bot package windows

@blender-bot package windows
Member

Package build started. Download here when ready.

Package build started. [Download here](https://builder.blender.org/download/patch/PR105538) 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

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
First-time contributor

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.

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.
Brecht Van Lommel force-pushed hip-raytracing from 9da482491a to e0ec48535b 2023-05-03 17:57:52 +02:00 Compare

@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 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.
First-time contributor

@BrianSavery
@brecht Indeed. Just tried the latest build & I'm getting exactly the same crashes & messed rendering unfortunately.

@BrianSavery @brecht Indeed. Just tried the latest build & I'm getting exactly the same crashes & messed rendering unfortunately.
First-time contributor

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)

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) `
First-time contributor

Neither offline nor windowed rendering renders results

Neither offline nor windowed rendering renders results
Brecht Van Lommel added 1 commit 2023-05-04 18:49:31 +02:00
buildbot/vexp-code-patch-coordinator Build done. Details
922d5a1a41
Try HIPRT SDK 2.0.3a134c7

There's a new build now at the same link, but I have not tested it myself yet.

There's a new build now at the same link, but I have not tested it myself yet.
First-time contributor

It now appears to render correctly, though it performs worse than HIP in the two renders I tried on my 6800XT.

It now appears to render correctly, though it performs worse than HIP in the two renders I tried on my 6800XT.
First-time contributor

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.

> 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.
First-time contributor

@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.

@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.
First-time contributor

@salipour @brecht Lighting/shading doesn't match in the Laundromat scene (see screenshots below)..

@salipour @brecht Lighting/shading doesn't match in the Laundromat scene (see screenshots below)..
Author
Member

@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.

> @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.
First-time contributor

@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.

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,)

> @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. 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,)
First-time contributor

@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/

@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/
Author
Member

@salipour some scenes crash when building the BHV (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?

> @salipour some scenes crash when building the BHV (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?
First-time contributor

@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

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).

> @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 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).
First-time contributor

@salipour some scenes crash when building the BHV (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?

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 some scenes crash when building the BHV (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? 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.
First-time contributor

@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

@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
Member

@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 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.

@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.
Author
Member

@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.

> @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:

  • List the remaining issues and todo's in #104110
  • Provide text for the release notes, including any mention of any issues that users should know about?
Can you: * List the remaining issues and todo's in #104110 * Provide text for the release notes, including any mention of any issues that users should know about?
Brecht Van Lommel added 1 commit 2023-05-15 12:12:40 +02:00
buildbot/vexp-code-patch-coordinator Build done. Details
4292347960
Merge branch 'main' into hiprt
Brecht Van Lommel merged commit b196109dac into main 2023-05-15 13:46:13 +02:00
Brecht Van Lommel deleted branch hip-raytracing 2023-05-15 13:46:14 +02:00

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

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
First-time contributor

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?

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.

We are looking to enable it for Blender 4.0 in Linux so users don't have to install the HIP RT SDK themselves.
Brecht Van Lommel removed this from the Render & Cycles project 2023-05-24 19:05:20 +02:00
Howard Trickey referenced this issue from a commit 2023-05-29 02:51:40 +02:00
First-time contributor

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.

> 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.

It's not meant to work on Linux yet, there are build instructions that can be given at this moment.
Sign in to join this conversation.
No reviewers
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
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
Overlay
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
Interest
Workbench
Interest: X11
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
No Milestone
No project
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
No description provided.