clang/test/OpenMP/target_map_names.cpp

   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 [[NAME_ARG:%.+]])
 205 // DEBUG: store ptr [[NAME_ARG]], ptr [[NAME_STACK:%.+]]
 206 // DEBUG: [[MAPPER_NAME:%.+]] = load ptr, ptr [[NAME_STACK]]
 207 // DEBUG: call void @__tgt_push_mapper_component(ptr %{{.*}}, ptr %{{.*}}, ptr %{{.*}}, i64 %{{.*}}, i64 %{{.*}}, ptr [[MAPPER_NAME]])
 208
 209 #endif