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