clang/test/OpenMP/target_map_codegen_34.cpp

   1 // expected-no-diagnostics
   2 #ifndef HEADER
   3 #define HEADER
   4
   5 ///==========================================================================///
   6 // RUN: %clang_cc1 -DCK34 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK34 --check-prefix CK34-64
   7 // RUN: %clang_cc1 -DCK34 -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 %s  --check-prefix CK34 --check-prefix CK34-64
   9 // RUN: %clang_cc1 -DCK34 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK34 --check-prefix CK34-32
  10 // RUN: %clang_cc1 -DCK34 -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 %s  --check-prefix CK34 --check-prefix CK34-32
  12
  13 // RUN: %clang_cc1 -DCK34 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
  14 // RUN: %clang_cc1 -DCK34 -fopenmp-simd -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-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 --check-prefix SIMD-ONLY32 %s
  16 // RUN: %clang_cc1 -DCK34 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s
  17 // RUN: %clang_cc1 -DCK34 -fopenmp-simd -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-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 --check-prefix SIMD-ONLY32 %s
  19 // SIMD-ONLY32-NOT: {{__kmpc|__tgt}}
  20 #ifdef CK34
  21
  22 class C {
  23 public:
  24   int a;
  25   double *b;
  26 };
  27
  28 #pragma omp declare mapper(C s) map(s.a, s.b[0:2])
  29
  30 class S {
  31   int a;
  32   C c;
  33   int b;
  34 public:
  35   void foo();
  36 };
  37
  38 // CK34-DAG: [[SIZE_TO:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 0, i64 {{16|8}}]
  39 // TARGET_PARAM = 0x20
  40 // MEMBER_OF_1 | TO = 0x1000000000001
  41 // MEMBER_OF_1 | IMPLICIT | TO = 0x1000000000201
  42 // CK34-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000001]], i64 [[#0x1000000000001]], i64 [[#0x1000000000201]]]
  43 // CK34-DAG: [[SIZE_FROM:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 0, i64 0, i64 {{16|8}}]
  44 // TARGET_PARAM = 0x20
  45 // MEMBER_OF_1 | FROM = 0x1000000000002
  46 // MEMBER_OF_1 | IMPLICIT | FROM = 0x1000000000202
  47 // CK34-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [4 x i64] [i64 [[#0x20]], i64 [[#0x1000000000002]], i64 [[#0x1000000000002]], i64 [[#0x1000000000202]]]
  48
  49 void default_mapper() {
  50   S s;
  51
  52   // CK34-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
  53   // CK34-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
  54   // CK34-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
  55   // CK34-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
  56   // CK34-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
  57   // CK34-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
  58   // CK34-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
  59   // CK34-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
  60   // CK34-DAG: store ptr [[MF:%.+]], ptr [[MARG]]
  61   // CK34-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
  62   // CK34-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
  63   // CK34-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
  64
  65   // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
  66
  67   // CK34-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
  68   // CK34-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
  69   // CK34-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
  70   // CK34-DAG: [[MF0:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 0
  71
  72   // CK34-DAG: store ptr [[S_ADDR:%.+]], ptr [[BP0]],
  73   // CK34-DAG: store ptr [[S_ADDR]], ptr [[P0]],
  74   // CK34-DAG: store i64 [[S_SIZE:%.+]], ptr [[S0]],
  75   // CK34-DAG: store ptr null, ptr [[MF0]],
  76
  77   // CK34-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
  78   // CK34-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
  79   // CK34-DAG: [[S_1_INTPTR]] = ptrtoint ptr [[S_1:%.+]] to i64
  80   // CK34-DAG: [[S_INTPTR]] = ptrtoint ptr [[S_ADDR]] to i64
  81   // CK34-DAG: [[S_1]] = getelementptr %class.S, ptr [[S_ADDR]], i32 1
  82
  83   // pass MEMBER_OF_1 | TO {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
  84
  85   // CK34-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
  86   // CK34-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
  87   // CK34-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
  88   // CK34-DAG: [[MF1:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 1
  89
  90   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP1]],
  91   // CK34-DAG: store ptr [[S_ADDR]], ptr [[P1]],
  92   // CK34-DAG: store i64 [[A_SIZE:%.+]], ptr [[S1]],
  93   // CK34-DAG: store ptr null, ptr [[MF1]],
  94
  95   // CK34-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
  96   // CK34-DAG: [[SZ]] = sub i64 [[C_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
  97   // CK34-DAG: [[S_INTPTR]] = ptrtoint ptr [[S_ADDR]] to i64
  98   // CK34-DAG: [[C_BEGIN_INTPTR]] = ptrtoint ptr [[C_ADDR:%.+]] to i64
  99   // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 2
 100   // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 1
 101
 102   // pass MEMBER_OF_1 | TO {&s, &s.c+1, ((void*)(&s)+31+1-(void*)(&s.c+1))} to copy the data of s.b.
 103
 104   // CK34-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
 105   // CK34-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 106   // CK34-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
 107   // CK34-DAG: [[MF2:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 2
 108
 109   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP2]],
 110   // CK34-DAG: store ptr [[C_END:%.+]], ptr [[P2]],
 111   // CK34-DAG: store i64 [[B_SIZE:%.+]], ptr [[S2]],
 112   // CK34-DAG: store ptr null, ptr [[MF2]],
 113
 114   // CK34-DAG: [[C_END]] = getelementptr %class.C, ptr [[C_ADDR]], i{{.+}} 1
 115
 116   // CK34-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 117   // CK34-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[C_END_INTPTR:%.+]]
 118   // CK34-DAG: [[C_END_INTPTR]] = ptrtoint ptr [[C_END]] to i64
 119   // CK34-DAG: [[S_END_INTPTR]] = ptrtoint ptr [[S_END_VOID:%.+]] to i64
 120   // CK34-DAG: [[S_END_VOID]] = getelementptr i8, ptr [[S_LAST:%.+]], i{{.+}} 1
 121   // CK34-64-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_ADDR]], i64 31
 122   // CK34-32-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_ADDR]], i32 15
 123
 124   // pass MEMBER_OF_1 | TO | IMPLICIT | MAPPER {&s, &s.c, 16} to copy the data of s.c.
 125
 126   // CK34-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
 127   // CK34-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
 128   // CK34-DAG: [[MF3:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 3
 129
 130   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP3]],
 131   // CK34-DAG: store ptr [[C_ADDR:%.+]], ptr [[P3]],
 132   // CK34-DAG: store ptr [[C_DEFAULT_MAPPER:@.+]], ptr [[MF3]],
 133
 134   // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 2
 135   // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 1
 136
 137   #pragma omp target map(to: s)
 138   s.foo();
 139
 140   // CK34 : call void
 141
 142   // CK34-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 143   // CK34-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 144   // CK34-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 145   // CK34-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 146   // CK34-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 147   // CK34-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 148   // CK34-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
 149   // CK34-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 150   // CK34-DAG: store ptr [[MF:%.+]], ptr [[MARG]]
 151   // CK34-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
 152   // CK34-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
 153   // CK34-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
 154
 155   // pass TARGET_PARAM {&s, &s, ((void*)(&s+1)-(void*)&s)}
 156
 157   // CK34-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
 158   // CK34-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
 159   // CK34-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
 160   // CK34-DAG: [[MF0:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 0
 161
 162   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP0]],
 163   // CK34-DAG: store ptr [[S_ADDR]], ptr [[P0]],
 164   // CK34-DAG: store i64 [[S_SIZE:%.+]], ptr [[S0]],
 165   // CK34-DAG: store ptr null, ptr [[MF0]],
 166
 167   // CK34-DAG: [[S_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 168   // CK34-DAG: [[SZ]] = sub i64 [[S_1_INTPTR:%.+]], [[S_INTPTR:%.+]]
 169   // CK34-DAG: [[S_1_INTPTR]] = ptrtoint ptr [[S_1:%.+]] to i64
 170   // CK34-DAG: [[S_INTPTR]] = ptrtoint ptr [[S_ADDR]] to i64
 171   // CK34-DAG: [[S_1]] = getelementptr %class.S, ptr [[S_ADDR]], i32 1
 172
 173   // pass MEMBER_OF_1 | FROM {&s, &s, ((void*)(&s.a+1)-(void*)&s)} to copy the data of s.a.
 174
 175   // CK34-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
 176   // CK34-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
 177   // CK34-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
 178   // CK34-DAG: [[MF1:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 1
 179
 180   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP1]],
 181   // CK34-DAG: store ptr [[S_ADDR]], ptr [[P1]],
 182   // CK34-DAG: store i64 [[A_SIZE:%.+]], ptr [[S1]],
 183   // CK34-DAG: store ptr null, ptr [[MF1]],
 184
 185   // CK34-DAG: [[A_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 186   // CK34-DAG: [[SZ]] = sub i64 [[C_BEGIN_INTPTR:%.+]], [[S_INTPTR:%.+]]
 187   // CK34-DAG: [[S_INTPTR]] = ptrtoint ptr [[S_ADDR]] to i64
 188   // CK34-DAG: [[C_BEGIN_INTPTR]] = ptrtoint ptr [[C_ADDR:%.+]] to i64
 189   // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 2
 190   // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 1
 191
 192   // pass MEMBER_OF_1 | FROM {&s, &s.c+1, ((void*)(&s)+31+1-(void*)(&s.c+1))} to copy the data of s.b.
 193
 194   // CK34-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
 195   // CK34-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
 196   // CK34-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
 197   // CK34-DAG: [[MF2:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 2
 198
 199   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP2]],
 200   // CK34-DAG: store ptr [[C_END:%.+]], ptr [[P2]],
 201   // CK34-DAG: store i64 [[B_SIZE:%.+]], ptr [[S2]],
 202   // CK34-DAG: store ptr null, ptr [[MF2]],
 203
 204   // CK34-DAG: [[C_END]] = getelementptr %class.C, ptr [[C_ADDR]], i{{.+}} 1
 205
 206   // CK34-DAG: [[B_SIZE]] = sdiv exact i64 [[SZ:%.+]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 207   // CK34-DAG: [[SZ]] = sub i64 [[S_END_INTPTR:%.+]], [[C_END_INTPTR:%.+]]
 208   // CK34-DAG: [[C_END_INTPTR]] = ptrtoint ptr [[C_END]] to i64
 209   // CK34-DAG: [[S_END_INTPTR]] = ptrtoint ptr [[S_END_VOID:%.+]] to i64
 210   // CK34-DAG: [[S_END_VOID]] = getelementptr i8, ptr [[S_LAST:%.+]], i{{.+}} 1
 211   // CK34-64-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_ADDR]], i64 31
 212   // CK34-32-DAG: [[S_LAST]] = getelementptr i8, ptr [[S_ADDR]], i32 15
 213
 214   // pass MEMBER_OF_1 | FROM | IMPLICIT | MAPPER {&s, &s.c, 16} to copy the data of s.c.
 215
 216   // CK34-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
 217   // CK34-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
 218   // CK34-DAG: [[MF3:%.+]] = getelementptr inbounds {{.+}}[[MF]], i{{.+}} 0, i{{.+}} 3
 219
 220   // CK34-DAG: store ptr [[S_ADDR]], ptr [[BP3]],
 221   // CK34-DAG: store ptr [[C_ADDR:%.+]], ptr [[P3]],
 222   // CK34-DAG: store ptr [[C_DEFAULT_MAPPER]], ptr [[MF3]],
 223
 224   // CK34-64-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 2
 225   // CK34-32-DAG: [[C_ADDR]] = getelementptr inbounds %class.S, ptr [[S_ADDR]], i32 0, i32 1
 226
 227   #pragma omp target map(from: s)
 228   s.foo();
 229 }
 230
 231 #endif // CK34
 232 #endif