1 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -debug-info-kind=limited -emit-llvm %s -o - | FileCheck %s --check-prefix DEBUG 2 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK 3 #ifndef HEADER 4 #define HEADER 5 6 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";d;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 7 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 8 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i[1:23];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 9 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 10 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p[1:24];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 11 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 12 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 13 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 14 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 15 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 16 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 17 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 18 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 19 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 20 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 21 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 22 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 23 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 24 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 25 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 26 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 27 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 28 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 29 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 30 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 31 32 struct S1 { 33 int i; 34 float f[50]; 35 }; 36 37 struct S2 { 38 int i; 39 float f[50]; 40 S1 s; 41 double *p; 42 struct S2 *ps; 43 }; 44 45 void foo() { 46 double d; 47 int i[100]; 48 float *p; 49 50 S2 s; 51 S2 *ps; 52 53 #pragma omp target map(d) 54 { } 55 #pragma omp target map(i) 56 { } 57 #pragma omp target map(i[1:23]) 58 { } 59 #pragma omp target map(p) 60 { } 61 #pragma omp target map(p[1:24]) 62 { } 63 #pragma omp target map(s) 64 { } 65 #pragma omp target map(s.i) 66 { } 67 #pragma omp target map(s.s.f) 68 { } 69 #pragma omp target map(s.p) 70 { } 71 #pragma omp target map(to: s.p[:22]) 72 { } 73 #pragma omp target map(s.ps) 74 { } 75 #pragma omp target map(from: s.ps->s.i) 76 { } 77 #pragma omp target map(to: s.ps->ps) 78 { } 79 #pragma omp target map(s.ps->ps->ps) 80 { } 81 #pragma omp target map(to: s.ps->ps->s.f[:22]) 82 { } 83 #pragma omp target map(ps) 84 { } 85 #pragma omp target map(ps->i) 86 { } 87 #pragma omp target map(ps->s.f) 88 { } 89 #pragma omp target map(from: ps->p) 90 { } 91 #pragma omp target map(to: ps->p[:22]) 92 { } 93 #pragma omp target map(ps->ps) 94 { } 95 #pragma omp target map(from: ps->ps->s.i) 96 { } 97 #pragma omp target map(from: ps->ps->ps) 98 { } 99 #pragma omp target map(ps->ps->ps->ps) 100 { } 101 #pragma omp target map(to: ps->ps->ps->s.f[:22]) 102 { } 103 #pragma omp target map(to: s.f[:22]) map(from: s.p[:33]) 104 { } 105 #pragma omp target map(from: s.f[:22]) map(to: ps->p[:33]) 106 { } 107 #pragma omp target map(from: s.f[:22], s.s) map(to: ps->p[:33]) 108 { } 109 } 110 111 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";B;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 112 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";unknown;unknown;0;0;;\00" 113 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";A;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 114 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 115 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";fn;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 116 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 117 // DEBUG: @{{.+}} = private constant [7 x ptr] [ptr @{{[0-9]+}}, ptr @{{[0-9]+}}, ptr @{{[0-9]+}}, ptr @{{[0-9]+}}, ptr @{{[0-9]+}}, ptr @{{[0-9]+}}, ptr @{{[0-9]+}}] 118 119 void bar(int N) { 120 double B[10]; 121 double A[N]; 122 double x; 123 S1 s; 124 auto fn = [&x]() { return x; }; 125 #pragma omp target 126 { 127 (void)B; 128 (void)A; 129 (void)fn(); 130 (void)s.f; 131 } 132 } 133 134 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";t;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 135 136 #pragma omp declare target 137 double t; 138 #pragma omp end declare target 139 140 void baz() { 141 #pragma omp target map(to:t) 142 { } 143 #pragma omp target map(to:t) nowait 144 { } 145 #pragma omp target teams map(to:t) 146 { } 147 #pragma omp target teams map(to:t) nowait 148 { } 149 #pragma omp target data map(to:t) 150 { } 151 #pragma omp target enter data map(to:t) 152 153 #pragma omp target enter data map(to:t) nowait 154 155 #pragma omp target exit data map(from:t) 156 157 #pragma omp target exit data map(from:t) nowait 158 159 #pragma omp target update from(t) 160 161 #pragma omp target update to(t) 162 163 #pragma omp target update from(t) nowait 164 165 #pragma omp target update to(t) nowait 166 } 167 168 struct S3 { 169 S3() { 170 #pragma omp target data map(alloc : Z[0:64]) 171 { } 172 } 173 double Z[64]; 174 }; 175 176 #pragma omp declare mapper(id: S3 s) map(s.Z[0:64]) 177 178 void qux() { 179 S3 s; 180 #pragma omp target map(mapper(id), to:s) 181 { } 182 } 183 184 185 186 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 187 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";this->Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00" 188 189 // Clang used to mistakenly generate the map name "x" for both x and y on this 190 // directive. Conditions to reproduce the bug: a single map clause has two 191 // variables, and at least the second is used in the associated statement. 192 // 193 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;[[@LINE+3]];7;;\00" 194 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";y;{{.*}}.cpp;[[@LINE+2]];10;;\00" 195 void secondMapNameInClause() { 196 int x, y; 197 #pragma omp target map(to: x, y) 198 x = y = 1; 199 } 200 201 // DEBUG: store ptr @[[NAME:.offload_mapnames.[0-9]+]], ptr %[[ARG:.+]] 202 // CHECK-NOT: store ptr @[[NAME:.offload_mapnames.[0-9]+]], ptr %[[ARG:.+]] 203 204 // DEBUG: void @.omp_mapper._ZTS2S3.id(ptr {{.*}}, ptr {{.*}}, ptr {{.*}}, i64 {{.*}}, i64 {{.*}}, ptr noundef [[MAPPER_NAME:%.+]]) 205 // DEBUG: call void @__tgt_push_mapper_component(ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}}, i64 %{{.*}}, i64 %{{.*}}, ptr [[MAPPER_NAME]]) 206 207 #endif 208