1 // RUN: %clang_cc1 %s -triple=amdgcn-amd-amdhsa -fcuda-is-device \
2 // RUN: -target-cpu gfx90a -Rpass=atomic-expand -S -o - 2>&1 | \
3 // RUN: FileCheck %s --check-prefix=GFX90A-CAS
4
5 // REQUIRES: amdgpu-registered-target
6
7 #include "Inputs/cuda.h"
8 #include <stdatomic.h>
9
10 // GFX90A-CAS: A compare and swap loop was generated for an atomic fadd operation at system memory scope
11 // GFX90A-CAS-LABEL: _Z14atomic_add_casPf
12 // GFX90A-CAS: flat_atomic_cmpswap
13 // GFX90A-CAS: s_cbranch_execnz
atomic_add_cas(float * p)14 __device__ float atomic_add_cas(float *p) {
15 return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
16 }
17