clang/test/OpenMP/target_update_depend_codegen.cpp

   1 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
   2 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
   3 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
   4 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
   5 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
   6 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
   7
   8 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
   9 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  10 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  11 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  12 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  13 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  14 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
  15
  16 // expected-no-diagnostics
  17 // CK1: [[ST:%.+]] = type { i32, ptr }
  18 // CK1: %struct.kmp_depend_info = type { i[[sz:64|32]],
  19 // CK1-SAME: i[[sz]], i8 }
  20 #ifndef HEADER
  21 #define HEADER
  22
  23 template <typename T>
  24 struct ST {
  25   T a;
  26   double *b;
  27 };
  28
  29 ST<int> gb;
  30 double gc[100];
  31
  32 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
  33 // CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  34
  35 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
  36 // CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
  37
  38 // CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  39
  40 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24]
  41 // CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
  42
  43 // CK1-LABEL: _Z3fooi
  44 void foo(int arg) {
  45   int la;
  46   float lb[arg];
  47
  48   // CK1: alloca [1 x %struct.kmp_depend_info],
  49   // CK1: alloca [3 x %struct.kmp_depend_info],
  50   // CK1: alloca [4 x %struct.kmp_depend_info],
  51   // CK1: alloca [5 x %struct.kmp_depend_info],
  52
  53   // Region 00
  54   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
  55   // CK1: store ptr @gc, ptr [[BP0]],
  56   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
  57   // CK1: store ptr @gc, ptr [[P0]],
  58   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
  59   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
  60   // CK1: [[CAP_DEVICE:%.+]] = getelementptr inbounds nuw %struct.anon, ptr [[CAPTURES:%.+]], i32 0, i32 0
  61   // CK1: [[DEVICE:%.+]] = load i32, ptr %{{.+}}
  62   // CK1: store i32 [[DEVICE]], ptr [[CAP_DEVICE]],
  63   // CK1: [[DEV1:%.+]] = load i32, ptr %{{.+}}
  64   // CK1: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64
  65   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 4, ptr [[TASK_ENTRY0:@.+]], i64 [[DEV2]])
  66   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates, ptr [[RES]], i32 0, i32 0
  67   // CK1: [[SHAREDS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t, ptr [[TASK_T]], i32 0, i32 0
  68   // CK1: [[SHAREDS_REF:%.+]] = load ptr, ptr [[SHAREDS]],
  69   // CK1: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 4 [[SHAREDS_REF]], ptr align 4 [[CAPTURES]], i[[sz]] 4, i1 false)
  70   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates, ptr [[RES]], i32 0, i32 1
  71   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 0
  72   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 8 [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
  73   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 1
  74   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 8 [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
  75   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 2
  76   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 8 [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE00]], i[[sz]] 8, i1 false)
  77   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 0
  78   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 4 [[PRIVS_SIZES]], ptr align 4 [[SIZE00]], i[[sz]] 8, i1 false)
  79   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 1
  80   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
  81   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 2
  82   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
  83   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
  84   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
  85   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
  86   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
  87   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
  88   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
  89   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
  90   // CK1: store i8 1, ptr [[DEP_ATTRS]]
  91   // CK1: = call i32 @__kmpc_omp_task_with_deps(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]], i32 1, ptr [[MAIN_DEP]], i32 0, ptr null)
  92
  93   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
  94   #pragma omp target update if(1+3-5) device(arg) from(gc) nowait depend(in: arg)
  95   {++arg;}
  96
  97   // Region 01
  98   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
  99   #pragma omp target update to(la) if(1+3-4) depend(in: la) depend(out: arg)
 100   {++arg;}
 101
 102   // Region 02
 103   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
 104   // CK1: [[IFTHEN]]
 105   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 106   // CK1: store ptr [[ARG:%.+]], ptr [[BP0]],
 107   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
 108   // CK1: store ptr [[ARG]], ptr [[P0]],
 109   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 110   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 111   // CK1: [[IF_DEVICE:%.+]] = getelementptr inbounds nuw %struct.anon{{.+}}, ptr [[CAPTURES:%.+]], i32 0, i32 0
 112   // CK1: [[IF:%.+]] = load i8, ptr %{{.+}}
 113   // CK1: [[IF_BOOL:%.+]] = trunc i8 [[IF]] to i1
 114   // CK1: [[IF:%.+]] = zext i1 [[IF_BOOL]] to i8
 115   // CK1: store i8 [[IF]], ptr [[IF_DEVICE]],
 116   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 1, ptr [[TASK_ENTRY2:@.+]])
 117   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
 118   // CK1: [[SHAREDS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t, ptr [[TASK_T]], i32 0, i32 0
 119   // CK1: [[SHAREDS_REF:%.+]] = load ptr, ptr [[SHAREDS]],
 120   // CK1: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 1 [[SHAREDS_REF]], ptr align 1 [[CAPTURES]], i[[sz]] 1, i1 false)
 121   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
 122   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 123   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
 124   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 125   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
 126   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 127   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE02]], i[[sz]] {{8|4}}, i1 false)
 128   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 129   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE02]], i[[sz]] {{8|4}}, i1 false)
 130   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 131   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
 132   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 133   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
 134   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 135   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
 136   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 137   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 138   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 139   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 140   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 141   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 142   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 143   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
 144   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 145   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 146   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 147   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 148   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 149   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 150   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
 151   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 152   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
 153   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 154   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
 155   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 156   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 157   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 3, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
 158   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 159   // CK1: = call i32 [[TASK_ENTRY2]](i32 %{{.+}}, ptr [[RES]])
 160   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 161
 162   // CK1: br label %[[IFEND:[^,]+]]
 163
 164   // CK1: [[IFELSE]]
 165   // CK1: br label %[[IFEND]]
 166   // CK1: [[IFEND]]
 167   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 168   #pragma omp target update to(arg) if(arg) device(4) depend(inout: arg, la, gc)
 169   {++arg;}
 170
 171   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 172   {++arg;}
 173
 174   // Region 03
 175   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 176   // CK1: store ptr [[VLA:%.+]], ptr [[BP0]],
 177   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
 178   // CK1: store ptr [[VLA]], ptr [[P0]],
 179   // CK1: [[S0:%.+]] = getelementptr inbounds [1 x i64], ptr [[S:%.+]], i32 0, i32 0
 180   // CK1: store i64 {{.+}}, ptr [[S0]],
 181   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 182   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 183   // CK1: [[GEPS0:%.+]] = getelementptr inbounds [1 x i64], ptr [[S]], i32 0, i32 0
 184   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 1, ptr [[TASK_ENTRY3:@.+]])
 185   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
 186   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
 187   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 188   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
 189   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 190   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
 191   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 192   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{8|4}}, i1 false)
 193   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 194   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{8|4}}, i1 false)
 195   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 196   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
 197   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 198   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
 199   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 200   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
 201   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 202   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 203   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 204   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
 205   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 206   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 207   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 208   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
 209   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 210   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 211   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 212   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 213   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 214   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 215   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 216   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
 217   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 218   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 219   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 220   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 221   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 222   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 223   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 3
 224   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 225   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
 226   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 227   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
 228   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 229   // CK1: store i8 3, ptr [[DEP_ATTRS]]
 230   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 4, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
 231   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 232   // CK1: = call i32 [[TASK_ENTRY3]](i32 %{{.+}}, ptr [[RES]])
 233   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 234   #pragma omp target update from(lb) depend(out: lb, arg, la, gc)
 235   {++arg;}
 236
 237   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 238   {++arg;}
 239
 240   // Region 04
 241   // CK1: [[DIV:%.+]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr (ptr, ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1) to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 242   // CK1: [[BP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 243   // CK1: store ptr @gb, ptr [[BP0]],
 244   // CK1: [[P0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P:%.+]], i32 0, i32 0
 245   // CK1: store ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), ptr [[P0]],
 246   // CK1: [[PS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS:%.+]], i32 0, i32 0
 247   // CK1: store i64 [[DIV]], ptr [[PS0]],
 248   // CK1: [[BP1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1
 249   // CK1: store ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), ptr [[BP1]],
 250   // CK1: [[P1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
 251   // CK1: store ptr %{{.+}}, ptr [[P1]],
 252   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0
 253   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0
 254   // CK1: [[GEPS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS]], i32 0, i32 0
 255   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{88|52}}, i[[sz]] 1, ptr [[TASK_ENTRY4:@.+]])
 256   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
 257   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
 258   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 259   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{16|8}}, i1 false)
 260   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 261   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{16|8}}, i1 false)
 262   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 263   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{16|8}}, i1 false)
 264   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
 265   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{16|8}}, i1 false)
 266   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
 267   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{16|8}}, i1 false)
 268   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
 269   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{16|8}}, i1 false)
 270   // CK1: %{{.+}} = sub nuw
 271   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 272   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
 273   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 274   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 275   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 276   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
 277   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 278   // CK1: store i8 1, ptr [[DEP_ATTRS]]
 279   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 280   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
 281   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 282   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 283   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 284   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 285   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 286   // CK1: store i8 1, ptr [[DEP_ATTRS]]
 287   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 288   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
 289   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 290   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 291   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 292   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
 293   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 294   // CK1: store i8 1, ptr [[DEP_ATTRS]]
 295   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 3
 296   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 297   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
 298   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 299   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
 300   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 301   // CK1: store i8 1, ptr [[DEP_ATTRS]]
 302   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
 303   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 4
 304   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
 305   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
 306   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
 307   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
 308   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
 309   // CK1: store i8 1, ptr [[DEP_ATTRS]]
 310   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 5, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
 311   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 312   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, ptr [[RES]])
 313   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
 314   #pragma omp target update to(gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
 315   {++arg;}
 316 }
 317
 318 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, ptr noalias noundef %1)
 319 // CK1-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE00]]{{.+}}, ptr null)
 320 // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
 321 // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
 322 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
 323 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
 324 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
 325 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
 326 // CK1: ret i32 0
 327 // CK1: }
 328
 329 // CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, ptr noalias noundef %1)
 330 // CK1-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 4, i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE02]]{{.+}}, ptr null)
 331 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
 332 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
 333 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
 334 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
 335 // CK1: ret i32 0
 336 // CK1: }
 337
 338 // CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, ptr noalias noundef %1)
 339 // CK1-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE03]]{{.+}}, ptr null)
 340 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
 341 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
 342 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
 343 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
 344 // CK1-NOT: __tgt_target_data_end
 345 // CK1: ret i32 0
 346 // CK1: }
 347
 348 // CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, ptr noalias noundef %1)
 349 // CK1-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE04]]{{.+}}, ptr null)
 350 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
 351 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
 352 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
 353 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
 354 // CK1-NOT: __tgt_target_data_end
 355 // CK1: ret i32 0
 356 // CK1: }
 357
 358 #endif