1 // RUN: %clang_cc1 -no-opaque-pointers %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
2 // RUN: -fcuda-is-device -target-cpu gfx906 -fnative-half-type \
3 // RUN: -fnative-half-arguments-and-returns | FileCheck %s
4
5 // REQUIRES: amdgpu-registered-target
6
7 #include "Inputs/cuda.h"
8 #include <stdatomic.h>
9
ffp1(float * p)10 __device__ float ffp1(float *p) {
11 // CHECK-LABEL: @_Z4ffp1Pf
12 // CHECK: atomicrmw fadd float* {{.*}} monotonic
13 return __atomic_fetch_add(p, 1.0f, memory_order_relaxed);
14 }
15
ffp2(double * p)16 __device__ double ffp2(double *p) {
17 // CHECK-LABEL: @_Z4ffp2Pd
18 // CHECK: atomicrmw fsub double* {{.*}} monotonic
19 return __atomic_fetch_sub(p, 1.0, memory_order_relaxed);
20 }
21
22 // long double is the same as double for amdgcn.
ffp3(long double * p)23 __device__ long double ffp3(long double *p) {
24 // CHECK-LABEL: @_Z4ffp3Pe
25 // CHECK: atomicrmw fsub double* {{.*}} monotonic
26 return __atomic_fetch_sub(p, 1.0L, memory_order_relaxed);
27 }
28
ffp4(double * p,float f)29 __device__ double ffp4(double *p, float f) {
30 // CHECK-LABEL: @_Z4ffp4Pdf
31 // CHECK: fpext float {{.*}} to double
32 // CHECK: atomicrmw fsub double* {{.*}} monotonic
33 return __atomic_fetch_sub(p, f, memory_order_relaxed);
34 }
35
ffp5(double * p,int i)36 __device__ double ffp5(double *p, int i) {
37 // CHECK-LABEL: @_Z4ffp5Pdi
38 // CHECK: sitofp i32 {{.*}} to double
39 // CHECK: atomicrmw fsub double* {{.*}} monotonic
40 return __atomic_fetch_sub(p, i, memory_order_relaxed);
41 }
42