This repository was archived by the owner on Jan 26, 2024. It is now read-only.

Description
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 5ff4b16. 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