32 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 42// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip \3// RUN: -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \4// RUN: -o - | FileCheck %s5 6#define __device__ __attribute__((device))7typedef __attribute__((address_space(3))) float *LP;8 9// CHECK-LABEL: define spir_func void @_Z22test_ds_atomic_add_f32Pff(10// CHECK-SAME: ptr addrspace(4) noundef [[ADDR:%.*]], float noundef [[VAL:%.*]]) addrspace(4) #[[ATTR0:[0-9]+]] {11// CHECK-NEXT: entry:12// CHECK-NEXT: [[ADDR_ADDR:%.*]] = alloca ptr addrspace(4), align 813// CHECK-NEXT: [[VAL_ADDR:%.*]] = alloca float, align 414// CHECK-NEXT: [[RTN:%.*]] = alloca ptr addrspace(4), align 815// CHECK-NEXT: [[ADDR_ADDR_ASCAST:%.*]] = addrspacecast ptr [[ADDR_ADDR]] to ptr addrspace(4)16// CHECK-NEXT: [[VAL_ADDR_ASCAST:%.*]] = addrspacecast ptr [[VAL_ADDR]] to ptr addrspace(4)17// CHECK-NEXT: [[RTN_ASCAST:%.*]] = addrspacecast ptr [[RTN]] to ptr addrspace(4)18// CHECK-NEXT: store ptr addrspace(4) [[ADDR]], ptr addrspace(4) [[ADDR_ADDR_ASCAST]], align 819// CHECK-NEXT: store float [[VAL]], ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 420// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[ADDR_ADDR_ASCAST]], align 821// CHECK-NEXT: [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr addrspace(3)22// CHECK-NEXT: [[TMP2:%.*]] = load float, ptr addrspace(4) [[VAL_ADDR_ASCAST]], align 423// CHECK-NEXT: [[TMP3:%.*]] = atomicrmw fadd ptr addrspace(3) [[TMP1]], float [[TMP2]] monotonic, align 424// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[RTN_ASCAST]], align 825// CHECK-NEXT: store float [[TMP3]], ptr addrspace(4) [[TMP4]], align 426// CHECK-NEXT: ret void27//28__device__ void test_ds_atomic_add_f32(float *addr, float val) {29 float *rtn;30 *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);31}32