Home | History | Annotate | Download | only in OpenMP 
      1 // Test host codegen.
      2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s
      3 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
      4 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
      5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s
      6 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
      7 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
      8 
      9 // Test target codegen - host bc file has to be created first.
     10 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
     11 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefix=TCHECK
     12 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
     13 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
     14 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
     15 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s -check-prefix=TCHECK
     16 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
     17 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefix=TCHECK
     18 
     19 // Check that no target code is emmitted if no omptests flag was provided.
     20 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s -check-prefix=CHECK-NTARGET
     21 
     22 // expected-no-diagnostics
     23 #ifndef HEADER
     24 #define HEADER
     25 
     26 // CHECK-DAG: [[SA:%.+]] = type { [4 x i32] }
     27 // CHECK-DAG: [[SB:%.+]] = type { [8 x i32] }
     28 // CHECK-DAG: [[SC:%.+]] = type { [16 x i32] }
     29 // CHECK-DAG: [[SD:%.+]] = type { [32 x i32] }
     30 // CHECK-DAG: [[SE:%.+]] = type { [64 x i32] }
     31 // CHECK-DAG: [[ST1:%.+]] = type { [228 x i32] }
     32 // CHECK-DAG: [[ST2:%.+]] = type { [1128 x i32] }
     33 // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] }
     34 // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
     35 // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
     36 
     37 // TCHECK:    [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]] }
     38 
     39 // CHECK-DAG: [[A1:@.+]] = internal global [[SA]]
     40 // CHECK-DAG: [[A2:@.+]] = global [[SA]]
     41 // CHECK-DAG: [[B1:@.+]] = global [[SB]]
     42 // CHECK-DAG: [[B2:@.+]] = global [[SB]]
     43 // CHECK-DAG: [[C1:@.+]] = internal global [[SC]]
     44 // CHECK-DAG: [[D1:@.+]] = global [[SD]]
     45 // CHECK-DAG: [[E1:@.+]] = global [[SE]]
     46 // CHECK-DAG: [[T1:@.+]] = global [[ST1]]
     47 // CHECK-DAG: [[T2:@.+]] = global [[ST2]]
     48 
     49 // CHECK-NTARGET-DAG: [[SA:%.+]] = type { [4 x i32] }
     50 // CHECK-NTARGET-DAG: [[SB:%.+]] = type { [8 x i32] }
     51 // CHECK-NTARGET-DAG: [[SC:%.+]] = type { [16 x i32] }
     52 // CHECK-NTARGET-DAG: [[SD:%.+]] = type { [32 x i32] }
     53 // CHECK-NTARGET-DAG: [[SE:%.+]] = type { [64 x i32] }
     54 // CHECK-NTARGET-DAG: [[ST1:%.+]] = type { [228 x i32] }
     55 // CHECK-NTARGET-DAG: [[ST2:%.+]] = type { [1128 x i32] }
     56 // CHECK-NTARGET-NOT: type { i8*,
     57 // CHECK-NTARGET-NOT: type { i32,
     58 
     59 // We have 7 target regions
     60 
     61 // CHECK-DAG: {{@.+}} = private constant i8 0
     62 // TCHECK-NOT: {{@.+}} = private constant i8 0
     63 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     64 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     65 // CHECK-DAG: {{@.+}} = private constant i8 0
     66 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     67 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     68 // CHECK-DAG: {{@.+}} = private constant i8 0
     69 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     70 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     71 // CHECK-DAG: {{@.+}} = private constant i8 0
     72 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     73 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     74 // CHECK-DAG: {{@.+}} = private constant i8 0
     75 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     76 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     77 // CHECK-DAG: {{@.+}} = private constant i8 0
     78 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     79 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     80 // CHECK-DAG: {{@.+}} = private constant i8 0
     81 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     82 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     83 // CHECK-DAG: {{@.+}} = private constant i8 0
     84 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     85 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     86 // CHECK-DAG: {{@.+}} = private constant i8 0
     87 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     88 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     89 // CHECK-DAG: {{@.+}} = private constant i8 0
     90 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     91 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     92 // CHECK-DAG: {{@.+}} = private constant i8 0
     93 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     94 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     95 // CHECK-DAG: {{@.+}} = private constant i8 0
     96 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 4]
     97 // CHECK-DAG: {{@.+}} = private unnamed_addr constant [1 x i32] [i32 288]
     98 
     99 // CHECK-NTARGET-NOT: private constant i8 0
    100 // CHECK-NTARGET-NOT: private unnamed_addr constant [1 x i
    101 
    102 // CHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
    103 // CHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    104 // CHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
    105 // CHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    106 // CHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
    107 // CHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    108 // CHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
    109 // CHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    110 // CHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
    111 // CHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    112 // CHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
    113 // CHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    114 // CHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
    115 // CHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    116 // CHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
    117 // CHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    118 // CHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
    119 // CHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    120 // CHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
    121 // CHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    122 // CHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
    123 // CHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    124 // CHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
    125 // CHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* @{{.*}}, i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    126 
    127 // TCHECK-DAG: [[NAMEPTR1:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME1:__omp_offloading_[0-9a-f]+_[0-9a-f]+__Z.+_l[0-9]+]]\00"
    128 // TCHECK-DAG: [[ENTRY1:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR1]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    129 // TCHECK-DAG: [[NAMEPTR2:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME2:.+]]\00"
    130 // TCHECK-DAG: [[ENTRY2:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR2]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    131 // TCHECK-DAG: [[NAMEPTR3:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME3:.+]]\00"
    132 // TCHECK-DAG: [[ENTRY3:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR3]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    133 // TCHECK-DAG: [[NAMEPTR4:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME4:.+]]\00"
    134 // TCHECK-DAG: [[ENTRY4:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR4]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    135 // TCHECK-DAG: [[NAMEPTR5:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME5:.+]]\00"
    136 // TCHECK-DAG: [[ENTRY5:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR5]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    137 // TCHECK-DAG: [[NAMEPTR6:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME6:.+]]\00"
    138 // TCHECK-DAG: [[ENTRY6:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR6]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    139 // TCHECK-DAG: [[NAMEPTR7:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME7:.+]]\00"
    140 // TCHECK-DAG: [[ENTRY7:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR7]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    141 // TCHECK-DAG: [[NAMEPTR8:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME8:.+]]\00"
    142 // TCHECK-DAG: [[ENTRY8:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR8]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    143 // TCHECK-DAG: [[NAMEPTR9:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME9:.+]]\00"
    144 // TCHECK-DAG: [[ENTRY9:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR9]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    145 // TCHECK-DAG: [[NAMEPTR10:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME10:.+]]\00"
    146 // TCHECK-DAG: [[ENTRY10:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR10]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    147 // TCHECK-DAG: [[NAMEPTR11:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME11:.+]]\00"
    148 // TCHECK-DAG: [[ENTRY11:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR11]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    149 // TCHECK-DAG: [[NAMEPTR12:@.+]] = internal unnamed_addr constant [{{.*}} x i8] c"[[NAME12:.+]]\00"
    150 // TCHECK-DAG: [[ENTRY12:@.+]] = constant [[ENTTY]] { i8* bitcast (void (i[[SZ]])* @{{.*}} to i8*), i8* getelementptr inbounds ([{{.*}} x i8], [{{.*}} x i8]* [[NAMEPTR12]], i32 0, i32 0), i[[SZ]] 0 }, section ".omp_offloading.entries", align 1
    151 
    152 // CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
    153 // CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
    154 // CHECK: [[DEVBEGIN:@.+]] = external constant i8
    155 // CHECK: [[DEVEND:@.+]] = external constant i8
    156 // CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }]
    157 // CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }
    158 
    159 // We have 4 initializers, one for the 500 priority, another one for 501, or more for the default priority, and the last one for the offloading registration function.
    160 // CHECK: @llvm.global_ctors = appending global [4 x { i32, void ()*, i8* }] [
    161 // CHECK-SAME: { i32, void ()*, i8* } { i32 500, void ()* [[P500:@[^,]+]], i8* null },
    162 // CHECK-SAME: { i32, void ()*, i8* } { i32 501, void ()* [[P501:@[^,]+]], i8* null },
    163 // CHECK-SAME: { i32, void ()*, i8* } { i32 65535, void ()* [[PMAX:@[^,]+]], i8* null },
    164 // CHECK-SAME: { i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* [[REGFN:@.+]] to void ()*), i8* null }]
    165 
    166 // CHECK-NTARGET: @llvm.global_ctors = appending global [3   x { i32, void ()*, i8* }] [
    167 
    168 extern int *R;
    169 
    170 struct SA {
    171   int arr[4];
    172   void foo() {
    173     int a = *R;
    174     a += 1;
    175     *R = a;
    176   }
    177   SA() {
    178     int a = *R;
    179     a += 2;
    180     *R = a;
    181   }
    182   ~SA() {
    183     int a = *R;
    184     a += 3;
    185     *R = a;
    186   }
    187 };
    188 
    189 struct SB {
    190   int arr[8];
    191   void foo() {
    192     int a = *R;
    193     #pragma omp target
    194     a += 4;
    195     *R = a;
    196   }
    197   SB() {
    198     int a = *R;
    199     a += 5;
    200     *R = a;
    201   }
    202   ~SB() {
    203     int a = *R;
    204     a += 6;
    205     *R = a;
    206   }
    207 };
    208 
    209 struct SC {
    210   int arr[16];
    211   void foo() {
    212     int a = *R;
    213     a += 7;
    214     *R = a;
    215   }
    216   SC() {
    217     int a = *R;
    218     #pragma omp target
    219     a += 8;
    220     *R = a;
    221   }
    222   ~SC() {
    223     int a = *R;
    224     a += 9;
    225     *R = a;
    226   }
    227 };
    228 
    229 struct SD {
    230   int arr[32];
    231   void foo() {
    232     int a = *R;
    233     a += 10;
    234     *R = a;
    235   }
    236   SD() {
    237     int a = *R;
    238     a += 11;
    239     *R = a;
    240   }
    241   ~SD() {
    242     int a = *R;
    243     #pragma omp target
    244     a += 12;
    245     *R = a;
    246   }
    247 };
    248 
    249 struct SE {
    250   int arr[64];
    251   void foo() {
    252     int a = *R;
    253     #pragma omp target if(0)
    254     a += 13;
    255     *R = a;
    256   }
    257   SE() {
    258     int a = *R;
    259     #pragma omp target
    260     a += 14;
    261     *R = a;
    262   }
    263   ~SE() {
    264     int a = *R;
    265     #pragma omp target
    266     a += 15;
    267     *R = a;
    268   }
    269 };
    270 
    271 template <int x>
    272 struct ST {
    273   int arr[128 + x];
    274   void foo() {
    275     int a = *R;
    276     #pragma omp target
    277     a += 16 + x;
    278     *R = a;
    279   }
    280   ST() {
    281     int a = *R;
    282     #pragma omp target
    283     a += 17 + x;
    284     *R = a;
    285   }
    286   ~ST() {
    287     int a = *R;
    288     #pragma omp target
    289     a += 18 + x;
    290     *R = a;
    291   }
    292 };
    293 
    294 // We have to make sure we us all the target regions:
    295 //CHECK-DAG: define internal void @[[NAME1]](
    296 //CHECK-DAG: call void @[[NAME1]](
    297 //CHECK-DAG: define internal void @[[NAME2]](
    298 //CHECK-DAG: call void @[[NAME2]](
    299 //CHECK-DAG: define internal void @[[NAME3]](
    300 //CHECK-DAG: call void @[[NAME3]](
    301 //CHECK-DAG: define internal void @[[NAME4]](
    302 //CHECK-DAG: call void @[[NAME4]](
    303 //CHECK-DAG: define internal void @[[NAME5]](
    304 //CHECK-DAG: call void @[[NAME5]](
    305 //CHECK-DAG: define internal void @[[NAME6]](
    306 //CHECK-DAG: call void @[[NAME6]](
    307 //CHECK-DAG: define internal void @[[NAME7]](
    308 //CHECK-DAG: call void @[[NAME7]](
    309 //CHECK-DAG: define internal void @[[NAME8]](
    310 //CHECK-DAG: call void @[[NAME8]](
    311 //CHECK-DAG: define internal void @[[NAME9]](
    312 //CHECK-DAG: call void @[[NAME9]](
    313 //CHECK-DAG: define internal void @[[NAME10]](
    314 //CHECK-DAG: call void @[[NAME10]](
    315 //CHECK-DAG: define internal void @[[NAME11]](
    316 //CHECK-DAG: call void @[[NAME11]](
    317 //CHECK-DAG: define internal void @[[NAME12]](
    318 //CHECK-DAG: call void @[[NAME12]](
    319 
    320 //TCHECK-DAG: define void @[[NAME1]](
    321 //TCHECK-DAG: define void @[[NAME2]](
    322 //TCHECK-DAG: define void @[[NAME3]](
    323 //TCHECK-DAG: define void @[[NAME4]](
    324 //TCHECK-DAG: define void @[[NAME5]](
    325 //TCHECK-DAG: define void @[[NAME6]](
    326 //TCHECK-DAG: define void @[[NAME7]](
    327 //TCHECK-DAG: define void @[[NAME8]](
    328 //TCHECK-DAG: define void @[[NAME9]](
    329 //TCHECK-DAG: define void @[[NAME10]](
    330 //TCHECK-DAG: define void @[[NAME11]](
    331 //TCHECK-DAG: define void @[[NAME12]](
    332 
    333 // CHECK-NTARGET-NOT: __tgt_target
    334 // CHECK-NTARGET-NOT: __tgt_register_lib
    335 // CHECK-NTARGET-NOT: __tgt_unregister_lib
    336 
    337 // TCHECK-NOT: __tgt_target
    338 // TCHECK-NOT: __tgt_register_lib
    339 // TCHECK-NOT: __tgt_unregister_lib
    340 
    341 // We have 2 initializers with priority 500
    342 //CHECK: define internal void [[P500]](
    343 //CHECK:     call void @{{.+}}()
    344 //CHECK:     call void @{{.+}}()
    345 //CHECK-NOT: call void @{{.+}}()
    346 //CHECK:     ret void
    347 
    348 // We have 1 initializers with priority 501
    349 //CHECK: define internal void [[P501]](
    350 //CHECK:     call void @{{.+}}()
    351 //CHECK-NOT: call void @{{.+}}()
    352 //CHECK:     ret void
    353 
    354 // We have 6 initializers with default priority
    355 //CHECK: define internal void [[PMAX]](
    356 //CHECK:     call void @{{.+}}()
    357 //CHECK:     call void @{{.+}}()
    358 //CHECK:     call void @{{.+}}()
    359 //CHECK:     call void @{{.+}}()
    360 //CHECK:     call void @{{.+}}()
    361 //CHECK:     call void @{{.+}}()
    362 //CHECK-NOT: call void @{{.+}}()
    363 //CHECK:     ret void
    364 
    365 // Check registration and unregistration
    366 
    367 //CHECK:     define internal void [[UNREGFN:@.+]](i8*)
    368 //CHECK:     call i32 @__tgt_unregister_lib([[DSCTY]]* [[DESC]])
    369 //CHECK:     ret void
    370 //CHECK:     declare i32 @__tgt_unregister_lib([[DSCTY]]*)
    371 
    372 //CHECK:     define internal void [[REGFN]](i8*)
    373 //CHECK:     call i32 @__tgt_register_lib([[DSCTY]]* [[DESC]])
    374 //CHECK:     call i32 @__cxa_atexit(void (i8*)* [[UNREGFN]], i8* bitcast ([[DSCTY]]* [[DESC]] to i8*),
    375 //CHECK:     ret void
    376 //CHECK:     declare i32 @__tgt_register_lib([[DSCTY]]*)
    377 
    378 static __attribute__((init_priority(500))) SA a1;
    379 SA a2;
    380 SB __attribute__((init_priority(500))) b1;
    381 SB __attribute__((init_priority(501))) b2;
    382 static SC c1;
    383 SD d1;
    384 SE e1;
    385 ST<100> t1;
    386 ST<1000> t2;
    387 
    388 
    389 int bar(int a){
    390   int r = a;
    391 
    392   a1.foo();
    393   a2.foo();
    394   b1.foo();
    395   b2.foo();
    396   c1.foo();
    397   d1.foo();
    398   e1.foo();
    399   t1.foo();
    400   t2.foo();
    401 
    402   #pragma omp target
    403   ++r;
    404 
    405   return r + *R;
    406 }
    407 
    408 // Check metadata is properly generated:
    409 // CHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
    410 // CHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 193, i32 {{[0-9]+}}}
    411 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 243, i32 {{[0-9]+}}}
    412 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 259, i32 {{[0-9]+}}}
    413 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 265, i32 {{[0-9]+}}}
    414 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 276, i32 {{[0-9]+}}}
    415 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 282, i32 {{[0-9]+}}}
    416 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 402, i32 {{[0-9]+}}}
    417 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 288, i32 {{[0-9]+}}}
    418 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 282, i32 {{[0-9]+}}}
    419 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 288, i32 {{[0-9]+}}}
    420 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 276, i32 {{[0-9]+}}}
    421 // CHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 218, i32 {{[0-9]+}}}
    422 
    423 // TCHECK:     !omp_offload.info = !{!{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}, !{{[0-9]+}}}
    424 // TCHECK-DAG: = !{i32 0, i32 [[DEVID:-?[0-9]+]], i32 [[FILEID:-?[0-9]+]], !"_ZN2SB3fooEv", i32 193, i32 {{[0-9]+}}}
    425 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SDD1Ev", i32 243, i32 {{[0-9]+}}}
    426 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SEC1Ev", i32 259, i32 {{[0-9]+}}}
    427 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SED1Ev", i32 265, i32 {{[0-9]+}}}
    428 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EE3fooEv", i32 276, i32 {{[0-9]+}}}
    429 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EEC1Ev", i32 282, i32 {{[0-9]+}}}
    430 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_Z3bari", i32 402, i32 {{[0-9]+}}}
    431 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EED1Ev", i32 288, i32 {{[0-9]+}}}
    432 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EEC1Ev", i32 282, i32 {{[0-9]+}}}
    433 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi1000EED1Ev", i32 288, i32 {{[0-9]+}}}
    434 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2STILi100EE3fooEv", i32 276, i32 {{[0-9]+}}}
    435 // TCHECK-DAG: = !{i32 0, i32 [[DEVID]], i32 [[FILEID]], !"_ZN2SCC1Ev", i32 218, i32 {{[0-9]+}}}
    436 
    437 #endif
    438