hipamd icon indicating copy to clipboard operation
hipamd copied to clipboard

Compile error in hip_cooperative_groups.h with ROCm 5.5

Open nolmoonen opened this issue 1 year ago • 1 comments

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

nolmoonen avatar May 12 '23 15:05 nolmoonen

can you true by just doing a make inside rocm-examples/HIP-Basic/cooperative_groups

jatinx avatar May 17 '23 16:05 jatinx