clang/test/CodeGen/amdgpu-abi-version.c

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals --version 3
   2 // RUN: %clang_cc1 -cc1 -triple amdgcn-amd-amdhsa -emit-llvm -mcode-object-version=none %s -o - | FileCheck %s
   3
   4 //.
   5 // CHECK: @__oclc_ABI_version = external addrspace(4) global i32
   6 //.
   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 ptr
  12 // CHECK-NEXT:    [[TMP0:%.*]] = load i32, ptr addrspace(4) @__oclc_ABI_version, align 4
  13 // CHECK-NEXT:    [[TMP1:%.*]] = icmp sge i32 [[TMP0]], 500
  14 // 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 12
  16 // 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 4
  18 // 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 !3, !noundef !3
  20 // CHECK-NEXT:    [[CONV:%.*]] = zext i16 [[TMP7]] to i32
  21 // CHECK-NEXT:    ret i32 [[CONV]]
  22 //
  23 int foo() { return __builtin_amdgcn_workgroup_size_x(); }