1  /* Test the atomic exchange expansion, execution.  */
       2  
       3  /* { dg-do run } */
       4  /* { dg-options "-Wno-long-long" } */
       5  
       6  /* We're trying to generate this type of store/exchange/load sequence:
       7       st.global.u32   [g32], %r60;
       8       atom.global.exch.b32    %r22, [g32], 2;
       9       ld.global.u32   %r23, [g32];
      10     with no insns inbetween.
      11  
      12     We compile this at -O0, to keep the compiler from optimizing out the
      13     "p = (P)" assignment.  If the assignment is optimized out we don't test
      14     the generic case, iow we generate for instance atom.global.exch.b32 instead
      15     of atom.exch.b32.
      16  
      17     Compiling at -O0 however does introduce loads and stores in the
      18     store/exchange/load sequence, so we fix that by using the register
      19     keyword.  */
      20  
      21  enum memmodel
      22  {
      23    MEMMODEL_RELAXED = 0,
      24  };
      25  
      26  unsigned int g32;
      27  unsigned long long int g64;
      28  
      29  unsigned int s32 __attribute__((shared));
      30  unsigned long long int s64 __attribute__((shared));
      31  
      32  #define TEST(P, V1, V2)						\
      33    {								\
      34      register typeof (*(P)) tmp;					\
      35      register typeof (*(P)) tmp2;				\
      36      __atomic_store_n ((P), (V1), MEMMODEL_RELAXED);		\
      37      tmp = __atomic_exchange_n ((P), (V2), MEMMODEL_RELAXED);	\
      38      tmp2 = __atomic_load_n ((P), MEMMODEL_RELAXED);		\
      39      if (tmp != (V1) || tmp2 != (V2))				\
      40        __builtin_abort ();					\
      41    }
      42  
      43  #define TEST2(P, V1, V2)					\
      44    {								\
      45      register typeof (*(P)) tmp;					\
      46      register typeof (*(P)) tmp2;				\
      47      *(P) = (V1);						\
      48      tmp = __atomic_exchange_n ((P), (V2), MEMMODEL_RELAXED);	\
      49      tmp2 = *(P);						\
      50      if (tmp != (V1) || tmp2 != (V2))				\
      51        __builtin_abort ();					\
      52    }
      53  
      54  #define TESTS(P)				\
      55    {						\
      56      TEST ((P), 1, 2);				\
      57      TEST2 ((P), 3, 4);				\
      58      {						\
      59        register typeof (*(P)) * p = (P);		\
      60        TEST (p, 1, 2);				\
      61        TEST2 (p, 3, 4);				\
      62      }						\
      63    }
      64  
      65  int
      66  main ()
      67  {
      68    TESTS (&g32);
      69    TESTS (&g64);
      70    TESTS (&s32);
      71    TESTS (&s64);
      72  
      73    return 0;
      74  }