summaryrefslogtreecommitdiffstats
path: root/clang
diff options
context:
space:
mode:
Diffstat (limited to 'clang')
-rw-r--r--clang/lib/CodeGen/CGOpenMPRuntime.cpp67
-rw-r--r--clang/test/OpenMP/target_enter_data_codegen.cpp4
-rw-r--r--clang/test/OpenMP/target_exit_data_codegen.cpp4
-rw-r--r--clang/test/OpenMP/target_update_codegen.cpp4
4 files changed, 70 insertions, 9 deletions
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 2359a32d561..5db29eb6004 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -691,12 +691,24 @@ enum OpenMPRTLFunction {
// Call to void __tgt_target_data_begin(int64_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
OMPRTL__tgt_target_data_begin,
+ // Call to void __tgt_target_data_begin_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ OMPRTL__tgt_target_data_begin_nowait,
// Call to void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
OMPRTL__tgt_target_data_end,
+ // Call to void __tgt_target_data_end_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ OMPRTL__tgt_target_data_end_nowait,
// Call to void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
OMPRTL__tgt_target_data_update,
+ // Call to void __tgt_target_data_update_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ OMPRTL__tgt_target_data_update_nowait,
};
/// A basic class for pre|post-action for advanced codegen sequence for OpenMP
@@ -2136,6 +2148,21 @@ CGOpenMPRuntime::createRuntimeFunction(unsigned Function) {
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_begin");
break;
}
+ case OMPRTL__tgt_target_data_begin_nowait: {
+ // Build void __tgt_target_data_begin_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.SizeTy->getPointerTo(),
+ CGM.Int64Ty->getPointerTo()};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_begin_nowait");
+ break;
+ }
case OMPRTL__tgt_target_data_end: {
// Build void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
@@ -2150,6 +2177,21 @@ CGOpenMPRuntime::createRuntimeFunction(unsigned Function) {
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end");
break;
}
+ case OMPRTL__tgt_target_data_end_nowait: {
+ // Build void __tgt_target_data_end_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.SizeTy->getPointerTo(),
+ CGM.Int64Ty->getPointerTo()};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_end_nowait");
+ break;
+ }
case OMPRTL__tgt_target_data_update: {
// Build void __tgt_target_data_update(int64_t device_id, int32_t arg_num,
// void** args_base, void **args, size_t *arg_sizes, int64_t *arg_types);
@@ -2164,6 +2206,21 @@ CGOpenMPRuntime::createRuntimeFunction(unsigned Function) {
RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update");
break;
}
+ case OMPRTL__tgt_target_data_update_nowait: {
+ // Build void __tgt_target_data_update_nowait(int64_t device_id, int32_t
+ // arg_num, void** args_base, void **args, size_t *arg_sizes, int64_t
+ // *arg_types);
+ llvm::Type *TypeParams[] = {CGM.Int64Ty,
+ CGM.Int32Ty,
+ CGM.VoidPtrPtrTy,
+ CGM.VoidPtrPtrTy,
+ CGM.SizeTy->getPointerTo(),
+ CGM.Int64Ty->getPointerTo()};
+ auto *FnTy =
+ llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+ RTLFn = CGM.CreateRuntimeFunction(FnTy, "__tgt_target_data_update_nowait");
+ break;
+ }
}
assert(RTLFn && "Unable to find OpenMP runtime function");
return RTLFn;
@@ -7524,19 +7581,23 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
auto &RT = CGF.CGM.getOpenMPRuntime();
// Select the right runtime function call for each expected standalone
// directive.
+ const bool HasNowait = D.hasClausesOfKind<OMPNowaitClause>();
OpenMPRTLFunction RTLFn;
switch (D.getDirectiveKind()) {
default:
llvm_unreachable("Unexpected standalone target data directive.");
break;
case OMPD_target_enter_data:
- RTLFn = OMPRTL__tgt_target_data_begin;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_begin_nowait
+ : OMPRTL__tgt_target_data_begin;
break;
case OMPD_target_exit_data:
- RTLFn = OMPRTL__tgt_target_data_end;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_end_nowait
+ : OMPRTL__tgt_target_data_end;
break;
case OMPD_target_update:
- RTLFn = OMPRTL__tgt_target_data_update;
+ RTLFn = HasNowait ? OMPRTL__tgt_target_data_update_nowait
+ : OMPRTL__tgt_target_data_update;
break;
}
CGF.EmitRuntimeCall(RT.createRuntimeFunction(RTLFn), OffloadingArgs);
diff --git a/clang/test/OpenMP/target_enter_data_codegen.cpp b/clang/test/OpenMP/target_enter_data_codegen.cpp
index fb916733f21..08b6c7c2702 100644
--- a/clang/test/OpenMP/target_enter_data_codegen.cpp
+++ b/clang/test/OpenMP/target_enter_data_codegen.cpp
@@ -38,7 +38,7 @@ void foo(int arg) {
float lb[arg];
// Region 00
- // CK1-DAG: call void @__tgt_target_data_begin(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // 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:%[^,]+]]
@@ -53,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
diff --git a/clang/test/OpenMP/target_exit_data_codegen.cpp b/clang/test/OpenMP/target_exit_data_codegen.cpp
index ce13c5202ff..9359e3be451 100644
--- a/clang/test/OpenMP/target_exit_data_codegen.cpp
+++ b/clang/test/OpenMP/target_exit_data_codegen.cpp
@@ -39,7 +39,7 @@ void foo(int arg) {
// Region 00
// CK1-NOT: __tgt_target_data_begin
- // CK1-DAG: call void @__tgt_target_data_end(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1-DAG: call void @__tgt_target_data_end_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:%[^,]+]]
@@ -53,7 +53,7 @@ void foo(int arg) {
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PC0]]
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- #pragma omp target exit data if(1+3-5) device(arg) map(from: gc)
+ #pragma omp target exit data if(1+3-5) device(arg) map(from: gc) nowait
{++arg;}
// Region 01
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
index ae2909ddbea..7f45c313a8a 100644
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -38,7 +38,7 @@ void foo(int arg) {
float lb[arg];
// Region 00
- // CK1-DAG: call void @__tgt_target_data_update(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK1-DAG: call void @__tgt_target_data_update_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:%[^,]+]]
@@ -52,7 +52,7 @@ void foo(int arg) {
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PC0]]
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
- #pragma omp target update if(1+3-5) device(arg) from(gc)
+ #pragma omp target update if(1+3-5) device(arg) from(gc) nowait
{++arg;}
// Region 01
OpenPOWER on IntegriCloud