clang/test/OpenMP/target_map_codegen_25.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-64
   7 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-64
   9 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-32
  10 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-32
  12
  13 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-64
  14 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-64
  16 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-32
  17 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-32
  19
  20 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-64
  21 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-64
  23 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 CK26 --check-prefix CK26-32
  24 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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 CK26 --check-prefix CK26-32
  26
  27 // RUN: %clang_cc1 -DCK26 -std=c++11 -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-ONLY25 %s
  28 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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-ONLY25 %s
  30 // RUN: %clang_cc1 -DCK26 -std=c++11 -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-ONLY25 %s
  31 // RUN: %clang_cc1 -DCK26 -std=c++11 -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 -std=c++11 -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-ONLY25 %s
  33 // SIMD-ONLY25-NOT: {{__kmpc|__tgt}}
  34 #ifdef CK26
  35 // CK26: [[ST:%.+]] = type { i32, ptr, i32, ptr }
  36
  37 // CK26-LABEL: @.__omp_offloading_{{.*}}CC{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  38 // CK26: [[SIZE00:@.+]] = private {{.*}}constant [2 x i64] [i64 {{32|16}}, i64 4]
  39 // CK26: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 547, i64 35]
  40
  41 // CK26-LABEL: @.__omp_offloading_{{.*}}CC{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  42 // CK26: [[SIZE01:@.+]] = private {{.*}}constant [2 x i64] [i64 {{32|16}}, i64 4]
  43 // CK26: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i64] [i64 547, i64 35]
  44
  45 // CK26-LABEL: @.__omp_offloading_{{.*}}CC{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  46 // CK26: [[SIZE02:@.+]] = private {{.*}}constant [2 x i64] [i64 {{32|16}}, i64 4]
  47 // CK26: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 547, i64 35]
  48
  49 // CK26-LABEL: @.__omp_offloading_{{.*}}CC{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
  50 // CK26: [[SIZE03:@.+]] = private {{.*}}constant [2 x i64] [i64 {{32|16}}, i64 4]
  51 // CK26: [[MTYPE03:@.+]] = private {{.*}}constant [2 x i64] [i64 547, i64 35]
  52
  53 // CK26-LABEL: explicit_maps_with_private_class_members{{.*}}(
  54
  55 struct CC {
  56   int fA;
  57   float &fB;
  58   int pA;
  59   float &pB;
  60
  61   CC(float &B) : fB(B), pB(B) {
  62
  63     // CK26: call {{.*}}@__kmpc_fork_call{{.*}} [[OUTCALL:@.+]]
  64     // define {{.*}}void [[OUTCALL]]
  65     #pragma omp parallel firstprivate(fA,fB) private(pA,pB)
  66     {
  67 // Region 00
  68 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  69 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  70 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  71 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  72 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  73 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  74 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  75
  76 // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  77 // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  78 // CK26-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
  79 // CK26-DAG: store ptr [[VAR0]], ptr [[P0]]
  80
  81 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
  82 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
  83 // CK26-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
  84 // CK26-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
  85 // CK26-DAG: [[VAR1]] = load ptr, ptr [[PVT:%.+]],
  86 // CK26-DAG: [[SEC1]] = load ptr, ptr [[PVT]],
  87
  88 // CK26: call void [[CALL00:@.+]](ptr {{[^,]+}}, ptr {{[^,]+}})
  89 #pragma omp target map(fA)
  90       {
  91         ++fA;
  92       }
  93
  94 // Region 01
  95 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  96 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  97 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  98 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  99 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 100 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 101 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 102
 103 // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 104 // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 105 // CK26-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 106 // CK26-DAG: store ptr [[VAR0]], ptr [[P0]]
 107
 108 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 109 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 110 // CK26-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
 111 // CK26-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 112 // CK26-DAG: [[VAR1]] = load ptr, ptr [[PVT:%.+]],
 113 // CK26-DAG: [[SEC1]] = load ptr, ptr [[PVT]],
 114
 115 // CK26: call void [[CALL01:@.+]](ptr {{[^,]+}}, ptr {{[^,]+}})
 116 #pragma omp target map(fB)
 117       {
 118         fB += 1.0;
 119       }
 120
 121 // Region 02
 122 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 123 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 124 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 125 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 126 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 127 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 128 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 129
 130 // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 131 // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 132 // CK26-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 133 // CK26-DAG: store ptr [[VAR0]], ptr [[P0]]
 134
 135 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 136 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 137 // CK26-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
 138 // CK26-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 139 // CK26-DAG: [[VAR1]] = load ptr, ptr [[PVT:%.+]],
 140 // CK26-DAG: [[SEC1]] = load ptr, ptr [[PVT]],
 141
 142 // CK26: call void [[CALL02:@.+]](ptr {{[^,]+}}, ptr {{[^,]+}})
 143 #pragma omp target map(pA)
 144       {
 145         ++pA;
 146       }
 147
 148 // Region 01
 149 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 150 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 151 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 152 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 153 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 154 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 155 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 156
 157 // CK26-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 158 // CK26-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 159 // CK26-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
 160 // CK26-DAG: store ptr [[VAR0]], ptr [[P0]]
 161
 162 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 163 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 164 // CK26-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
 165 // CK26-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
 166 // CK26-DAG: [[VAR1]] = load ptr, ptr [[PVT:%.+]],
 167 // CK26-DAG: [[SEC1]] = load ptr, ptr [[PVT]],
 168
 169 // CK26: call void [[CALL03:@.+]](ptr {{[^,]+}}, ptr {{[^,]+}})
 170 #pragma omp target map(pB)
 171       {
 172         pB += 1.0;
 173       }
 174     }
 175   }
 176
 177   int foo() {
 178     return fA + pA;
 179   }
 180 };
 181
 182 // Make sure the private instance is used in all target regions.
 183 // CK26: define {{.+}}[[CALL00]]({{.*}}ptr{{.*}}[[PVTARG:%.+]])
 184 // CK26: store ptr [[PVTARG]], ptr [[PVTADDR:%.+]],
 185 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 186 // CK26: store ptr [[ADDR]], ptr [[PVTADDR:%.+]],
 187 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 188 // CK26: [[VAL:%.+]] = load i32, ptr [[ADDR]],
 189 // CK26: add nsw i32 [[VAL]], 1
 190
 191 // CK26: define {{.+}}[[CALL01]]({{.*}}ptr{{.*}}[[PVTARG:%.+]])
 192 // CK26: store ptr [[PVTARG]], ptr [[PVTADDR:%.+]],
 193 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 194 // CK26: store ptr [[ADDR]], ptr [[PVTADDR:%.+]],
 195 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 196 // CK26: [[VAL:%.+]] = load float, ptr [[ADDR]],
 197 // CK26: [[EXT:%.+]] = fpext float [[VAL]] to double
 198 // CK26: fadd double [[EXT]], 1.000000e+00
 199
 200 // CK26: define {{.+}}[[CALL02]]({{.*}}ptr{{.*}}[[PVTARG:%.+]])
 201 // CK26: store ptr [[PVTARG]], ptr [[PVTADDR:%.+]],
 202 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 203 // CK26: store ptr [[ADDR]], ptr [[PVTADDR:%.+]],
 204 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 205 // CK26: [[VAL:%.+]] = load i32, ptr [[ADDR]],
 206 // CK26: add nsw i32 [[VAL]], 1
 207
 208 // CK26: define {{.+}}[[CALL03]]({{.*}}ptr{{.*}}[[PVTARG:%.+]])
 209 // CK26: store ptr [[PVTARG]], ptr [[PVTADDR:%.+]],
 210 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 211 // CK26: store ptr [[ADDR]], ptr [[PVTADDR:%.+]],
 212 // CK26: [[ADDR:%.+]] = load ptr, ptr [[PVTADDR]],
 213 // CK26: [[VAL:%.+]] = load float, ptr [[ADDR]],
 214 // CK26: [[EXT:%.+]] = fpext float [[VAL]] to double
 215 // CK26: fadd double [[EXT]], 1.000000e+00
 216
 217 int explicit_maps_with_private_class_members(){
 218   float B;
 219   CC c(B);
 220   return c.foo();
 221 }
 222 #endif // CK26
 223 #endif