hipamd
hipamd copied to clipboard
Compile error in hip_cooperative_groups.h with ROCm 5.5
With ROCm 5.5, the inclusion of hip/hip_cooperative_groups.h
results in a compile error on debug builds:
In file included from /opt/rocm-5.5.0/include/hip/hip_cooperative_groups.h:38:
/opt/rocm-5.5.0/include/hip/amd_detail/amd_hip_cooperative_groups.h:455:7: error: invalid instruction, did you mean: s_trap?
__hip_assert(false && "invalid cooperative group type");
^
/opt/rocm-5.5.0/include/hip/amd_detail/amd_hip_cooperative_groups.h:48:7: note: expanded from macro '__hip_assert'
__hip_abort(); \
^
/opt/rocm-5.5.0/include/hip/amd_detail/amd_hip_cooperative_groups.h:41:9: note: expanded from macro '__hip_abort'
{ asm("trap;"); }
^
<inline asm>:1:2: note: instantiated into assembly here
trap;
^
This issue does not occur in ROCm 5.4, where this code used assert
instead. The problem seems to have been introduced by commit https://github.com/ROCm-Developer-Tools/hipamd/commit/5ff4b16db0cdebd250fc637eb209927dec04b4cb. Note that it only occurs if NDEBUG
is not defined. Steps to reproduce:
git clone https://github.com/amd/rocm-examples.git
cd rocm-examples/HIP-Basic/cooperative_groups
cmake -S . -B build
cmake --build build
can you true by just doing a make
inside rocm-examples/HIP-Basic/cooperative_groups