[Bug target/96964] New: [nvptx] Implement __atomic_test_and_set

[vries at gcc dot gnu.org]

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

            Bug ID: 96964
           Summary: [nvptx] Implement __atomic_test_and_set
           Product: gcc
           Version: 11.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

Currently __atomic_test_and_set for nvptx falls back onto the "Failing all
else, assume a single threaded environment and simply perform the operation"
case in expand_atomic_test_and_set, so it doesn't map onto an actual atomic
operation.

So, for test-case test.c:
...
int a;

int
main (void)
{
  int res = __atomic_test_and_set (&a, __ATOMIC_SEQ_CST);
  return res;
}
...
we get:
...
$ gcc test.c -S -o-
// BEGIN PREAMBLE
        .version        3.1
        .target sm_30
        .address_size 64
// END PREAMBLE


// BEGIN GLOBAL FUNCTION DECL: main
.visible .func (.param.u32 %value_out) main (.param.u32 %in_ar0, .param.u64
%in_ar1);

// BEGIN GLOBAL FUNCTION DEF: main
.visible .func (.param.u32 %value_out) main (.param.u32 %in_ar0, .param.u64
%in_ar1)
{
        .reg.u32 %value;
        .local .align 16 .b8 %frame_ar[16];
        .reg.u64 %frame;
        cvta.local.u64 %frame, %frame_ar;
        .reg.u32 %r22;
        .reg.u32 %r23;
        .reg.u32 %r24;
        .reg.u32 %r25;
        .reg.u32 %r26;
                ld.global.u8    %r25, [a];
                mov.u32 %r26, 1;
                st.global.u8    [a], %r26;
                cvt.u32.u8      %r22, %r25;
                st.u32  [%frame], %r22;
                ld.u32  %r23, [%frame];
                mov.u32 %r24, %r23;
                mov.u32 %value, %r24;
        st.param.u32    [%value_out], %value;
        ret;
}


// BEGIN GLOBAL VAR DEF: a
        .visible .global .align 4 .u32 a[1];
...