RenderKit / oidn

Intel® Open Image Denoise library
https://www.openimagedenoise.org/
Apache License 2.0
1.73k stars 160 forks source link

HIP compilation error on RX6800XT (RDNA2) on windows10 #193

Closed feniksa closed 5 months ago

feniksa commented 5 months ago

Hi guys,

By default, master branch doesn't compiles on RX6800XT video card with HIP 5.7.1 , Windows10.

Compilation error is:

In file included from D:/workspace/oidn/devices/hip/../../external/composable_kernel/include\ck/utility/common_header.hpp:46:
D:/workspace/oidn/devices/hip/../../external/composable_kernel/include\ck/utility/amd_inline_asm.hpp:212:21: error: instruction not supported on this GPU
    asm volatile("\n \
                    ^
<inline asm>:2:14: note: instantiated into assembly here
             v_dot4_i32_i8 v5, v6, v7, v5
atafra commented 5 months ago

Was HIP 5.7.1 just released? The HIP SDK download page seems to be down, and last week the most recent version was HIP 5.7.0, which compiles without issues.

feniksa commented 5 months ago

@atafra Due to page: https://www.amd.com/de/developer/resources/rocm-hub/hip-sdk.html

image

After installation, hipsdk version in hip/hip_version.h:

#define HIP_VERSION_MAJOR 5
#define HIP_VERSION_MINOR 7
#define HIP_VERSION_PATCH 32000
#define HIP_VERSION_GITHASH "193a0b56e"
#define HIP_VERSION_BUILD_ID 0
#define HIP_VERSION_BUILD_NAME ""
atafra commented 5 months ago

None of the download links on this page works for me (404).

feniksa commented 5 months ago

Same for me now. Will ping AMD guys.

When issue will be resolved, I will notify in this ticket. (I was lucky to download HIP SDK from working link).

feniksa commented 5 months ago

@atafra HIP download links fixed

atafra commented 5 months ago

Thanks. This is actually identical to what I downloaded last week. It's installed in the "5.7" subdirectory, which confused me.

I can reproduce this issue but only in Debug mode. Building in Release mode works without issues. We can't really do anything about this issue. This seems to be a compiler bug. The code for which the compiler fails is an AMD library (Composable Kernels). Perhaps you could try to report this issue to AMD.

feniksa commented 5 months ago

Hi @atafra

Thanks. I will report this issue to AMD, and close this issue. When AMD (or me? hehe) will fix issue in composable kernels repo, I will create PR for submodule update.

Thanks

feniksa commented 5 months ago

Close issue as not issue in oidn

feniksa commented 5 months ago

Linked issue: https://github.com/ROCm/composable_kernel/issues/1146