clang/test/OpenMP/target_map_codegen_26.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK27 --check-prefix CK27-64
   7 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
   8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-64
   9 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  10 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  12
  13 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK27 --check-prefix CK27-64
  14 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  15 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-64
  16 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  17 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  18 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  19
  20 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefix CK27 --check-prefix CK27-64
  21 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  22 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-64
  23 // RUN: %clang_cc1 -DCK27 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  24 // RUN: %clang_cc1 -DCK27 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  25 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefix CK27 --check-prefix CK27-32
  26
  27 // RUN: %clang_cc1 -DCK27 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
  28 // RUN: %clang_cc1 -DCK27 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
  29 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
  30 // RUN: %clang_cc1 -DCK27 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
  31 // RUN: %clang_cc1 -DCK27 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
  32 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY26 %s
  33 // SIMD-ONLY26-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK27
  35
  36 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  37 // CK27: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
  38 // CK27: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 544]
  39
  40 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  41 // CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
  42 // CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  43
  44 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  45 // CK27: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
  46 // CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  47
  48 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  49 // CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
  50 // CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  51
  52 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  53 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  54 // CK27: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
  55 // CK27: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
  56
  57 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  58 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  59 // CK27: [[SIZE07:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  60 // CK27: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 288]
  61
  62 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  63 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  64 // CK27: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 40]
  65 // CK27: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 161]
  66
  67 // CK27-LABEL: zero_size_section_and_private_maps{{.*}}(
  68 void zero_size_section_and_private_maps (int ii){
  69
  70   // Map of a pointer.
  71   int *pa;
  72
  73 // Region 00
  74 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  75 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  76 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  77 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  78 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  79 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  80 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  81
  82 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  83 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  84 // CK27-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
  85 // CK27-DAG: store ptr [[VAR0]], ptr [[P0]]
  86
  87 // CK27: call void [[CALL00:@.+]](ptr {{[^,]+}})
  88 #pragma omp target
  89   {
  90     pa[50]++;
  91   }
  92
  93 // Region 01
  94 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  95 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  96 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  97 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  98 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  99 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 100 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 101
 102 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 103 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 104 // CK27-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
 105 // CK27-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 106 // CK27-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
 107 // CK27-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0
 108 // CK27-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
 109
 110 // CK27: call void [[CALL01:@.+]](ptr {{[^,]+}})
 111 #pragma omp target map(pa[:0])
 112   {
 113     pa[50]++;
 114   }
 115
 116 // Region 02
 117 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 118 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 119 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 120 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 121 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 122 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 123 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 124
 125 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 126 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 127 // CK27-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
 128 // CK27-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 129 // CK27-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
 130 // CK27-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0
 131 // CK27-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
 132
 133 // CK27: call void [[CALL02:@.+]](ptr {{[^,]+}})
 134 #pragma omp target map(pa [0:0])
 135   {
 136     pa[50]++;
 137   }
 138
 139 // Region 03
 140 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 141 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 142 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 143 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 144 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 145 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 146 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 147
 148 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 149 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 150 // CK27-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
 151 // CK27-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 152 // CK27-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
 153 // CK27-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} %{{.+}}
 154 // CK27-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
 155
 156 // CK27: call void [[CALL03:@.+]](ptr {{[^,]+}})
 157 #pragma omp target map(pa [ii:0])
 158   {
 159     pa[50]++;
 160   }
 161
 162   int *pvtPtr;
 163   int pvtScl;
 164   int pvtArr[10];
 165
 166 // Region 04
 167 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 168 // CK27: call void [[CALL04:@.+]]()
 169 #pragma omp target private(pvtPtr)
 170   {
 171     pvtPtr[5]++;
 172   }
 173
 174 // Region 05
 175 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 176 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 177 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 178 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 179 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 180 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 181 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 182
 183 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 184 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 185 // CK27-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 186 // CK27-DAG: store ptr [[VAR0]], ptr [[P0]]
 187
 188 // CK27: call void [[CALL05:@.+]](ptr {{[^,]+}})
 189 #pragma omp target firstprivate(pvtPtr)
 190   {
 191     pvtPtr[5]++;
 192   }
 193
 194 // Region 06
 195 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 196 // CK27: call void [[CALL06:@.+]]()
 197 #pragma omp target private(pvtScl)
 198   {
 199     pvtScl++;
 200   }
 201
 202 // Region 07
 203 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 204 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 205 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 206 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 207 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 208 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 209 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 210 // CK27-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 211 // CK27-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 212 // CK27-DAG: store i[[Z:64|32]] [[VAL:%.+]], ptr [[BP1]]
 213 // CK27-DAG: store i[[Z]] [[VAL]], ptr [[P1]]
 214 // CK27-DAG: [[VAL]] = load i[[Z]], ptr [[ADDR:%.+]],
 215 // CK27-64-DAG: store i32 {{.+}}, ptr [[ADDR]],
 216
 217 // CK27: call void [[CALL07:@.+]](i[[Z]] [[VAL]])
 218 #pragma omp target firstprivate(pvtScl)
 219   {
 220     pvtScl++;
 221   }
 222
 223 // Region 08
 224 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 225 // CK27: call void [[CALL08:@.+]]()
 226 #pragma omp target private(pvtArr)
 227   {
 228     pvtArr[5]++;
 229   }
 230
 231 // Region 09
 232 // CK27-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 233 // CK27-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 234 // CK27-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 235 // CK27-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 236 // CK27-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 237 // CK27-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 238 // CK27-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 239
 240 // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 241 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 242 // CK27-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 243 // CK27-DAG: store ptr [[VAR0]], ptr [[P0]]
 244
 245 // CK27: call void [[CALL09:@.+]](ptr {{[^,]+}})
 246 #pragma omp target firstprivate(pvtArr)
 247   {
 248     pvtArr[5]++;
 249   }
 250 }
 251
 252 // CK27: define {{.+}}[[CALL00]]
 253 // CK27: define {{.+}}[[CALL01]]
 254 // CK27: define {{.+}}[[CALL02]]
 255 // CK27: define {{.+}}[[CALL03]]
 256 // CK27: define {{.+}}[[CALL04]]
 257 // CK27: define {{.+}}[[CALL05]]
 258 // CK27: define {{.+}}[[CALL06]]
 259 // CK27: define {{.+}}[[CALL07]]
 260 #endif // CK27
 261 #endif