llvm/test/CodeGen/AMDGPU/lower-ctor-dtor.ll

   1 ; RUN: opt -S -mtriple=amdgcn-- -amdgpu-lower-ctor-dtor < %s | FileCheck %s
   2 ; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-ctor-dtor < %s | FileCheck %s
   3
   4 ; Make sure we get the same result if we run multiple times
   5 ; RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-ctor-dtor,amdgpu-lower-ctor-dtor < %s | FileCheck %s
   6 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -s - 2>&1 | FileCheck %s -check-prefix=VISIBILITY
   7 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -S - 2>&1 | FileCheck %s -check-prefix=SECTION
   8 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -amdgpu-lower-global-ctor-dtor=0 -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf -s - 2>&1 | FileCheck %s -check-prefix=DISABLED
   9 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf --notes - 2>&1 | FileCheck %s -check-prefix=METADATA
  10
  11 @llvm.global_ctors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @foo, ptr null }]
  12 @llvm.global_dtors = appending addrspace(1) global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 1, ptr @bar, ptr null }]
  13
  14 ; CHECK: @__init_array_start = external addrspace(1) constant [0 x ptr addrspace(1)]
  15 ; CHECK: @__init_array_end = external addrspace(1) constant [0 x ptr addrspace(1)]
  16 ; CHECK: @__fini_array_start = external addrspace(1) constant [0 x ptr addrspace(1)]
  17 ; CHECK: @__fini_array_end = external addrspace(1) constant [0 x ptr addrspace(1)]
  18 ; CHECK: @llvm.used = appending global [2 x ptr] [ptr @amdgcn.device.init, ptr @amdgcn.device.fini]
  19
  20 ; CHECK-LABEL: define weak_odr amdgpu_kernel void @amdgcn.device.init() #0
  21 ; CHECK-NEXT:  entry:
  22 ; CHECK-NEXT:    br i1 icmp ne (ptr addrspace(1) @__init_array_start, ptr addrspace(1) @__init_array_end), label [[WHILE_ENTRY:%.*]], label [[WHILE_END:%.*]]
  23 ; CHECK:       while.entry:
  24 ; CHECK-NEXT:    [[PTR:%.*]] = phi ptr addrspace(1) [ @__init_array_start, [[ENTRY:%.*]] ], [ [[NEXT:%.*]], [[WHILE_ENTRY]] ]
  25 ; CHECK-NEXT:    [[CALLBACK:%.*]] = load ptr, ptr addrspace(1) [[PTR]], align 8
  26 ; CHECK-NEXT:    call void [[CALLBACK]]()
  27 ; CHECK-NEXT:    [[NEXT]] = getelementptr ptr addrspace(1), ptr addrspace(1) [[PTR]], i64 1
  28 ; CHECK-NEXT:    [[END:%.*]] = icmp eq ptr addrspace(1) [[NEXT]], @__init_array_end
  29 ; CHECK-NEXT:    br i1 [[END]], label [[WHILE_END]], label [[WHILE_ENTRY]]
  30 ; CHECK:       while.end:
  31 ; CHECK-NEXT:    ret void
  32
  33 ; CHECK-LABEL: define weak_odr amdgpu_kernel void @amdgcn.device.fini() #1
  34 ; CHECK-NEXT:  entry:
  35 ; CHECK-NEXT:    br i1 icmp ne (ptr addrspace(1) @__fini_array_start, ptr addrspace(1) @__fini_array_end), label [[WHILE_ENTRY:%.*]], label [[WHILE_END:%.*]]
  36 ; CHECK:       while.entry:
  37 ; CHECK-NEXT:    [[PTR:%.*]] = phi ptr addrspace(1) [ @__fini_array_start, [[ENTRY:%.*]] ], [ [[NEXT:%.*]], [[WHILE_ENTRY]] ]
  38 ; CHECK-NEXT:    [[CALLBACK:%.*]] = load ptr, ptr addrspace(1) [[PTR]], align 8
  39 ; CHECK-NEXT:    call void [[CALLBACK]]()
  40 ; CHECK-NEXT:    [[NEXT]] = getelementptr ptr addrspace(1), ptr addrspace(1) [[PTR]], i64 1
  41 ; CHECK-NEXT:    [[END:%.*]] = icmp eq ptr addrspace(1) [[NEXT]], @__fini_array_end
  42 ; CHECK-NEXT:    br i1 [[END]], label [[WHILE_END]], label [[WHILE_ENTRY]]
  43 ; CHECK:       while.end:
  44 ; CHECK-NEXT:    ret void
  45
  46 ; CHECK-NOT: amdgcn.device.
  47
  48 ; VISIBILITY: FUNC   WEAK PROTECTED {{.*}} amdgcn.device.init
  49 ; VISIBILITY: OBJECT WEAK DEFAULT {{.*}} amdgcn.device.init.kd
  50 ; VISIBILITY: FUNC   WEAK PROTECTED {{.*}} amdgcn.device.fini
  51 ; VISIBILITY: OBJECT   WEAK DEFAULT {{.*}} amdgcn.device.fini.kd
  52 ; SECTION: .init_array.1     INIT_ARRAY      {{.*}} {{.*}} 000008 00  WA  0   0  8
  53 ; SECTION: .fini_array.1     FINI_ARRAY      {{.*}} {{.*}} 000008 00  WA  0   0  8
  54 ; DISABLED-NOT: FUNC   GLOBAL PROTECTED {{.*}} amdgcn.device.init
  55 ; DISABLED-NOT: OBJECT GLOBAL DEFAULT {{.*}} amdgcn.device.init.kd
  56 ; DISABLED-NOT: FUNC   GLOBAL PROTECTED {{.*}} amdgcn.device.fini
  57 ; DISABLED-NOT: OBJECT   GLOBAL DEFAULT {{.*}} amdgcn.device.fini.kd
  58 ; METADATA:  amdhsa.kernels:
  59 ; METADATA:    .kind:           init
  60 ; METADATA:    .max_flat_workgroup_size: 1
  61 ; METADATA:    .name:           amdgcn.device.init
  62 ; METADATA:    .symbol:         amdgcn.device.init.kd
  63 ; METADATA:    .kind:           fini
  64 ; METADATA:    .max_flat_workgroup_size: 1
  65 ; METADATA:    .name:           amdgcn.device.fini
  66 ; METADATA:    .symbol:         amdgcn.device.fini.kd
  67
  68 define internal void @foo() {
  69   ret void
  70 }
  71
  72 define internal void @bar() {
  73   ret void
  74 }
  75
  76 ; CHECK: attributes #0 = { "amdgpu-flat-work-group-size"="1,1" "device-init" }
  77 ; CHECK: attributes #1 = { "amdgpu-flat-work-group-size"="1,1" "device-fini" }