99   __attribute__((shared)) __UINT32_TYPE__ val;  in test_shared32()  local
104   val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_shared32()
109   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_shared32()
114   __attribute__((shared)) __UINT64_TYPE__ val;  in test_shared64()  local
119   val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_shared64()
124   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_shared64()
189   __attribute__((shared)) __UINT32_TYPE__ val;  in test_order32()  local
192   val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_ACQUIRE, "workgroup");  in test_order32()
195   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_RELEASE, "workgroup");  in test_order32()
198   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_ACQ_REL, "workgroup");  in test_order32()
201   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_order32()
206   __attribute__((shared)) __UINT64_TYPE__ val;  in test_order64()  local
209   val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_ACQUIRE, "workgroup");  in test_order64()
212   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_RELEASE, "workgroup");  in test_order64()
215   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_ACQ_REL, "workgroup");  in test_order64()
218   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_order64()
223   __attribute__((shared)) __UINT32_TYPE__ val;  in test_scope32()  local
226   val = __builtin_amdgcn_atomic_inc32(&val, val, __ATOMIC_SEQ_CST, "");  in test_scope32()
229   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_scope32()
232   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "agent");  in test_scope32()
235   val = __builtin_amdgcn_atomic_dec32(&val, val, __ATOMIC_SEQ_CST, "wavefront");  in test_scope32()
240   __attribute__((shared)) __UINT64_TYPE__ val;  in test_scope64()  local
243   val = __builtin_amdgcn_atomic_inc64(&val, val, __ATOMIC_SEQ_CST, "");  in test_scope64()
246   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "workgroup");  in test_scope64()
249   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "agent");  in test_scope64()
252   val = __builtin_amdgcn_atomic_dec64(&val, val, __ATOMIC_SEQ_CST, "wavefront");  in test_scope64()