ROCmSoftwarePlatform
[Issue]: Compilation error in Visual Studio debug mode , Win10
Problem Description
Compilation error in debug on Visual Studio compiler
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
Same time everything is fine in release mode
Operating System
Windows 10 22H2
CPU
AMD Epyc
GPU
AMD Radeon RX 6800 XT
Other
No response
ROCm Version
ROCm 5.7.1
ROCm Component
HIP
Steps to Reproduce
- Fetch OIDN repository
git clone --recurse-submodules [email protected]:OpenImageDenoise/oidn.git - Configure cmake project with enabled HIP backend:
cmake -DOIDN_DEVICE_HIP=1 .. - Run build in debug mode
P.S. If you have CLion IDE, it is much easier to configure project. Open project after fetch, open Settings -> Build, Execution, Deployment -> CMake, Select Build Type "Debug", Toolchain Visual Studio, Generator "Ninja", in "CMake Options" insert -DOIDN_DEVICE_HIP=1
(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support
No response
Additional Information
No response
Unfortunately, still observing errors when trying to build in debug mode with ROCm6.1 compiler.
@feniksa Apologies for the lack of response. Internal ticket has been created to investigate this issue. Thanks!
Hi @feniksa , there is another similar issue reported in the past: https://github.com/ROCm/composable_kernel/issues/508#issuecomment-1334549872 Please try compiling with -O1 instead of -O0.
The root cause of your issue is probably similar to that of the linked ticket, though I am not sure exactly what your build toolchain is doing. Note that your GPU (AMD Radeon RX 6800 XT = GFX1030) does theoretically support instruction v_dot4_i32_i8 as noted in the LLVM GFX1030 ISA reference. However, you still may encounter that instruction not supported on this GPU error if you're compiling for older GPU architectures.
So, another idea may be to target only GFX1030+ architectures - I'm not sure what compiler you're using, if it's LLVM Clang you can try setting the offload-arch parameter to build for just your GPU architecture. However, it's probably simpler to build with -O1.
