Asking @brecht some hint: we are testing a compute based realtime render engine on HIP, but got stuck on the fact seems latest release of HIP allow to read but not write textures. Also all the writes to ROCm buffer are not visible from Vulkan or D3D12.
Did you solved this problem with AMD on the closed source driver, or used some other way around.
1 Like
If you take a look at recent commits in the HIP repos, you can see AMD still working on interop and as far as I know they are clearly not supporting it at the moment.
I believe/hope that it may change next ROCm release.
For example: https://github.com/ROCm-Developer-Tools/HIP/pull/2454
2 Likes
But there is no AMD comment on that coming to the 5.0 ROCm release?
I would be surprised. This is just my impression having put my nose quite a lot in the stack last weeks.
1 Like
This will be a breaker for us, so needing to float in CUDA direction.
We was almost there being able to use HIP (even with all the involved risks).
I understood Cycles X didn’t need this feature after all, even if I’m a bit surprised.
Not a dev anyway so sure missing many things.
We stuck with no texture support (not possible to create a mipmapped array at this moment) and no interop. Interestingly, tests with interop for Vulkan and mipmapped array tests are present inside HIP.git. So we are wondering if there is any magic driver version that supports these features.
I don’t know the stack deep enough and Blender even less, so hopefully Brecht can help you more.
There is no real magic in the drivers, the stack is almost completely open on github and the packages available for download. Only one or two obscure packages are proprietary.
Building the stack was a few years ago a nightmare, now it is almost acceptable - still very rough.
Maybe this can help you, on top of the official drivers:
- this script from an AMD engineer builds almost all the stack with just cmake.
- I adapted his script to build on Debian with vanilla LLVM - AMD have forked LLVM
- I also made a builder for Debian packages here.
From there enabling the tests and tinkering with them is pretty easy.
There are also Gentoo and Arch packages.
Yep, we are using these packages for Linux.
We have finished our HIP backend. Everything is working great, including GLSL->HIP translation. It’s working under Linux and Windows (with the latest driver). But the absence of Vulkan/D3D12 interop for buffers and textures makes HIP useless at this moment for us. That is sad.
2 Likes
Interestingly, tests with interop for Vulkan and mipmapped array tests are present inside HIP
I executed all HIP tests 3 days ago and had no issue with this one…
…until I realized that the test is skipped on Linux. ROC Common Language runtime routes to rocr-runtime on linux, and PAL on windows.
I don’t understand why it is disabled on Linux though, since there seems to be no compute/graphics interop in there, the kernel output is memcopied back from the device (GPU) to the host (CPU), instead of staying on the GPU until it reaches the scanout line.
Just to understand, on the contrary the test was correctly executed on Windows? I’m really confused since from system requirements ROCm is not even supported to work on Windows.
ROCm does work on windows as far as I know, although working on a different runtime as explained above.
Blender uses ROCm HIP for Cycles kernels on windows in its current 3.0 release and it works.
P.S. ah indeed the whole ROCm stack is not supported on windows. I suspect Blender/ROCm on windows of translating to openCL calls…?
By the way, I got to wonder: can one write their own compute/graphics interop? Or is there lack of hardware capabilities exposed by the firmware?
I was wondering exactly the same.
Would like @brecht can give us some hint about how they sorted out working with AMD and if adding this feature will be a feasible path forward.
HIP.git/samples/2_Cookbook/20_hip_vulkan test is crashing on Linux at hipImportExternalMemory().
There is no crash on Windows under D3D12/VK, but HIP buffer updates are not visible for graphics (only for HIP). After memory import, the virtual buffer address is the same at Vulkan and HIP.
Both hipMipmappedArrayCreate() and hipMallocMipmappedArray() are returning errors at any trivial parameter configurations.
The HIP is working fine with the Adrenalin 21.12.1 driver on Windows (not all API functions are present). But texture and interop problems are the same with the Linux ROCm 4.5.2 runtime. Basic tests with radix/bitonic sort are passing. The only issue is faulty warp function on new Radeons where warp size is 32 (there are no problems on old Vega). We will investigate it more.
1 Like
I’m really not the person to provide support for others using HIP in their applications. All I know is graphics interop is not working on Windows and Linux for Cycles and currently disabled, but AMD was working to fix it for a future driver version.
Moved this to “Other development topics” subforum. In general, keep in mind though, that this forum (devtalk) is only for Blender development.
@brecht @ThomasDinges thanks, my intent was to gather info about if and how this limitation was managed in Cycles X, the only GUI based app I was able to find using HIP.
1 Like