composable_kernel icon indicating copy to clipboard operation
composable_kernel copied to clipboard

[Issue]: Can't compile for debugging with -O0

Open cameronshinn opened this issue 3 months ago • 1 comments

Problem Description

I currently am unable to compile for debugging with -O0. When I add that flag, I get an error "error: Illegal instruction detected: Operand has incorrect register class.". The errors span various lines within include/ck/utility/amd_buffer_addressing.hpp.

I found a comment about this in another issue here. The fix is to disable CK_USE_AMD_BUFFER_LOAD and CK_USE_AMD_BUFFER_STORE in ck.hpp. That allows me to successfully compile, but when I run my program I get segfaults.

Operating System

Ubuntu 22.04.4 LTS (Jammy Jellyfish)

CPU

AMD EPYC 7402 24-Core Processor

GPU

AMD Instinct MI210

Other

No response

ROCm Version

ROCm 6.0.0

ROCm Component

Composable Kernel

Steps to Reproduce

Commands:

mkdir build && cd build
cmake -D CMAKE_PREFIX_PATH=/opt/rocm -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc -D CMAKE_BUILD_TYPE=Debug -D GPU_TARGETS=gfx90a -D CMAKE_CXX_FLAGS_DEBUG="-g -O0" ..  # Specifically the CMAKE_CXX_FLAGS_DEBUG
make -j example_gemm_xdl

Output: compile_example_gemm_xdl.txt

(Optional for Linux users) Output of /opt/rocm/bin/rocminfo --support

No response

Additional Information

No response

cameronshinn avatar Mar 17 '24 20:03 cameronshinn