[Bug target/98321] New: [nvptx] 'atom.add.f32' for atomic add of 32-bit 'float'

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

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

            Bug ID: 98321
           Summary: [nvptx] 'atom.add.f32' for atomic add of 32-bit
                    'float'
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tschwinge at gcc dot gnu.org
                CC: vries at gcc dot gnu.org
  Target Milestone: ---
            Target: nvptx

Consider:

    TYPE f(TYPE a, TYPE b)
    {
      #pragma acc atomic update
      a += b;

      return a;
    }

Compiling always with '-fopenacc', for '-DTYPE=int'/'-DTYPE=long' I do see the
expected 'atom.add.u32'/'atom.add.u64', but for '-DTYPE=float' I do not see the
expected 'atom.add.f32' but instead an 'atom.cas.b32' loop.  (I understand that
'-DTYPE=double': 'atom.add.f64' depends on PTX 5.0, SM 6.0 support.)