43 lines · plain
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature2 // REQUIRES: amdgpu-registered-target3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s4 5#define __shared__ __attribute__((shared))6 7__shared__ __amdgpu_named_workgroup_barrier_t bar;8__shared__ __amdgpu_named_workgroup_barrier_t arr[2];9__shared__ struct {10 __amdgpu_named_workgroup_barrier_t x;11 __amdgpu_named_workgroup_barrier_t y;12} str;13 14__amdgpu_named_workgroup_barrier_t *getBar();15void useBar(__amdgpu_named_workgroup_barrier_t *);16 17// CHECK-LABEL: define {{[^@]+}}@_Z7testSemPu34__amdgpu_named_workgroup_barrier_t18// CHECK-SAME: (ptr noundef [[P:%.*]]) #[[ATTR0:[0-9]+]] {19// CHECK-NEXT: entry:20// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr, align 8, addrspace(5)21// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)22// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr23// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[P_ADDR]] to ptr24// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 825// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 826// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[TMP0]]) #[[ATTR2:[0-9]+]]27// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef addrspacecast (ptr addrspace(1) @bar to ptr)) #[[ATTR2]]28// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds ([2 x target("amdgcn.named.barrier", 0)], ptr addrspacecast (ptr addrspace(1) @arr to ptr), i64 0, i64 1)) #[[ATTR2]]29// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef getelementptr inbounds nuw ([[STRUCT_ANON:%.*]], ptr addrspacecast (ptr addrspace(1) @str to ptr), i32 0, i32 1)) #[[ATTR2]]30// CHECK-NEXT: [[CALL:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]]31// CHECK-NEXT: call void @_Z6useBarPu34__amdgpu_named_workgroup_barrier_t(ptr noundef [[CALL]]) #[[ATTR2]]32// CHECK-NEXT: [[CALL1:%.*]] = call noundef ptr @_Z6getBarv() #[[ATTR2]]33// CHECK-NEXT: ret ptr [[CALL1]]34//35__amdgpu_named_workgroup_barrier_t *testSem(__amdgpu_named_workgroup_barrier_t *p) {36 useBar(p);37 useBar(&bar);38 useBar(&arr[1]);39 useBar(&str.y);40 useBar(getBar());41 return getBar();42}43