17 lines · plain
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-CAS4 5// REQUIRES: amdgpu-registered-target6 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 scope11// GFX90A-CAS-LABEL: _Z14atomic_add_casPf12// GFX90A-CAS: flat_atomic_cmpswap13// GFX90A-CAS: s_cbranch_execnz14__device__ float atomic_add_cas(float *p) {15 return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);16}17