[Bug target/97444] [nvptx] stack atomics

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

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

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
Created attachment 52169
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=52169&action=edit
Tentative patch, __atomic_exchange only

Code generated for the generic case:
...
        { // Atomic exchange - generic - start                                  
          .reg .pred %local_p;                                                  
          isspacep.local %local_p, %r26;                                        
          @%local_p bra $L51;                                                   
          .reg .pred %shared_p;                                                 
          isspacep.shared %shared_p, %r26;                                      
          @%local_p bra $L52;                                                   
          // Atomic exchange - generic - global                                 
          atom.exch.b32 %r27, [%r26], 2;                                        
          bra $L53;                                                             
$L51:
          // Atomic exchange - generic - local                                  
          {
            .reg .b32 %tmp;                                                     
            ld.b32 %tmp, [%r26];                                                
            st.b32 [%r26], 2;                                                   
            mov.b32 %r27, %tmp;                                                 
          }
          bra $L53;                                                             
$L52:
          // Atomic exchange - generic - shared                                 
          atom.exch.b32 %r27, [%r26], 2;                                        
$L53:
        } // Atomic exchange - generic - end                                    
...