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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits