diff options
Diffstat (limited to 'test/OpenMP/target_enter_data_codegen.cpp')
| -rw-r--r-- | test/OpenMP/target_enter_data_codegen.cpp | 28 |
1 files changed, 15 insertions, 13 deletions
diff --git a/test/OpenMP/target_enter_data_codegen.cpp b/test/OpenMP/target_enter_data_codegen.cpp index 683cd976d395..08b6c7c2702e 100644 --- a/test/OpenMP/target_enter_data_codegen.cpp +++ b/test/OpenMP/target_enter_data_codegen.cpp @@ -22,15 +22,15 @@ ST<int> gb; double gc[100]; // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800] -// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i32] [i32 32] +// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32] // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] -// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i32] [i32 33] +// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33] -// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i32] [i32 37] +// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37] // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] 24] -// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i32] [i32 33, i32 17] +// CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] // CK1-LABEL: _Z3fooi void foo(int arg) { @@ -38,8 +38,9 @@ void foo(int arg) { float lb[arg]; // Region 00 - // CK1-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) - // CK1-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, + // CK1-DAG: call void @__tgt_target_data_begin_nowait(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) + // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 + // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}}, // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -52,7 +53,7 @@ void foo(int arg) { // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1-NOT: __tgt_target_data_end - #pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc) + #pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc) nowait {++arg;} // Region 01 @@ -63,7 +64,7 @@ void foo(int arg) { // Region 02 // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] // CK1: [[IFTHEN]] - // CK1-DAG: call void @__tgt_target_data_begin(i32 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) + // CK1-DAG: call void @__tgt_target_data_begin(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -87,7 +88,7 @@ void foo(int arg) { {++arg;} // Region 03 - // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) + // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] @@ -110,7 +111,7 @@ void foo(int arg) { {++arg;} // Region 04 - // CK1-DAG: call void @__tgt_target_data_begin(i32 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}) + // CK1-DAG: call void @__tgt_target_data_begin(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}) // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] @@ -161,7 +162,7 @@ struct ST { }; // CK2: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 24] -// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i32] [i32 37, i32 21] +// CK2: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 37, i64 21] // CK2-LABEL: _Z3bari int bar(int arg){ @@ -172,8 +173,9 @@ int bar(int arg){ // Region 00 // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] // CK2: [[IFTHEN]] -// CK2-DAG: call void @__tgt_target_data_begin(i32 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) -// CK2-DAG: [[DEV]] = load i32, i32* %{{[^,]+}}, +// CK2-DAG: call void @__tgt_target_data_begin(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) +// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 +// CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}}, // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] |