33 lines · c
1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 52// RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck %s3 4//.5// CHECK: @__oclc_ABI_version = external addrspace(4) global i326//.7// CHECK-LABEL: define dso_local i32 @foo(8// CHECK-SAME: ) #[[ATTR0:[0-9]+]] {9// CHECK-NEXT: [[ENTRY:.*:]]10// CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4, addrspace(5)11// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[RETVAL]] to ptr12// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 413// CHECK-NEXT: [[TMP1:%.*]] = icmp sge i32 [[TMP0]], 50014// CHECK-NEXT: [[TMP2:%.*]] = call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()15// CHECK-NEXT: [[TMP3:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP2]], i32 1216// CHECK-NEXT: [[TMP4:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()17// CHECK-NEXT: [[TMP5:%.*]] = getelementptr i8, ptr addrspace(4) [[TMP4]], i32 418// CHECK-NEXT: [[TMP6:%.*]] = select i1 [[TMP1]], ptr addrspace(4) [[TMP3]], ptr addrspace(4) [[TMP5]]19// CHECK-NEXT: [[TMP7:%.*]] = load i16, ptr addrspace(4) [[TMP6]], align 2, !range [[RNG2:![0-9]+]], !invariant.load [[META3:![0-9]+]], !noundef [[META3]]20// CHECK-NEXT: [[CONV:%.*]] = zext i16 [[TMP7]] to i3221// CHECK-NEXT: ret i32 [[CONV]]22//23int foo() { return __builtin_amdgcn_workgroup_size_x(); }24//.25// CHECK: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" }26// CHECK: attributes #[[ATTR1:[0-9]+]] = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }27//.28// CHECK: [[META0:![0-9]+]] = !{i32 1, !"wchar_size", i32 4}29// CHECK: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"}30// CHECK: [[RNG2]] = !{i16 1, i16 1025}31// CHECK: [[META3]] = !{}32//.33