clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu

   1 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \
   2 // RUN:  -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \
   3 // RUN:  -o - | FileCheck %s
   4
   5 #define __device__ __attribute__((device))
   6 typedef __attribute__((address_space(3))) float *LP;
   7
   8 // CHECK-LABEL: test_ds_atomic_add_f32
   9 // CHECK: %[[ADDR_ADDR:.*]] = alloca float*, align 8, addrspace(5)
  10 // CHECK: %[[ADDR_ADDR_ASCAST_PTR:.*]] = addrspacecast float* addrspace(5)* %[[ADDR_ADDR]] to float**
  11 // CHECK: store float* %addr, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
  12 // CHECK: %[[ADDR_ADDR_ASCAST:.*]] = load float*, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
  13 // CHECK: %[[AS_CAST:.*]] = addrspacecast float* %[[ADDR_ADDR_ASCAST]] to float addrspace(3)*
  14 // CHECK: %3 = call contract float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %[[AS_CAST]]
  15 // CHECK: %4 = load float*, float** %rtn.ascast, align 8
  16 // CHECK: store float %3, float* %4, align 4
  17 __device__ void test_ds_atomic_add_f32(float *addr, float val) {
  18   float *rtn;
  19   *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
  20 }