r320717 - [OPENMP] Add codegen for target data constructs with `nowait` clause.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Thu Dec 14 09:00:18 PST 2017
Author: abataev
Date: Thu Dec 14 09:00:17 2017
New Revision: 320717
URL: http://llvm.org/viewvc/llvm-project?rev=320717&view=rev
Log:
[OPENMP] Add codegen for target data constructs with `nowait` clause.
Added codegen for the `nowait` clause in target data constructs.
Modified:
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp
cfe/trunk/test/OpenMP/target_exit_data_codegen.cpp
cfe/trunk/test/OpenMP/target_update_codegen.cpp
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=320717&r1=320716&r2=320717&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Dec 14 09:00:17 2017
@@ -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(u
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(u
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(u
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::emitTargetDataStan
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);
Modified: cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp?rev=320717&r1=320716&r2=320717&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_enter_data_codegen.cpp Thu Dec 14 09:00:17 2017
@@ -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
Modified: cfe/trunk/test/OpenMP/target_exit_data_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_exit_data_codegen.cpp?rev=320717&r1=320716&r2=320717&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_exit_data_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_exit_data_codegen.cpp Thu Dec 14 09:00:17 2017
@@ -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
Modified: cfe/trunk/test/OpenMP/target_update_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_update_codegen.cpp?rev=320717&r1=320716&r2=320717&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_update_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_update_codegen.cpp Thu Dec 14 09:00:17 2017
@@ -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
More information about the cfe-commits
mailing list