clang/test/OpenMP/target_codegen.cpp

   1 // Test host codegen.
   2 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP45
   3 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP45
   5 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP45
   6 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -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 -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP45
   8
   9 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP50
  10 // RUN: %clang_cc1 -fopenmp -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 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP50
  12 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP50
  13 // RUN: %clang_cc1 -fopenmp -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 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP50
  15
  16 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP51
  17 // RUN: %clang_cc1 -fopenmp -x c++ -fopenmp-version=51 -D_DOMP51 -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
  18 // RUN: %clang_cc1 -fopenmp -x c++ -fopenmp-version=51 -D_DOMP51 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64 --check-prefix OMP51
  19 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP51
  20 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
  21 // RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -D_DOMP51 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32 --check-prefix OMP51
  22
  23 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  24 // 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
  25 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  26 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  27 // 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
  28 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  29 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  30
  31 // Test target codegen - host bc file has to be created first.
  32 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  33 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
  34 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  35 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
  36 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  37 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
  38 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  39 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
  40
  41 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  42 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
  43 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  44 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
  45 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  46 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
  47 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  48 // RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
  49
  50 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
  51 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  52 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
  53 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  54 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
  55 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  56 // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
  57 // RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY1 %s
  58 // SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
  59
  60 // expected-no-diagnostics
  61 #ifndef HEADER
  62 #define HEADER
  63
  64 // CHECK-DAG: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }
  65 // CHECK-DAG: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%.+]], [[KMP_PRIVATES_T:%.+]] }
  66 // CHECK-DAG: [[KMP_TASK_T]] = type { ptr, ptr, i32, {{%.+}}, {{%.+}} }
  67 // CHECK-DAG: [[TT:%.+]] = type { i64, i8 }
  68 // CHECK-DAG: [[S1:%.+]] = type { double }
  69 // CHECK-DAG: [[S2:%.+]] = type { i32, i32, i32 }
  70 // CHECK-DAG: [[ENTTY:%.+]] = type { ptr, ptr, i[[SZ:32|64]], i32, i32 }
  71 // CHECK-DAG: [[ANON_T:%.+]] = type { ptr, i32, i32 }
  72 // CHECK-32-DAG: [[KMP_PRIVATES_T]] = type { [2 x i64], ptr, i32, [2 x ptr], [2 x ptr] }
  73 // CHECK-64-DAG: [[KMP_PRIVATES_T]] = type { ptr, [2 x ptr], [2 x ptr], [2 x i64], i32 }
  74
  75 // TCHECK: [[ENTTY:%.+]] = type { ptr, ptr, i{{32|64}}, i32, i32 }
  76
  77 // We have 9 target regions, but only 8 that actually will generate offloading
  78 // code and have mapped arguments, and only 6 have all-constant map sizes.
  79
  80 // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 4]
  81 // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800]
  82 // CHECK-DAG: [[SIZET2:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] [i64 2]
  83 // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i64] [i64 800]
  84 // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2]
  85 // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
  86 // CHECK-DAG: [[SIZET4:@.+]] = private unnamed_addr constant [9 x i64] [i64 4, i64 40, i64 {{4|8}}, i64 0, i64 400, i64 {{4|8}}, i64 {{4|8}}, i64 0, i64 {{12|16}}]
  87 // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547]
  88 // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40]
  89 // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547]
  90 // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40]
  91 // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547]
  92 // CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [5 x i64] [i64 8, i64 4, i64 {{4|8}}, i64 {{4|8}}, i64 0]
  93 // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547]
  94 // CHECK-DAG: [[SIZET9:@.+]] = private unnamed_addr constant [1 x i64] [i64 12]
  95 // CHECK-DAG: [[MAPT10:@.+]] = private unnamed_addr constant [1 x i64] [i64 35]
  96 // CHECK-DAG: @{{.*}} = weak constant i8 0
  97 // CHECK-DAG: @{{.*}} = weak constant i8 0
  98 // CHECK-DAG: @{{.*}} = weak constant i8 0
  99 // CHECK-DAG: @{{.*}} = weak constant i8 0
 100 // CHECK-DAG: @{{.*}} = weak constant i8 0
 101 // CHECK-DAG: @{{.*}} = weak constant i8 0
 102 // CHECK-DAG: @{{.*}} = weak constant i8 0
 103 // CHECK-DAG: @{{.*}} = weak constant i8 0
 104
 105 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 106 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 107 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 108 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 109 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 110 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 111 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 112 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 113 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 114 // TCHECK: @{{.+}} = weak constant [[ENTTY]]
 115 // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]]
 116
 117 // Check target registration is registered as a Ctor.
 118 // CHECK: appending global [1 x { i32, ptr, ptr }] [{ i32, ptr, ptr } { i32 0, ptr @.omp_offloading.requires_reg, ptr null }]
 119
 120
 121 template<typename tx, typename ty>
 122 struct TT{
 123   tx X;
 124   ty Y;
 125 };
 126
 127 int global;
 128 extern int global;
 129
 130 // CHECK: define {{.*}}[[FOO:@.+]](
 131 int foo(int n) {
 132   // CHECK: [[OFFLOADBPTR:%.+]] = alloca [2 x ptr], align
 133   // CHECK: [[OFFLOADPTR:%.+]] = alloca [2 x ptr], align
 134   // CHECK: [[OFFLOADMAPPER:%.+]] = alloca [2 x ptr], align
 135   int a = 0;
 136   short aa = 0;
 137   float b[10];
 138   float bn[n];
 139   double c[5][10];
 140   double cn[5][n];
 141   TT<long long, char> d;
 142   static long *plocal;
 143
 144 // CHECK:       [[ADD:%.+]] = add nsw i32
 145 // CHECK:       store i32 [[ADD]], ptr [[DEVICE_CAP:%.+]],
 146 // CHECK:       [[DEV:%.+]] = load i32, ptr [[DEVICE_CAP]],
 147 // CHECK:       [[DEVICE:%.+]] = sext i32 [[DEV]] to i64
 148 // CHECK:       [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE]], i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 149 // CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 150 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 151 // CHECK:       [[FAIL]]
 152 // CHECK:       call void [[HVT0:@.+]]()
 153 // CHECK-NEXT:  br label %[[END]]
 154 // CHECK:       [[END]]
 155 #pragma omp target device(global + a)
 156   {
 157   }
 158
 159   // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0
 160   // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0
 161   // CHECK: [[BPRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADBPTR]], i32 0, i32 0
 162   // CHECK: [[PRGEP:%.+]] = getelementptr inbounds [2 x ptr], ptr [[OFFLOADPTR]], i32 0, i32 0
 163   // CHECK: [[DEVICE:%.+]] = sext i32 {{%.+}} to i64
 164   // CHECK-32: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 %0, i32 1, i32 60, i32 12, ptr [[OMP_TASK_ENTRY:@.+]], i64 [[DEVICE]])
 165   // CHECK-64: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 %0, i32 1, i64 104, i64 16, ptr [[OMP_TASK_ENTRY:@.+]], i64 [[DEVICE]])
 166   // CHECK: [[TASK_WITH_PRIVATES_GEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK]], i32 0, i32 1
 167   // CHECK-32: [[SIZEGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 0
 168   // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[SIZEGEP]], ptr align 4 [[SIZET]], i32 16, i1 false)
 169   // CHECK-32: [[FPBPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 3
 170   // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[FPBPRGEP]], ptr align 4 [[BPRGEP]], i32 8, i1 false)
 171   // CHECK-32: [[FPPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 4
 172   // CHECK-32: call void @llvm.memcpy.p0.p0.i32(ptr align 4 [[FPPRGEP]], ptr align 4 [[PRGEP]], i32 8, i1 false)
 173   // CHECK-64: [[FPBPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 1
 174   // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[FPBPRGEP]], ptr align 8 [[BPRGEP]], i64 16, i1 false)
 175   // CHECK-64: [[FPPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 2
 176   // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[FPPRGEP]], ptr align 8 [[PRGEP]], i64 16, i1 false)
 177   // CHECK-64: [[SIZEGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], ptr [[TASK_WITH_PRIVATES_GEP]], i32 0, i32 3
 178   // CHECK-64: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[SIZEGEP]], ptr align 8 [[SIZET]], i64 16, i1 false)
 179   // CHECK: call i32 @__kmpc_omp_task(ptr {{.+}}, i32 {{.+}}, ptr [[TASK]])
 180   #pragma omp target device(global + a) nowait
 181   {
 182     static int local1;
 183     *plocal = global;
 184     local1 = global;
 185   }
 186
 187   // CHECK:       call void [[HVT1:@.+]](i[[SZ]] {{[^,]+}})
 188   #pragma omp target if(0) firstprivate(global)
 189   {
 190     global += 1;
 191   }
 192
 193 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 194 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 195 // CHECK-DAG:   store ptr [[BP:%.+]], ptr [[BPARG]]
 196 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 197 // CHECK-DAG:   store ptr [[P:%.+]], ptr [[PARG]]
 198 // CHECK-DAG:   [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPR:%[^,]+]], i32 0, i32 0
 199 // CHECK-DAG:   [[P]] = getelementptr inbounds [1 x ptr], ptr [[PR:%[^,]+]], i32 0, i32 0
 200 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BPR]], i32 0, i32 [[IDX0:[0-9]+]]
 201 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[PR]], i32 0, i32 [[IDX0]]
 202 // CHECK-DAG:   store i[[SZ]] [[BP0:%[^,]+]], ptr [[BPADDR0]]
 203 // CHECK-DAG:   store i[[SZ]] [[P0:%[^,]+]], ptr [[PADDR0]]
 204
 205 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 206 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 207 // CHECK:       [[FAIL]]
 208 // CHECK:       call void [[HVT2:@.+]](i[[SZ]] {{[^,]+}})
 209 // CHECK-NEXT:  br label %[[END]]
 210 // CHECK:       [[END]]
 211 #pragma omp target if (1)
 212   {
 213     aa += 1;
 214   }
 215
 216 // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 10
 217 // CHECK:       br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 218 // CHECK:       [[IFTHEN]]
 219 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 220 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 221 // CHECK-DAG:   store ptr [[BPR:%.+]], ptr [[BPARG]]
 222 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 223 // CHECK-DAG:   store ptr [[PR:%.+]], ptr [[PARG]]
 224 // CHECK-DAG:   [[BPR]] = getelementptr inbounds [2 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
 225 // CHECK-DAG:   [[PR]] = getelementptr inbounds [2 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
 226
 227 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0
 228 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0
 229 // CHECK-DAG:   store i[[SZ]] [[BP0:%[^,]+]], ptr [[BPADDR0]]
 230 // CHECK-DAG:   store i[[SZ]] [[P0:%[^,]+]], ptr [[PADDR0]]
 231
 232 // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1
 233 // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
 234 // CHECK-DAG:   store i[[SZ]] [[BP1:%[^,]+]], ptr [[BPADDR1]]
 235 // CHECK-DAG:   store i[[SZ]] [[P1:%[^,]+]], ptr [[PADDR1]]
 236 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 237 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
 238 // CHECK:       [[FAIL]]
 239 // CHECK:       call void [[HVT3:@.+]]({{[^,]+}}, {{[^,]+}})
 240 // CHECK-NEXT:  br label %[[END]]
 241 // CHECK:       [[END]]
 242 // CHECK-NEXT:  br label %[[IFEND:.+]]
 243 // CHECK:       [[IFELSE]]
 244 // CHECK:       call void [[HVT3]]({{[^,]+}}, {{[^,]+}})
 245 // CHECK-NEXT:  br label %[[IFEND]]
 246
 247 // CHECK:       [[IFEND]]
 248 #pragma omp target if (n > 10)
 249   {
 250     a += 1;
 251     aa += 1;
 252   }
 253
 254   // We capture 3 VLA sizes in this target region
 255   // CHECK-64:       [[A_VAL:%.+]] = load i32, ptr %{{.+}},
 256   // CHECK-64:       store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],
 257   // CHECK-64:       [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],
 258
 259   // CHECK-32:       [[A_VAL:%.+]] = load i32, ptr %{{.+}},
 260   // CHECK-32:       store i32 [[A_VAL]], ptr [[A_CADDR:%.+]],
 261   // CHECK-32:       [[A_CVAL:%.+]] = load i[[SZ]], ptr [[A_CADDR]],
 262
 263   // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 20
 264   // CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]
 265   // CHECK:       [[TRY]]
 266   // CHECK-64:    [[BNSIZE:%.+]] = mul nuw i64 [[VLA0:%.+]], 4
 267   // CHECK-32:    [[BNSZSIZE:%.+]] = mul nuw i32 [[VLA0:%.+]], 4
 268   // CHECK-32:    [[BNSIZE:%.+]] = sext i32 [[BNSZSIZE]] to i64
 269   // CHECK:       [[CNELEMSIZE2:%.+]] = mul nuw i[[SZ]] 5, [[VLA1:%.+]]
 270   // CHECK-64:    [[CNSIZE:%.+]] = mul nuw i64 [[CNELEMSIZE2]], 8
 271   // CHECK-32:    [[CNSZSIZE:%.+]] = mul nuw i32 [[CNELEMSIZE2]], 8
 272   // CHECK-32:    [[CNSIZE:%.+]] = sext i32 [[CNSZSIZE]] to i64
 273
 274 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 275 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 276 // CHECK-DAG:   store ptr [[BPR:%.+]], ptr [[BPARG]]
 277 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 278 // CHECK-DAG:   store ptr [[PR:%.+]], ptr [[PARG]]
 279 // CHECK-DAG:   [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 280 // CHECK-DAG:   store ptr [[SZ4:%.+]], ptr [[SARG]]
 281 // CHECK-DAG:   [[BPR]] = getelementptr inbounds [9 x ptr], ptr [[BP:%[^,]+]], i32 0, i32 0
 282 // CHECK-DAG:   [[PR]] = getelementptr inbounds [9 x ptr], ptr [[P:%[^,]+]], i32 0, i32 0
 283 // CHECK-DAG:   [[SZ4]] = getelementptr inbounds [9 x i64], ptr [[PSZ:%[^,]+]], i32 0, i32 0
 284
 285 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:0]]
 286 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]
 287 // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:1]]
 288 // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]
 289 // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:2]]
 290 // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]
 291 // CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX3:3]]
 292 // CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]
 293 // CHECK-DAG:   [[PSZ3:%.+]] = getelementptr inbounds [9 x i64], ptr [[PSZ]], i32 0, i32 [[IDX3]]
 294 // CHECK-DAG:   [[BPADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:4]]
 295 // CHECK-DAG:   [[PADDR4:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]
 296 // CHECK-DAG:   [[BPADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX5:5]]
 297 // CHECK-DAG:   [[PADDR5:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX5]]
 298 // CHECK-DAG:   [[BPADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX6:6]]
 299 // CHECK-DAG:   [[PADDR6:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX6]]
 300 // CHECK-DAG:   [[BPADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX7:7]]
 301 // CHECK-DAG:   [[PADDR7:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX7]]
 302 // CHECK-DAG:   [[PSZ7:%.+]] = getelementptr inbounds [9 x i64], ptr [[PSZ]], i32 0, i32 [[IDX7]]
 303 // CHECK-DAG:   [[BPADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[BP]], i32 0, i32 [[IDX8:8]]
 304 // CHECK-DAG:   [[PADDR8:%.+]] = getelementptr inbounds [9 x ptr], ptr [[P]], i32 0, i32 [[IDX8]]
 305
 306 // The names below are not necessarily consistent with the names used for the
 307 // addresses above as some are repeated.
 308 // CHECK-DAG:   store i[[SZ]] [[VLA0]], ptr [[BPADDR2]]
 309 // CHECK-DAG:   store i[[SZ]] [[VLA0]], ptr [[PADDR2]]
 310
 311 // CHECK-DAG:   store i[[SZ]] [[VLA1]], ptr [[BPADDR6]]
 312 // CHECK-DAG:   store i[[SZ]] [[VLA1]], ptr [[PADDR6]]
 313
 314 // CHECK-DAG:   store i[[SZ]] 5, ptr [[BPADDR5]]
 315 // CHECK-DAG:   store i[[SZ]] 5, ptr [[PADDR5]]
 316
 317 // CHECK-DAG:   store i[[SZ]] [[A_CVAL]], ptr [[BPADDR0]]
 318 // CHECK-DAG:   store i[[SZ]] [[A_CVAL]], ptr [[PADDR0]]
 319
 320 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR1]]
 321 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR1]]
 322
 323 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR3]]
 324 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR3]]
 325 // CHECK-DAG:   store i64 [[BNSIZE]], ptr [[PSZ3]]
 326
 327 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR4]]
 328 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR4]]
 329
 330 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR7]]
 331 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR7]]
 332 // CHECK-DAG:   store i64 [[CNSIZE]], ptr [[PSZ7]]
 333
 334 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR8]]
 335 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR8]]
 336
 337 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 338 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
 339 // CHECK:       [[FAIL]]
 340 // CHECK:       call void [[HVT4:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 341 // CHECK-NEXT:  br label %[[END]]
 342 // CHECK:       [[END]]
 343 // CHECK-NEXT:  br label %[[IFEND:.+]]
 344 // CHECK:       [[IFELSE]]
 345 // CHECK:       call void [[HVT4]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 346 // CHECK-NEXT:  br label %[[IFEND]]
 347
 348 // CHECK:       [[IFEND]]
 349 #pragma omp target if (n > 20)
 350   {
 351     a += 1;
 352     b[2] += 1.0;
 353     bn[3] += 1.0;
 354     c[1][2] += 1.0;
 355     cn[1][3] += 1.0;
 356     d.X += 1;
 357     d.Y += 1;
 358   }
 359
 360   return a;
 361 }
 362
 363 // Check that the offloading functions are emitted and that the arguments are
 364 // correct and loaded correctly for the target regions in foo().
 365
 366 // CHECK:       define internal void [[HVT0]]()
 367
 368 // CHECK: define internal void [[HVT0_:@.+]](ptr noundef {{%[^,]+}}, i[[SZ]] noundef {{%[^,]+}})
 369 // CHECK: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
 370 // CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:%.+]], i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 371 // CHECK-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 372 // CHECK-DAG: store ptr [[BPR:%.+]], ptr [[BPARG]]
 373 // CHECK-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 374 // CHECK-DAG: store ptr [[PR:%.+]], ptr [[PARG]]
 375 // CHECK-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 376 // CHECK-DAG: store ptr [[SIZE:%.+]], ptr [[SARG]]
 377 // CHECK-DAG: [[DEVICE]] = sext i32 [[DEV:%.+]] to i64
 378 // CHECK-DAG: [[DEV]] = load i32, ptr [[DEVADDR:%.+]], align
 379 // CHECK-DAG: [[DEVADDR]] = getelementptr inbounds [[ANON_T]], ptr {{%.+}}, i32 0, i32 2
 380 // CHECK-DAG: [[BPR]] = load ptr, ptr [[FPPTR_BPR:%.+]], align
 381 // CHECK-DAG: [[PR]] = load ptr, ptr [[FPPTR_PR:%.+]], align
 382 // CHECK-DAG: [[SIZE]] = load ptr, ptr [[FPPTR_SIZE:%.+]], align
 383 // CHECK-DAG: call void {{%[0-9]+}}(ptr {{%[^,]+}}, ptr [[FPPTR_PLOCAL:%.+]], ptr [[FPPTR_GLOBAL:%.+]], ptr [[FPPTR_BPR]], ptr [[FPPTR_PR]], ptr [[FPPTR_SIZE]])
 384 // CHECK-DAG: [[PLOCALADDR:%.+]] = load ptr, ptr [[FPPTR_PLOCAL]], align
 385 // CHECK-DAG: {{%.+}} = load ptr, ptr [[FPPTR_GLOBAL:%.+]], align
 386 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 387 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 388 // CHECK: [[FAIL]]
 389 // CHECK: [[PLOCAL:%.+]] = load ptr, ptr [[PLOCALADDR]], align
 390 // CHECK: [[GLOBAL:%.+]] = load i32, ptr {{@.+}}, align
 391 // CHECK-32: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align
 392 // CHECK-64: store i32 [[GLOBAL]], ptr [[GLOBALCAST:%.+]], align
 393 // CHECK: [[GLOBAL:%.+]] = load i[[SZ]], ptr [[GLOBALCAST]], align
 394 // CHECK: call void [[HVT0_]](ptr [[PLOCAL]], i[[SZ]] [[GLOBAL]])
 395 // CHECK-NEXT: br label %[[END]]
 396 // CHECK: [[END]]
 397
 398 // CHECK:       define internal void [[HVT1]](i[[SZ]] noundef %{{.+}})
 399 // Create stack storage and store argument in there.
 400 // CHECK:       [[AA_ADDR:%.+]] = alloca i[[SZ]], align
 401 // CHECK:       store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
 402 // CHECK-64:    load i32, ptr [[AA_ADDR]], align
 403 // CHECK-32:    load i32, ptr [[AA_ADDR]], align
 404
 405 // CHECK:       define internal void [[HVT2]](i[[SZ]] noundef %{{.+}})
 406 // Create stack storage and store argument in there.
 407 // CHECK:       [[AA_ADDR:%.+]] = alloca i[[SZ]], align
 408 // CHECK:       store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
 409 // CHECK:       load i16, ptr [[AA_ADDR]], align
 410
 411 // CHECK:       define internal void [[HVT3]]
 412 // Create stack storage and store argument in there.
 413 // CHECK:       [[A_ADDR:%.+]] = alloca i[[SZ]], align
 414 // CHECK:       [[AA_ADDR:%.+]] = alloca i[[SZ]], align
 415 // CHECK-DAG:   store i[[SZ]] %{{.+}}, ptr [[A_ADDR]], align
 416 // CHECK-DAG:   store i[[SZ]] %{{.+}}, ptr [[AA_ADDR]], align
 417 // CHECK-64-DAG:load i32, ptr [[A_ADDR]], align
 418 // CHECK-32-DAG:load i32, ptr [[A_ADDR]], align
 419 // CHECK-DAG:   load i16, ptr [[AA_ADDR]], align
 420
 421 // CHECK:       define internal void [[HVT4]]
 422 // Create local storage for each capture.
 423 // CHECK:       [[LOCAL_A:%.+]] = alloca i[[SZ]]
 424 // CHECK:       [[LOCAL_B:%.+]] = alloca ptr
 425 // CHECK:       [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
 426 // CHECK:       [[LOCAL_BN:%.+]] = alloca ptr
 427 // CHECK:       [[LOCAL_C:%.+]] = alloca ptr
 428 // CHECK:       [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
 429 // CHECK:       [[LOCAL_VLA3:%.+]] = alloca i[[SZ]]
 430 // CHECK:       [[LOCAL_CN:%.+]] = alloca ptr
 431 // CHECK:       [[LOCAL_D:%.+]] = alloca ptr
 432 // CHECK-DAG:   store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
 433 // CHECK-DAG:   store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
 434 // CHECK-DAG:   store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]
 435 // CHECK-DAG:   store ptr [[ARG_BN:%.+]], ptr [[LOCAL_BN]]
 436 // CHECK-DAG:   store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]
 437 // CHECK-DAG:   store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]
 438 // CHECK-DAG:   store i[[SZ]] [[ARG_VLA3:%.+]], ptr [[LOCAL_VLA3]]
 439 // CHECK-DAG:   store ptr [[ARG_CN:%.+]], ptr [[LOCAL_CN]]
 440 // CHECK-DAG:   store ptr [[ARG_D:%.+]], ptr [[LOCAL_D]]
 441
 442 // CHECK-DAG:   [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
 443 // CHECK-DAG:   [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],
 444 // CHECK-DAG:   [[REF_BN:%.+]] = load ptr, ptr [[LOCAL_BN]],
 445 // CHECK-DAG:   [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],
 446 // CHECK-DAG:   [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],
 447 // CHECK-DAG:   [[VAL_VLA3:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA3]],
 448 // CHECK-DAG:   [[REF_CN:%.+]] = load ptr, ptr [[LOCAL_CN]],
 449 // CHECK-DAG:   [[REF_D:%.+]] = load ptr, ptr [[LOCAL_D]],
 450
 451 // Use captures.
 452 // CHECK-64-DAG:   load i32, ptr [[LOCAL_A]]
 453 // CHECK-32-DAG:   load i32, ptr [[LOCAL_A]]
 454 // CHECK-DAG:   getelementptr inbounds [10 x float], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
 455 // CHECK-DAG:   getelementptr inbounds float, ptr [[REF_BN]], i[[SZ]] 3
 456 // CHECK-DAG:   getelementptr inbounds [5 x [10 x double]], ptr [[REF_C]], i[[SZ]] 0, i[[SZ]] 1
 457 // CHECK-DAG:   getelementptr inbounds double, ptr [[REF_CN]], i[[SZ]] %{{.+}}
 458 // CHECK-DAG:   getelementptr inbounds [[TT]], ptr [[REF_D]], i32 0, i32 0
 459
 460 template<typename tx>
 461 tx ftemplate(int n) {
 462   tx a = 0;
 463   short aa = 0;
 464   tx b[10];
 465
 466   #pragma omp target if(n>40)
 467   {
 468     a += 1;
 469     aa += 1;
 470     b[2] += 1;
 471   }
 472
 473   return a;
 474 }
 475
 476 static
 477 int fstatic(int n) {
 478   int a = 0;
 479   short aa = 0;
 480   char aaa = 0;
 481   int b[10];
 482
 483   #pragma omp target if(n>50)
 484   {
 485     a += 1;
 486     aa += 1;
 487     aaa += 1;
 488     b[2] += 1;
 489   }
 490
 491   return a;
 492 }
 493
 494 struct S1 {
 495   double a;
 496
 497   int r1(int n){
 498     int b = n+1;
 499     short int c[2][n];
 500
 501     #pragma omp target if(n>60)
 502     {
 503       this->a = (double)b + 1.5;
 504       c[1][1] = ++a;
 505     }
 506
 507     return c[1][1] + (int)b;
 508   }
 509 };
 510
 511 // CHECK: define {{.*}}@{{.*}}bar{{.*}}
 512 int bar(int n){
 513   int a = 0;
 514
 515   // CHECK: call {{.*}}i32 [[FOO]](i32 {{.*}})
 516   a += foo(n);
 517
 518   S1 S;
 519   // CHECK: call {{.*}}i32 [[FS1:@.+]](ptr {{.*}}, i32 {{.*}})
 520   a += S.r1(n);
 521
 522   // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
 523   a += fstatic(n);
 524
 525   // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
 526   a += ftemplate<int>(n);
 527
 528   return a;
 529 }
 530
 531 //
 532 // CHECK: define {{.*}}[[FS1]]
 533 //
 534 // CHECK:          ptr @llvm.stacksave.p0()
 535 // CHECK-64:       store i32 %{{.+}}, ptr [[B_CADDR:%.+]],
 536 // CHECK-64:       [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_CADDR]],
 537
 538 // CHECK-32:       store i32 %{{.+}}, ptr %__vla_expr
 539 // CHECK-32:       store i32 %{{.+}}, ptr [[B_CADDR:%.+]],
 540 // CHECK-32:       [[B_CVAL:%.+]] = load i[[SZ]], ptr [[B_CADDR]],
 541
 542 // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 60
 543 // CHECK:       br i1 [[IF]], label %[[TRY:[^,]+]], label %[[IFELSE:[^,]+]]
 544 // CHECK:       [[TRY]]
 545 // We capture 2 VLA sizes in this target region
 546 // CHECK:       [[CELEMSIZE2:%.+]] = mul nuw i[[SZ]] 2, [[VLA0:%.+]]
 547 // CHECK-64:    [[CSIZE:%.+]] = mul nuw i64 [[CELEMSIZE2]], 2
 548 // CHECK-32:    [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2
 549 // CHECK-32:    [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64
 550
 551 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 552 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 553 // CHECK-DAG:   store ptr [[BPR:%.+]], ptr [[BPARG]]
 554 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 555 // CHECK-DAG:   store ptr [[PR:%.+]], ptr [[PARG]]
 556 // CHECK-DAG:   [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 557 // CHECK-DAG:   store ptr [[SZ7:%.+]], ptr [[SARG]]
 558 // CHECK-DAG:   [[BPR]] = getelementptr inbounds [5 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 559 // CHECK-DAG:   [[PR]] = getelementptr inbounds [5 x ptr], ptr [[P:%.+]], i32 0, i32 0
 560 // CHECK-DAG:   [[SZ7]] = getelementptr inbounds [5 x i64], ptr [[PSZ:%.+]], i32 0, i32 0
 561 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX0:0]]
 562 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX0]]
 563 // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX1:1]]
 564 // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX1]]
 565 // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX2:2]]
 566 // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX2]]
 567 // CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX3:3]]
 568 // CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX3]]
 569 // CHECK-DAG:   [[BPADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BP]], i32 0, i32 [[IDX4:4]]
 570 // CHECK-DAG:   [[PADDR4:%.+]] = getelementptr inbounds [5 x ptr], ptr [[P]], i32 0, i32 [[IDX4]]
 571 // CHECK-DAG:   [[PSZ4:%.+]] = getelementptr inbounds [5 x i64], ptr [[PSZ:%.+]], i32 0, i32 [[IDX4]]
 572
 573 // The names below are not necessarily consistent with the names used for the
 574 // addresses above as some are repeated.
 575 // CHECK-DAG:   store ptr %{{.+}}, ptr [[BPADDR4]]
 576 // CHECK-DAG:   store ptr %{{.+}}, ptr [[PADDR4]]
 577 // CHECK-DAG:   store i64 [[CSIZE]], ptr [[PSZ4]]
 578
 579 // CHECK-DAG:   store i[[SZ]] [[VLA0]], ptr [[BPADDR3]]
 580 // CHECK-DAG:   store i[[SZ]] [[VLA0]], ptr [[PADDR3]]
 581
 582 // CHECK-DAG:   store i[[SZ]] 2, ptr [[BPADDR2]]
 583 // CHECK-DAG:   store i[[SZ]] 2, ptr [[PADDR2]]
 584
 585 // CHECK-DAG:   store i[[SZ]] [[B_CVAL]], ptr [[BPADDR1]]
 586 // CHECK-DAG:   store i[[SZ]] [[B_CVAL]], ptr [[PADDR1]]
 587
 588 // CHECK-DAG:   store ptr [[THIS:%.+]], ptr [[BPADDR0]]
 589 // CHECK-DAG:   store ptr [[A:%.+]], ptr [[PADDR0]]
 590
 591 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 592 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
 593 // CHECK:       [[FAIL]]
 594 // CHECK:       call void [[HVT7:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 595 // CHECK-NEXT:  br label %[[END]]
 596 // CHECK:       [[END]]
 597 // CHECK-NEXT:  br label %[[IFEND:.+]]
 598 // CHECK:       [[IFELSE]]
 599 // CHECK:       call void [[HVT7]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 600 // CHECK-NEXT:  br label %[[IFEND]]
 601
 602 // CHECK:       [[IFEND]]
 603
 604 //
 605 // CHECK: define {{.*}}[[FSTATIC]]
 606 //
 607 // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 50
 608 // CHECK:       br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 609 // CHECK:       [[IFTHEN]]
 610 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 611 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 612 // CHECK-DAG:   store ptr [[BPR:%.+]], ptr [[BPARG]]
 613 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 614 // CHECK-DAG:   store ptr [[PR:%.+]], ptr [[PARG]]
 615 // CHECK-DAG:   [[BPR]] = getelementptr inbounds [4 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 616 // CHECK-DAG:   [[PR]] = getelementptr inbounds [4 x ptr], ptr [[P:%.+]], i32 0, i32 0
 617
 618 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 0
 619 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 0
 620 // CHECK-DAG:   store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]
 621 // CHECK-DAG:   store i[[SZ]] [[VAL0]], ptr [[PADDR0]]
 622
 623 // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 1
 624 // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 1
 625 // CHECK-DAG:   store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]
 626 // CHECK-DAG:   store i[[SZ]] [[VAL1]], ptr [[PADDR1]]
 627
 628 // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 2
 629 // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 2
 630 // CHECK-DAG:   store i[[SZ]] [[VAL2:%[^,]+]], ptr [[BPADDR2]]
 631 // CHECK-DAG:   store i[[SZ]] [[VAL2]], ptr [[PADDR2]]
 632
 633 // CHECK-DAG:   [[BPADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[BP]], i32 0, i32 3
 634 // CHECK-DAG:   [[PADDR3:%.+]] = getelementptr inbounds [4 x ptr], ptr [[P]], i32 0, i32 3
 635 // CHECK-DAG:   store ptr [[VAL3:%[^,]+]], ptr [[BPADDR3]]
 636 // CHECK-DAG:   store ptr [[VAL3]], ptr [[PADDR3]]
 637
 638 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 639 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
 640 // CHECK:       [[FAIL]]
 641 // CHECK:       call void [[HVT6:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 642 // CHECK-NEXT:  br label %[[END]]
 643 // CHECK:       [[END]]
 644 // CHECK-NEXT:  br label %[[IFEND:.+]]
 645 // CHECK:       [[IFELSE]]
 646 // CHECK:       call void [[HVT6]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 647 // CHECK-NEXT:  br label %[[IFEND]]
 648
 649 // CHECK:       [[IFEND]]
 650
 651 //
 652 // CHECK: define {{.*}}[[FTEMPLATE]]
 653 //
 654 // CHECK:       [[IF:%.+]] = icmp sgt i32 {{[^,]+}}, 40
 655 // CHECK:       br i1 [[IF]], label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 656 // CHECK:       [[IFTHEN]]
 657 // CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 {{.+}}, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 658 // CHECK-DAG:   [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 659 // CHECK-DAG:   store ptr [[BPR:%.+]], ptr [[BPARG]]
 660 // CHECK-DAG:   [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 661 // CHECK-DAG:   store ptr [[PR:%.+]], ptr [[PARG]]
 662 // CHECK-DAG:   [[BPR]] = getelementptr inbounds [3 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 663 // CHECK-DAG:   [[PR]] = getelementptr inbounds [3 x ptr], ptr [[P:%.+]], i32 0, i32 0
 664
 665 // CHECK-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 0
 666 // CHECK-DAG:   [[PADDR0:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 0
 667 // CHECK-DAG:   store i[[SZ]] [[VAL0:%[^,]+]], ptr [[BPADDR0]]
 668 // CHECK-DAG:   store i[[SZ]] [[VAL0]], ptr [[PADDR0]]
 669
 670 // CHECK-DAG:   [[BPADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 1
 671 // CHECK-DAG:   [[PADDR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 1
 672 // CHECK-DAG:   store i[[SZ]] [[VAL1:%[^,]+]], ptr [[BPADDR1]]
 673 // CHECK-DAG:   store i[[SZ]] [[VAL1]], ptr [[PADDR1]]
 674
 675 // CHECK-DAG:   [[BPADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BP]], i32 0, i32 2
 676 // CHECK-DAG:   [[PADDR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[P]], i32 0, i32 2
 677 // CHECK-DAG:   store ptr [[VAL2:%[^,]+]], ptr [[BPADDR2]]
 678 // CHECK-DAG:   store ptr [[VAL2]], ptr [[PADDR2]]
 679
 680 // CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 681 // CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
 682 // CHECK:       [[FAIL]]
 683 // CHECK:       call void [[HVT5:@.+]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 684 // CHECK-NEXT:  br label %[[END]]
 685 // CHECK:       [[END]]
 686 // CHECK-NEXT:  br label %[[IFEND:.+]]
 687 // CHECK:       [[IFELSE]]
 688 // CHECK:       call void [[HVT5]]({{[^,]+}}, {{[^,]+}}, {{[^,]+}})
 689 // CHECK-NEXT:  br label %[[IFEND]]
 690
 691 // CHECK:       [[IFEND]]
 692
 693 // OMP45: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})
 694
 695 // OMP45: define {{.*}}@{{.*}}zee{{.*}}
 696
 697 // OMP45:       [[LOCAL_THIS:%.+]] = alloca ptr
 698 // OMP45:       [[BP:%.+]] = alloca [1 x ptr]
 699 // OMP45:       [[P:%.+]] = alloca [1 x ptr]
 700 // OMP45:       [[LOCAL_THIS1:%.+]] = load ptr, ptr [[LOCAL_THIS]]
 701
 702 // OMP45:       call void @__kmpc_critical(
 703 // OMP45:       [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
 704 // OMP45:       [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
 705
 706 // OMP45-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 707 // OMP45-DAG:   [[PADDR0:%.+]] =  getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 708 // OMP45-DAG:   store ptr [[ARR_IDX]], ptr [[BPADDR0]]
 709 // OMP45-DAG:   store ptr [[ARR_IDX2]], ptr [[PADDR0]]
 710
 711 // OMP45:       [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 712 // OMP45:       [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 713 // OMP45:       [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 714 // OMP45-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 715 // OMP45-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 716 // OMP45:       [[FAIL]]
 717 // OMP45:       call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])
 718 // OMP45-NEXT:  br label %[[END]]
 719 // OMP45:       [[END]]
 720 // OMP45:       call void @__kmpc_end_critical(
 721
 722 // Check that the offloading functions are emitted and that the arguments are
 723 // correct and loaded correctly for the target regions of the callees of bar().
 724
 725 // CHECK:       define internal void [[HVT7]]
 726 // Create local storage for each capture.
 727 // CHECK:       [[LOCAL_THIS:%.+]] = alloca ptr
 728 // CHECK:       [[LOCAL_B:%.+]] = alloca i[[SZ]]
 729 // CHECK:       [[LOCAL_VLA1:%.+]] = alloca i[[SZ]]
 730 // CHECK:       [[LOCAL_VLA2:%.+]] = alloca i[[SZ]]
 731 // CHECK:       [[LOCAL_C:%.+]] = alloca ptr
 732 // CHECK-DAG:   store ptr [[ARG_THIS:%.+]], ptr [[LOCAL_THIS]]
 733 // CHECK-DAG:   store i[[SZ]] [[ARG_B:%.+]], ptr [[LOCAL_B]]
 734 // CHECK-DAG:   store i[[SZ]] [[ARG_VLA1:%.+]], ptr [[LOCAL_VLA1]]
 735 // CHECK-DAG:   store i[[SZ]] [[ARG_VLA2:%.+]], ptr [[LOCAL_VLA2]]
 736 // CHECK-DAG:   store ptr [[ARG_C:%.+]], ptr [[LOCAL_C]]
 737 // Store captures in the context.
 738 // CHECK-DAG:   [[REF_THIS:%.+]] = load ptr, ptr [[LOCAL_THIS]],
 739 // CHECK-DAG:   [[VAL_VLA1:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA1]],
 740 // CHECK-DAG:   [[VAL_VLA2:%.+]] = load i[[SZ]], ptr [[LOCAL_VLA2]],
 741 // CHECK-DAG:   [[REF_C:%.+]] = load ptr, ptr [[LOCAL_C]],
 742 // Use captures.
 743 // CHECK-DAG:   getelementptr inbounds [[S1]], ptr [[REF_THIS]], i32 0, i32 0
 744 // CHECK-64-DAG:load i32, ptr [[LOCAL_B]]
 745 // CHECK-32-DAG:load i32, ptr [[LOCAL_B]]
 746 // CHECK-DAG:   getelementptr inbounds i16, ptr [[REF_C]], i[[SZ]] %{{.+}}
 747
 748
 749 // CHECK:       define internal void [[HVT6]]
 750 // Create local storage for each capture.
 751 // CHECK:       [[LOCAL_A:%.+]] = alloca i[[SZ]]
 752 // CHECK:       [[LOCAL_AA:%.+]] = alloca i[[SZ]]
 753 // CHECK:       [[LOCAL_AAA:%.+]] = alloca i[[SZ]]
 754 // CHECK:       [[LOCAL_B:%.+]] = alloca ptr
 755 // CHECK-DAG:   store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
 756 // CHECK-DAG:   store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]
 757 // CHECK-DAG:   store i[[SZ]] [[ARG_AAA:%.+]], ptr [[LOCAL_AAA]]
 758 // CHECK-DAG:   store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
 759 // Store captures in the context.
 760 // CHECK-DAG:      [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
 761 // Use captures.
 762 // CHECK-64-DAG:   load i32, ptr [[LOCAL_A]]
 763 // CHECK-DAG:      load i16, ptr [[LOCAL_AA]]
 764 // CHECK-DAG:      load i8, ptr [[LOCAL_AAA]]
 765 // CHECK-32-DAG:   load i32, ptr [[LOCAL_A]]
 766 // CHECK-DAG:      getelementptr inbounds [10 x i32], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
 767
 768 // CHECK:       define internal void [[HVT5]]
 769 // Create local storage for each capture.
 770 // CHECK:       [[LOCAL_A:%.+]] = alloca i[[SZ]]
 771 // CHECK:       [[LOCAL_AA:%.+]] = alloca i[[SZ]]
 772 // CHECK:       [[LOCAL_B:%.+]] = alloca ptr
 773 // CHECK-DAG:   store i[[SZ]] [[ARG_A:%.+]], ptr [[LOCAL_A]]
 774 // CHECK-DAG:   store i[[SZ]] [[ARG_AA:%.+]], ptr [[LOCAL_AA]]
 775 // CHECK-DAG:   store ptr [[ARG_B:%.+]], ptr [[LOCAL_B]]
 776 // Store captures in the context.
 777 // CHECK-DAG:   [[REF_B:%.+]] = load ptr, ptr [[LOCAL_B]],
 778 // Use captures.
 779 // CHECK-64-DAG:   load i32, ptr [[LOCAL_A]]
 780 // CHECK-32-DAG:   load i32, ptr [[LOCAL_A]]
 781 // CHECK-DAG:   load i16, ptr [[LOCAL_AA]]
 782 // CHECK-DAG:   getelementptr inbounds [10 x i32], ptr [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
 783
 784 // OMP50: define internal void @__omp_offloading_{{.+}}_{{.+}}bar{{.+}}_l{{[0-9]+}}(i[[SZ]] noundef %{{.+}})
 785
 786 // OMP50: define {{.*}}@{{.*}}zee{{.*}}
 787
 788 // OMP50:       [[LOCAL_THIS:%.+]] = alloca ptr
 789 // OMP50:       [[BP:%.+]] = alloca [1 x ptr]
 790 // OMP50:       [[P:%.+]] = alloca [1 x ptr]
 791 // OMP50:       [[LOCAL_THIS1:%.+]] = load ptr, ptr [[LOCAL_THIS]]
 792 // OMP50:       [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
 793 // OMP50:       [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], ptr [[LOCAL_THIS1]], i[[SZ]] 0
 794
 795 // OMP50-DAG:   [[BPADDR0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 796 // OMP50-DAG:   [[PADDR0:%.+]] =  getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 797 // OMP50-DAG:   store ptr [[ARR_IDX]], ptr [[BPADDR0]]
 798 // OMP50-DAG:   store ptr [[ARR_IDX2]], ptr [[PADDR0]]
 799
 800 // OMP50:       [[BPR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 801 // OMP50:       [[PR:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 802 // OMP50:       [[RET:%.+]] = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 {{.+}}, i32 {{.+}}, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 803 // OMP50-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET]], 0
 804 // OMP50-NEXT:  br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]]
 805 // OMP50:       [[FAIL]]
 806 // OMP50:       call void [[HVT0:@.+]](ptr [[LOCAL_THIS1]])
 807 // OMP50-NEXT:  br label %[[END]]
 808 // OMP50:       [[END]]
 809
 810 void bar () {
 811 #define pragma_target _Pragma("omp target")
 812 pragma_target
 813 {
 814   global = 0;
 815 #pragma omp parallel shared(global)
 816   global = 1;
 817 }
 818 }
 819
 820 class S2 {
 821   int a, b, c;
 822
 823 public:
 824   void zee() {
 825 #pragma omp critical
 826     #pragma omp target map(this[0])
 827       a++;
 828   }
 829 };
 830
 831 #ifdef _DOMP51
 832 void thread_limit_target(int TargetTL, int TeamsTL) {
 833
 834 #pragma omp target
 835 {}
 836 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1, i32 0,
 837
 838 #pragma omp target
 839 #pragma omp teams
 840 {}
 841 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 0,
 842
 843 #pragma omp target thread_limit(TargetTL)
 844 {}
 845 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
 846 // OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]
 847 // OMP51: load {{.*}} [[CEA]]
 848 // OMP51: [[CE:%.*]] = load {{.*}} [[CEA]]
 849 // OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
 850 // OMP51: call i32 [[OMP_TASK_ENTRY]]
 851
 852 #pragma omp target thread_limit(TargetTL)
 853 #pragma omp teams
 854 {}
 855 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
 856 // OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]
 857 // OMP51: load {{.*}} [[CEA]]
 858 // OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
 859 // OMP51: call i32 [[OMP_TASK_ENTRY]]
 860
 861 #pragma omp target
 862 #pragma omp teams thread_limit(TeamsTL)
 863 {}
 864 // OMP51: load {{.*}} %TeamsTL.addr
 865 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
 866 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[TeamsL]],
 867
 868 #pragma omp target thread_limit(TargetTL)
 869 #pragma omp teams thread_limit(TeamsTL)
 870 {}
 871 // OMP51: load {{.*}} %TeamsTL.addr
 872 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
 873 // OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
 874 // OMP51: call i32 [[OMP_TASK_ENTRY]]
 875
 876 }
 877 #endif
 878 // Check that the offloading functions are called after setting thread_limit in the task entry functions
 879
 880 // OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
 881 // OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
 882 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1,
 883
 884 // OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
 885 // OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
 886 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,
 887
 888 // OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
 889 // OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
 890 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,
 891
 892
 893 // CHECK:     define internal void @.omp_offloading.requires_reg()
 894 // CHECK:     call void @__tgt_register_requires(i64 1)
 895 // CHECK:     ret void
 896
 897 int main () {
 898   S2 bar;
 899   bar.zee();
 900 }
 901
 902 #endif