[Bug target/127542] New: [GCN] ICE: too few operands for instruction v_bcnt_u32_b32 v4, v4

[burnus at gcc dot gnu.org via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=127542

            Bug ID: 127542
           Summary: [GCN] ICE: too few operands for instruction
                    v_bcnt_u32_b32  v4, v4
           Product: gcc
           Version: 17.0
            Status: UNCONFIRMED
          Keywords: ice-on-valid-code
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: ams at gcc dot gnu.org
  Target Milestone: ---
            Target: gcn

gcc -fopenmp -foffload=amdgcn-amdhsa -O2 inp2.ii

fails as follows. It works with nvptx offloading or with -O0 / -O1, failing
with -O3 or higher.

The fail is with the LLVM-mc assembler (here: llvm 23 but 18 is identical):


.../libexec/gcc/x86_64-pc-linux-gnu/17.0.0/accel/amdgcn-amdhsa/as
-triple=amdgcn--amdhsa -mcpu=gfx90a --amdhsa-code-object-version=4
-filetype=obj -o
./a.xamdgcn-amdhsa.mkoffload.hsaco-a.xamdgcn-amdhsa.mkoffload.2.o
./a.xamdgcn-amdhsa.mkoffload.2.s

./a.xamdgcn-amdhsa.mkoffload.2.s:198:2: error: too few operands for instruction
        v_bcnt_u32_b32  v4, v4
        ^
./a.xamdgcn-amdhsa.mkoffload.2.s:199:2: error: too few operands for instruction
        v_bcnt_u32_b32  v3, v0


Reduced testcase from the C++ testcase
https://github.com/ORNL/HeCBench/tree/master/src/popcount-omp

-------------------- testcase.cc -------------
int atoi_length;
int *atoi_result;
void atoi() {
#pragma omp target
  for (int i; i < atoi_length; i++) {
    unsigned long x = i;
    x -= x >> 1 & 6148914691236517205;
    x = (x & 3689348814741910323) + (x >> 2 & 3689348814741910323);
    x = x + (x >> 4) & 1085102592571150095;
    atoi_result[i] = x * 72340172838076673 >> 56;
  }
}