clang/test/OpenMP/target_map_codegen_22.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-64
   7 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-64
   9 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-32
  10 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-32
  12
  13 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-64
  14 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-64
  16 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-32
  17 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-32
  19
  20 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-64
  21 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-64
  23 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 CK23 --check-prefix CK23-32
  24 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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 CK23 --check-prefix CK23-32
  26
  27 // RUN: %clang_cc1 -std=c++11 -DCK23 -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-ONLY22 %s
  28 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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-ONLY22 %s
  30 // RUN: %clang_cc1 -std=c++11 -DCK23 -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-ONLY22 %s
  31 // RUN: %clang_cc1 -std=c++11 -DCK23 -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 -std=c++11 -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-ONLY22 %s
  33 // SIMD-ONLY22-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK23
  35
  36 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  37 // CK23: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  38 // CK23: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  39
  40 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  41 // CK23: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
  42 // CK23: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  43
  44 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  45 // CK23: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
  46 // CK23: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  47
  48 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  49 // CK23: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
  50 // CK23: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  51
  52 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  53 // CK23: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
  54 // CK23: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  55
  56 // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  57 // CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
  58 // CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
  59
  60 // CK23-LABEL: explicit_maps_inside_captured{{.*}}(
  61 int explicit_maps_inside_captured(int a){
  62   float b;
  63   float c[100];
  64   float *d;
  65
  66   // CK23: call void @{{.*}}explicit_maps_inside_captured{{.*}}(ptr {{.*}})
  67   // CK23: define {{.*}}explicit_maps_inside_captured{{.*}}
  68   [&](void){
  69 // Region 00
  70 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  71 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  72 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  73 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  74 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  75 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  76 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  77
  78 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  79 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  80 // CK23-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
  81 // CK23-DAG: store ptr [[VAR00:%.+]], ptr [[P0]]
  82 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
  83 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
  84 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
  85 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
  86
  87 // CK23: call void [[CALL00:@.+]](ptr {{[^,]+}})
  88 #pragma omp target map(a)
  89       { a+=1; }
  90 // Region 01
  91 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  92 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  93 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  94 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  95 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  96 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  97 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  98
  99 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 100 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 101 // CK23-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 102 // CK23-DAG: store ptr [[VAR00:%.+]], ptr [[P0]]
 103 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
 104 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
 105 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
 106 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
 107
 108 // CK23: call void [[CALL01:@.+]](ptr {{[^,]+}})
 109 #pragma omp target map(b)
 110       { b+=1; }
 111 // Region 02
 112 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 113 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 114 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 115 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 116 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 117 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 118 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 119
 120 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 121 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 122 // CK23-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 123 // CK23-DAG: store ptr [[VAR00:%.+]], ptr [[P0]]
 124 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
 125 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
 126 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
 127 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
 128
 129 // CK23: call void [[CALL02:@.+]](ptr {{[^,]+}})
 130 #pragma omp target map(c)
 131       { c[3]+=1; }
 132
 133 // Region 03
 134 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 135 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 136 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 137 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 138 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 139 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 140 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 141
 142 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 143 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 144 // CK23-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 145 // CK23-DAG: store ptr [[VAR00:%.+]], ptr [[P0]]
 146 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
 147 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
 148 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
 149 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
 150
 151 // CK23: call void [[CALL03:@.+]](ptr {{[^,]+}})
 152 #pragma omp target map(d)
 153       { d[3]+=1; }
 154 // Region 04
 155 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 156 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 157 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 158 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 159 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 160 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 161 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 162
 163 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 164 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 165 // CK23-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 166 // CK23-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 167 // CK23-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
 168 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
 169 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
 170 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
 171 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
 172
 173 // CK23: call void [[CALL04:@.+]](ptr {{[^,]+}})
 174 #pragma omp target map(c [2:4])
 175       { c[3]+=1; }
 176
 177 // Region 05
 178 // CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 179 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 180 // CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 181 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 182 // CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 183 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 184 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 185
 186 // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 187 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 188 // CK23-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
 189 // CK23-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
 190 // CK23-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
 191 // CK23-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 2
 192 // CK23-DAG: [[RVAR00]] = load ptr, ptr [[VAR00:%[^,]+]]
 193 // CK23-DAG: [[VAR0]] = load ptr, ptr [[CAP0:%[^,]+]]
 194 // CK23-DAG: [[CAP0]] = getelementptr inbounds %class.anon,
 195 // CK23-DAG: [[VAR00]] = load ptr, ptr [[CAP00:%[^,]+]]
 196 // CK23-DAG: [[CAP00]] = getelementptr inbounds %class.anon,
 197
 198 // CK23: call void [[CALL05:@.+]](ptr {{[^,]+}})
 199 #pragma omp target map(d [2:4])
 200       { d[3]+=1; }
 201   }();
 202   return b;
 203 }
 204
 205 // CK23: define {{.+}}[[CALL00]]
 206 // CK23: define {{.+}}[[CALL01]]
 207 // CK23: define {{.+}}[[CALL02]]
 208 // CK23: define {{.+}}[[CALL03]]
 209 // CK23: define {{.+}}[[CALL04]]
 210 // CK23: define {{.+}}[[CALL05]]
 211 #endif // CK23
 212 #endif