ABataev created this revision.
ABataev added a reviewer: jdoerfert.
Herald added subscribers: guansong, yaxunl.
Herald added a project: clang.
ABataev requested review of this revision.
OMP_MAP_TARGET_PARAM flag is used to mark the data that shoud be passed
as arguments to the target kernels, nothing else. But the compiler still
marks the data with OMP_MAP_TARGET_PARAM flags even if the data is
passed to the data movement directives, like target data, target update
etc. This flag is just ignored for this directives and the compiler does
not need to emit it.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D91261
Files:
clang/lib/CodeGen/CGOpenMPRuntime.cpp
clang/test/OpenMP/declare_mapper_codegen.cpp
clang/test/OpenMP/target_data_codegen.cpp
clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
clang/test/OpenMP/target_data_use_device_ptr_if_codegen.cpp
clang/test/OpenMP/target_enter_data_codegen.cpp
clang/test/OpenMP/target_enter_data_depend_codegen.cpp
clang/test/OpenMP/target_exit_data_codegen.cpp
clang/test/OpenMP/target_exit_data_depend_codegen.cpp
clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
clang/test/OpenMP/target_update_codegen.cpp
clang/test/OpenMP/target_update_depend_codegen.cpp
Index: clang/test/OpenMP/target_update_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_update_depend_codegen.cpp
+++ clang/test/OpenMP/target_update_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
Index: clang/test/OpenMP/target_update_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_update_codegen.cpp
+++ clang/test/OpenMP/target_update_codegen.cpp
@@ -36,15 +36,15 @@
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -205,7 +205,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710674]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -342,7 +342,7 @@
#ifdef CK5
// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK5-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -384,7 +384,7 @@
#ifdef CK6
// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK6-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -431,7 +431,7 @@
#ifdef CK7
// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
-// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK7-LABEL: lvalue
void lvalue(int *B, int l, int e) {
@@ -481,7 +481,7 @@
#ifdef CK8
// CK8: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4]
-// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 1, i64 17]
// CK8-LABEL: lvalue
void lvalue(int **B, int l, int e) {
@@ -535,7 +535,7 @@
double *p;
};
-// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK9-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -585,7 +585,7 @@
double *p;
};
-// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK10-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -635,7 +635,7 @@
struct S {
double *p;
};
-// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK11-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -687,7 +687,7 @@
double *p;
struct S *sp;
};
-// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 0, i64 281474976710672, i64 17]
// CK12-LABEL: lvalue
void lvalue(struct S *s, int l, int e) {
@@ -749,7 +749,7 @@
#ifdef CK13
// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK13-LABEL: lvalue
void lvalue(int **BB, int a, int b) {
@@ -796,7 +796,7 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK14
-// CK14: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK14: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 281474976710673]
struct SSA {
double *p;
@@ -869,7 +869,7 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK15
-// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
struct SSA {
double *p;
@@ -935,7 +935,7 @@
#ifdef CK16
// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
//CK16-LABEL: lvalue_find_base
void lvalue_find_base(float *f, int *i) {
@@ -980,7 +980,7 @@
#ifdef CK17
// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 1]
struct SSA {
int i;
@@ -1038,8 +1038,8 @@
// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
#ifdef CK18
-// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33]
-// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 1]
+// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 2]
//CK18-LABEL: array_shaping
void array_shaping(float *f, int sa) {
@@ -1113,11 +1113,11 @@
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
#ifdef CK19
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1021
+// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
-// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+// PRESENT=0x1000 | FROM=0x2 = 0x1002
+// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
// CK19-LABEL: _Z13check_presenti
void check_present(int arg) {
@@ -1186,7 +1186,7 @@
// CK20: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
// CK20: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3]
-// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044417]
// CK20-LABEL: _Z3foo
void foo(int arg) {
@@ -1254,7 +1254,7 @@
// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x double*]]] }
// CK21: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
-// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 299067162755073]
+// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 299067162755073]
struct ST {
double *dptr[10][10][10];
@@ -1333,7 +1333,7 @@
// CK22: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
// CK22: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044417]
struct ST {
// CK22: _ZN2ST3fooEPA10_Pi
@@ -1409,7 +1409,7 @@
// CK23: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
// CK23: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK23: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+// CK23: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044417]
// CK23: foo
void foo(int arg) {
@@ -1486,7 +1486,7 @@
// CK24: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
// CK24: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK24: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+// CK24: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044417]
// CK24: foo
void foo(int arg) {
@@ -1563,7 +1563,7 @@
// CK25: [[STRUCT_DESCRIPTOR:%.+]] = type { i64, i64, i64 }
// CK25: [[MSIZE:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4, i64 3]
-// CK25: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044449, i64 33, i64 17592186044449]
+// CK25: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044417, i64 1, i64 17592186044417]
// CK25-LABEL: _Z3foo
void foo(int arg) {
Index: clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
+++ clang/test/OpenMP/target_map_member_expr_array_section_codegen.cpp
@@ -5,11 +5,11 @@
#ifndef HEADER
#define HEADER
-// 32 = 0x20 = OMP_MAP_TARGET_PARAM
+// 0 = OMP_MAP_NONE
// 281474976710656 = 0x1000000000000 = OMP_MAP_MEMBER_OF of 1-st element
-// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710656]
+// CHECK: [[MAP_ENTER:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 281474976710656]
// 281474976710664 = 0x1000000000008 = OMP_MAP_MEMBER_OF of 1-st element | OMP_MAP_DELETE
-// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 281474976710664]
+// CHECK: [[MAP_EXIT:@.+]] = private unnamed_addr constant [2 x i64] [i64 0, i64 281474976710664]
template <typename T>
struct S {
constexpr static int size = 6;
Index: clang/test/OpenMP/target_exit_data_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_exit_data_depend_codegen.cpp
+++ clang/test/OpenMP/target_exit_data_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 40]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710674]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
Index: clang/test/OpenMP/target_exit_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_exit_data_codegen.cpp
+++ clang/test/OpenMP/target_exit_data_codegen.cpp
@@ -36,19 +36,19 @@
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 38]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 6]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710672]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710672]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1058]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1026]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1062]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1030]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -265,7 +265,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710676]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710676]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -369,7 +369,7 @@
}
};
-// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711700]
+// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711700]
// CK4-LABEL: _Z3bari
int bar(int arg){
Index: clang/test/OpenMP/target_enter_data_depend_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_depend_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_depend_codegen.cpp
@@ -30,15 +30,15 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
Index: clang/test/OpenMP/target_enter_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_enter_data_codegen.cpp
+++ clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -37,19 +37,19 @@
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] zeroinitializer
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 5]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1025]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1029]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -264,11 +264,11 @@
ST<int> gb;
double gc[100];
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1425
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
// CK1A-LABEL: _Z3fooi
void foo(int arg) {
@@ -354,7 +354,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710677]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710677]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -504,7 +504,7 @@
}
};
-// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711701]
// CK5-LABEL: _Z3bari
int bar(int arg){
Index: clang/test/OpenMP/target_data_use_device_ptr_if_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_use_device_ptr_if_codegen.cpp
+++ clang/test/OpenMP/target_data_use_device_ptr_if_codegen.cpp
@@ -18,7 +18,7 @@
// SIMD-ONLY1-NOT: {{__kmpc|__tgt}}
#ifdef CK1
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [1 x i64] [i64 99]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [1 x i64] [i64 67]
// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 288]
// CK1: [[MTYPE02:@.+]] = {{.*}}constant [1 x i64] [i64 288]
Index: clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
+++ clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp
@@ -22,18 +22,18 @@
double *g;
// CK1: @g = global double*
-// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96]
-// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 99]
-// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 35]
-// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 99]
-// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
-// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 35, i64 96]
+// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 19, i64 64]
+// CK1: [[MTYPE01:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE03:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE04:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE05:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE06:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE07:@.+]] = {{.*}}constant [1 x i64] [i64 67]
+// CK1: [[MTYPE08:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 3]
+// CK1: [[MTYPE09:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
+// CK1: [[MTYPE10:@.+]] = {{.*}}constant [2 x i64] [i64 67, i64 67]
+// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
+// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 3, i64 64]
// CK1-LABEL: @_Z3foo
template<typename T>
@@ -346,10 +346,10 @@
#ifdef CK2
// CK2: [[ST:%.+]] = type { double*, double** }
-// 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]
+// CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
+// CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 0, i64 281474976710739]
+// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 3, i64 0, i64 562949953421392]
+// CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 0, i64 281474976710739, i64 281474976710736]
template <typename T>
struct ST {
Index: clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
+++ clang/test/OpenMP/target_data_use_device_addr_codegen.cpp
@@ -12,11 +12,11 @@
#define HEADER
// CHECK-DAG: [[SIZES1:@.+]] = private unnamed_addr constant [5 x i64] zeroinitializer
-// 96 = 0x60 = OMP_MAP_TARGET_PARAM | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 96, i64 96, i64 96, i64 96, i64 96]
-// 32 = 0x20 = OMP_MAP_TARGET_PARAM
+// 64 = 0x40 = OMP_MAP_RETURN_PARAM
+// CHECK-DAG: [[MAPTYPES1:@.+]] = private unnamed_addr constant [5 x i64] [i64 64, i64 64, i64 64, i64 64, i64 64]
+// 0 = OMP_MAP_NONE
// 281474976710720 = 0x1000000000040 = OMP_MAP_MEMBER_OF | OMP_MAP_RETURN_PARAM
-// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 32, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
+// CHECK-DAG: [[MAPTYPES2:@.+]] = private unnamed_addr constant [5 x i64] [i64 0, i64 281474976710720, i64 281474976710720, i64 281474976710720, i64 281474976710720]
struct S {
int a = 0;
int *ptr = &a;
Index: clang/test/OpenMP/target_data_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_data_codegen.cpp
+++ clang/test/OpenMP/target_data_codegen.cpp
@@ -30,19 +30,19 @@
double gc[100];
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
-// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
-// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 1]
-// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37]
+// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 5]
// CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 sdiv exact (i64 sub (i64 ptrtoint (double** getelementptr (double*, double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (double** getelementptr inbounds (%struct.ST, %struct.ST* @gb, i32 0, i32 1) to i64)), i64 ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)), i64 24]
-// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710673]
-// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1057]
+// CK1: [[MTYPE05:@.+]] = {{.+}}constant [1 x i64] [i64 1025]
-// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1061]
+// CK1: [[MTYPE06:@.+]] = {{.+}}constant [1 x i64] [i64 1029]
// CK1-LABEL: _Z3fooi
void foo(int arg) {
@@ -255,17 +255,17 @@
ST<int> gb;
double gc[100];
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK1A: [[MTYPE00Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
-// TARGET_PARAM=0x20 | TO=0x1 = 0x21
-// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x21]]]
+// TO=0x1 = 0x1
+// CK1A: [[MTYPE00End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1]]]
-// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
-// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+// PRESENT=0x1000 | CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x1405
+// CK1A: [[MTYPE01Begin:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1405]]]
-// CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x425
-// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x425]]]
+// CLOSE=0x400 | ALWAYS=0x4 | TO=0x1 = 0x405
+// CK1A: [[MTYPE01End:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x405]]]
// CK1A-LABEL: _Z3fooi
void foo(int arg) {
@@ -357,7 +357,7 @@
}
};
-// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710677]
+// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710677]
// CK2-LABEL: _Z3bari
int bar(int arg){
@@ -475,7 +475,7 @@
}
};
-// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976711701]
+// CK4: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976711701]
// CK4-LABEL: _Z3bari
int bar(int arg){
@@ -561,7 +561,7 @@
void test_close_modifier(int arg) {
S2 *ps;
- // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, i64 16, i64 1043]
+ // CK5: private unnamed_addr constant [5 x i64] [i64 1027, i64 0, i64 562949953421328, i64 16, i64 1043]
#pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s)
{
++(arg);
@@ -585,7 +585,7 @@
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
#ifdef CK6
void test_close_modifier(int arg) {
- // CK6: private unnamed_addr constant [1 x i64] [i64 1059]
+ // CK6: private unnamed_addr constant [1 x i64] [i64 1027]
#pragma omp target data map(close,tofrom: arg)
{++arg;}
}
@@ -644,30 +644,30 @@
// ps1
//
- // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+ // PRESENT=0x1000 = 0x1000
// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
// PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
//
- // CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]],
+ // CK8-SAME: {{^}} [i64 [[#0x1000]], i64 [[#0x1000000000003]],
// CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
// arg
//
- // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
+ // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
//
- // CK8-SAME: {{^}} i64 [[#0x1023]],
+ // CK8-SAME: {{^}} i64 [[#0x1003]],
// ps2
//
- // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+ // PRESENT=0x1000 = 0x1000
// MEMBER_OF_7=0x7000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003
// MEMBER_OF_7=0x7000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010
// PTR_AND_OBJ=0x10 = 0x10
// PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
//
- // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x7000000001003]],
+ // CK8-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x7000000001003]],
// CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
#pragma omp target data map(tofrom: ps1->s) \
map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
@@ -694,8 +694,8 @@
// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
#ifdef CK9
void test_present_modifier(int arg) {
- // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
- // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1023]]]
+ // PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1003
+ // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1003]]]
#pragma omp target data map(present,tofrom: arg)
{++arg;}
}
Index: clang/test/OpenMP/declare_mapper_codegen.cpp
===================================================================
--- clang/test/OpenMP/declare_mapper_codegen.cpp
+++ clang/test/OpenMP/declare_mapper_codegen.cpp
@@ -58,25 +58,25 @@
// CK0: [[TEAMNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
// CK0-64: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+// CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
// CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
// CK0-64: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK0-32: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34]
+// CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
class C {
public:
@@ -141,10 +141,10 @@
// CK0-DAG: [[MEMBER]]
// CK0-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK0-DAG: [[MEMBERCOM]]
-// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK0-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK0-DAG: br label %[[LTYPE]]
// CK0-DAG: [[LTYPE]]
-// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK0-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -701,10 +701,10 @@
// CK1-DAG: [[MEMBER]]
// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK1-DAG: [[MEMBERCOM]]
-// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK1-DAG: br label %[[LTYPE]]
// CK1-DAG: [[LTYPE]]
-// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -864,10 +864,10 @@
// CK2-DAG: [[MEMBER]]
// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK2-DAG: [[MEMBERCOM]]
-// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK2-DAG: br label %[[LTYPE]]
// CK2-DAG: [[LTYPE]]
-// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
@@ -1065,13 +1065,13 @@
// CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
-// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+// PRESENT=0x1000 | TO=0x1 = 0x1001
+// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
// CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
// CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
-// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022
-// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]]
+// PRESENT=0x1000 | FROM=0x2 = 0x1002
+// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
class C {
public:
@@ -1136,10 +1136,10 @@
// CK4-DAG: [[MEMBER]]
// CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]]
// CK4-DAG: [[MEMBERCOM]]
-// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]]
+// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
// CK4-DAG: br label %[[LTYPE]]
// CK4-DAG: [[LTYPE]]
-// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
+// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 0, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ]
// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -8187,7 +8187,7 @@
void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
MapFlagsArrayTy &CurTypes,
const StructRangeInfoTy &PartialStruct,
- bool NotTargetParams = false) const {
+ bool NotTargetParams = true) const {
// Base is the base of the struct
CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer());
// Pointer is the address of the lowest element
@@ -8232,7 +8232,7 @@
/// pair of the relevant declaration and index where it occurs is appended to
/// the device pointers info array.
void generateAllInfo(
- MapCombinedInfoTy &CombinedInfo, bool NotTargetParams = false,
+ MapCombinedInfoTy &CombinedInfo,
const llvm::DenseSet<CanonicalDeclPtr<const Decl>> &SkipVarSet =
llvm::DenseSet<CanonicalDeclPtr<const Decl>>()) const {
// We have to process the component lists that relate with the same
@@ -8355,9 +8355,7 @@
UseDevicePtrCombinedInfo.Pointers.push_back(Ptr);
UseDevicePtrCombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- UseDevicePtrCombinedInfo.Types.push_back(
- OMP_MAP_RETURN_PARAM |
- (NotTargetParams ? OMP_MAP_NONE : OMP_MAP_TARGET_PARAM));
+ UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
UseDevicePtrCombinedInfo.Mappers.push_back(nullptr);
}
}
@@ -8424,19 +8422,13 @@
CombinedInfo.Pointers.push_back(Ptr);
CombinedInfo.Sizes.push_back(
llvm::Constant::getNullValue(CGF.Int64Ty));
- CombinedInfo.Types.push_back(
- OMP_MAP_RETURN_PARAM |
- (NotTargetParams ? OMP_MAP_NONE : OMP_MAP_TARGET_PARAM));
+ CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM);
CombinedInfo.Mappers.push_back(nullptr);
}
}
}
for (const auto &M : Info) {
- // We need to know when we generate information for the first component
- // associated with a capture, because the mapping flags depend on it.
- bool IsFirstComponentList = !NotTargetParams;
-
// Temporary generated information.
MapCombinedInfoTy CurInfo;
StructRangeInfoTy PartialStruct;
@@ -8449,10 +8441,10 @@
unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size();
CurInfo.NonContigInfo.IsNonContiguous =
L.Components.back().isNonContiguous();
- generateInfoForComponentList(L.MapType, L.MapModifiers,
- L.MotionModifiers, L.Components, CurInfo,
- PartialStruct, IsFirstComponentList,
- L.IsImplicit, L.Mapper, L.ForDeviceAddr);
+ generateInfoForComponentList(
+ L.MapType, L.MapModifiers, L.MotionModifiers, L.Components, CurInfo,
+ PartialStruct, /*IsFirstComponentList=*/false, L.IsImplicit,
+ L.Mapper, L.ForDeviceAddr);
// If this entry relates with a device pointer, set the relevant
// declaration and add the 'return pointer' flag.
@@ -8469,7 +8461,6 @@
RelevantVD);
CurInfo.Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PARAM;
}
- IsFirstComponentList = false;
}
// Append any pending zero-length pointers which are struct members and
@@ -8510,8 +8501,7 @@
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid())
- emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
- NotTargetParams);
+ emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
@@ -9957,7 +9947,8 @@
// If there is an entry in PartialStruct it means we have a struct with
// individual members mapped. Emit an extra combined entry.
if (PartialStruct.Base.isValid())
- MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
+ MEHandler.emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct,
+ /*NoTargetParam=*/false);
// We need to append the results of this capture to what we already have.
CombinedInfo.append(CurInfo);
@@ -9968,8 +9959,7 @@
CombinedInfo.Types);
// Map any list items in a map clause that were not captures because they
// weren't referenced within the construct.
- MEHandler.generateAllInfo(CombinedInfo, /*NotTargetParams=*/true,
- MappedVarSet);
+ MEHandler.generateAllInfo(CombinedInfo, MappedVarSet);
TargetDataInfo Info;
// Fill up the arrays and create the arguments.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
