The patch adds support for the new map interface between clang and libomptarget. The changes in the interface are the following:

device IDs are now 64-bit integers (as opposed to 32-bit)
map flags are 64-bit long (used to be 32-bit)
mappings for partially mapped structs are now calculated at compile time and members of partially mapped structs are flagged using the MEMBER_OF field
Support for is_device_ptr on struct members was dropped - this functionality is not supported by the OpenMP standard and its implementation is technically infeasible (however, use_device_ptr on struct members works as a non-standard extension of the compiler)

llvm-svn: 337468

1 files changed, 29 insertions, 27 deletions

diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
index d6f18cca36c..a3d8043b6b4 100644
--- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -346,10 +346,10 @@ void bar(float *&a, int *&b) {
 #ifdef CK2
 
 // CK2: [[ST:%.+]] = type { double*, double** }
-// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 83]
-// CK2: [[MTYPE01:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 19, i64 83]
-// CK2: [[MTYPE02:@.+]] = {{.*}}constant [2 x i64] [i64 96, i64 35]
-// CK2: [[MTYPE03:@.+]] = {{.*}}constant [4 x i64] [i64 96, i64 32, i64 19, i64 83]
+// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
+// CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 35, i64 32, i64 562949953421392]
+// CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736]
 
 template <typename T>
 struct ST {
@@ -382,7 +382,7 @@ struct ST {
     // CK2:     getelementptr inbounds double, double* [[TTT]], i32 1
     a++;
 
-    // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
+    // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1
     // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
     // CK2:     store double** [[RVAL:%.+]], double*** [[CBP]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE01]]
@@ -404,11 +404,12 @@ struct ST {
     // CK2:     getelementptr inbounds double, double* [[TTTT]], i32 1
     b++;
 
-    // CK2:     [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0
-    // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double**
-    // CK2:     store double* [[RVAL:%.+]], double** [[CBP]],
+    // CK2:     [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
+    // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double***
+    // CK2:     store double** [[RVAL:%.+]], double*** [[CBP]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE02]]
-    // CK2:     [[VAL:%.+]] = load double*, double** [[CBP]],
+    // CK2:     [[CVAL:%.+]] = bitcast double*** [[CBP]] to double**
+    // CK2:     [[VAL:%.+]] = load double*, double** [[CVAL]],
     // CK2:     store double* [[VAL]], double** [[PVT:%.+]],
     // CK2:     store double** [[PVT]], double*** [[PVT2:%.+]],
     // CK2:     [[TT1:%.+]] = load double**, double*** [[PVT2]],
@@ -426,26 +427,27 @@ struct ST {
     a++;
     la++;
 
-    // CK2:     [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 0
-    // CK2:     [[CBP:%.+]] = bitcast i8** [[BP]] to double**
-    // CK2:     store double* [[RVAL:%.+]], double** [[CBP]],
-    // CK2:     [[_BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 3
-    // CK2:     [[_CBP:%.+]] = bitcast i8** [[_BP]] to double***
-    // CK2:     store double** [[_RVAL:%.+]], double*** [[_CBP]],
+    // CK2:     [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 1
+    // CK2:     [[CBP1:%.+]] = bitcast i8** [[BP1]] to double***
+    // CK2:     store double** [[RVAL1:%.+]], double*** [[CBP1]],
+    // CK2:     [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2
+    // CK2:     [[CBP2:%.+]] = bitcast i8** [[BP2]] to double***
+    // CK2:     store double** [[RVAL2:%.+]], double*** [[CBP2]],
     // CK2:     call void @__tgt_target_data_begin{{.+}}[[MTYPE03]]
-    // CK2:     [[VAL:%.+]] = load double*, double** [[CBP]],
-    // CK2:     store double* [[VAL]], double** [[PVT:%.+]],
-    // CK2:     store double** [[PVT]], double*** [[PVT2:%.+]],
-    // CK2:     [[_CBP1:%.+]] = bitcast double*** [[_CBP]] to double**
-    // CK2:     [[_VAL:%.+]] = load double*, double** [[_CBP1]],
-    // CK2:     store double* [[_VAL]], double** [[_PVT:%.+]],
-    // CK2:     store double** [[_PVT]], double*** [[_PVT2:%.+]],
-    // CK2:     [[TT1:%.+]] = load double**, double*** [[PVT2]],
-    // CK2:     [[TT2:%.+]] = load double*, double** [[TT1]],
-    // CK2:     getelementptr inbounds double, double* [[TT2]], i32 1
-    // CK2:     [[_TT1:%.+]] = load double**, double*** [[_PVT2]],
-    // CK2:     [[_TT2:%.+]] = load double*, double** [[_TT1]],
+    // CK2:     [[_CBP2:%.+]] = bitcast double*** [[CBP2]] to double**
+    // CK2:     [[VAL2:%.+]] = load double*, double** [[_CBP2]],
+    // CK2:     store double* [[VAL2]], double** [[PVT2:%.+]],
+    // CK2:     store double** [[PVT2]], double*** [[_PVT2:%.+]],
+    // CK2:     [[_CBP1:%.+]] = bitcast double*** [[CBP1]] to double**
+    // CK2:     [[VAL1:%.+]] = load double*, double** [[_CBP1]],
+    // CK2:     store double* [[VAL1]], double** [[PVT1:%.+]],
+    // CK2:     store double** [[PVT1]], double*** [[_PVT1:%.+]],
+    // CK2:     [[TT2:%.+]] = load double**, double*** [[_PVT2]],
+    // CK2:     [[_TT2:%.+]] = load double*, double** [[TT2]],
     // CK2:     getelementptr inbounds double, double* [[_TT2]], i32 1
+    // CK2:     [[TT1:%.+]] = load double**, double*** [[_PVT1]],
+    // CK2:     [[_TT1:%.+]] = load double*, double** [[TT1]],
+    // CK2:     getelementptr inbounds double, double* [[_TT1]], i32 1
     #pragma omp target data map(b[:10]) use_device_ptr(a, b)
     {
       a++;