[OpenMP] Clang Support for taskwait nowait clause

[koops]

Summary

Support for taskwait nowait clause with placeholder for runtime changes.

[tianshilei1992]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

5817

better to use consistent name as declaration. how about HasNowaitClause?

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

634 ↗

(On Diff #452924)

I'd keep the style consistent. Here there is a default argument bool HaveNoTaskWaitClause = false but no for emitTaskwaitCall. It's better to use default argument (potentially save a couple of changes) for both, or no default for both.

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

1232

Is the OpenMPVer necessary as an argument? Do we have alternative way to tell the version? If not, I think it makes more sense to take it as a data member of OpenMPIRBuilder instead of passing as a function argument.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

5817

I would recommend to make it a member of OMPTaskDataTy

[tianshilei1992]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

5817

Yeah, that sounds better.

[koops]

In the latest revision that was submitted the following changes were made:

1. Changed HaveNoTaskWaitClause to HasNowaitClause.
2. HasNowaitClause is now a data member of OMPTaskDataTy.
3. OpenMPVer is now a data member of OpenMPIRBuilder.

[ABataev]

clang/lib/CodeGen/CGStmtOpenMP.cpp

5266

Data.HasNowaitClause = S.hasClausesOfKind<OMPNowaitClause>();

[koops]

Contains changes suggested for:

Comment at: clang/lib/CodeGen/CGStmtOpenMP.cpp:5215
+ if (S.hasClausesOfKind<OMPNowaitClause>()) Data.HasNowaitClause = true;
Replaced with:
Data.HasNowaitClause = S.hasClausesOfKind<OMPNowaitClause>();

[koops]

I seem to have made changes for all the feedback which were provided. Can someone please look at the changes I have uploaded?

[ABataev]

clang/lib/CodeGen/CGStmtOpenMP.cpp

5266–5267

No need for if here, just

Data.HasNowaitClause = S.hasClausesOfKind<OMPNowaitClause>();

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

1418 ↗

(On Diff #462397)

int OpenMPVersion = 0?

1421 ↗

(On Diff #462397)

int getOpenMPVersion() const {return OpenMPVersion; }
Plus, do you really need an OpenMP version check here? Just call the new function always

[koops]

The current patch has the following changes:

1. "if" statement removed for Data.HasNowaitClause = S.hasClausesOfKind<OMPNowaitClause>();
2. OpenMPVer replaced with OpenMPVersion.

[koops]

Reply

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

1421 ↗

(On Diff #462397)

The new function __kmpc_omp_taskwait_deps_51, in openmp/runtime/src/kmp_taskdeps.cpp, is a placeholder for taking into account "nowait" clause from Spec 5.1. The scheduling of tasks might change during runtime due to "nowait" clause. Hence a check for OpenMP version is needed.

[koops]

Changes for the build failure due to git-clang-format. Some minor changes are also because of git pull.

[jdoerfert]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

5873

Same as above.

5878

Can we avoid the duplication, always make the array 7 long, and use the ArrayRef substring capability to pass in either 6 or 7 elements?

Conceptually something like this:

Fill all
emitRuntimeCall(..., ArrayRef(DepWaitTaskArgs, OpenMP >= 51 ? 7: 6)

On second thought, why do we need to ever emit the old runtime call at all? Always emit the new one.

llvm/include/llvm/Frontend/OpenMP/OMPIRBuilder.h

80 ↗

(On Diff #469120)

Not sure a default is wise here. Users should pick. And we might want to have a "latest" constant somewhere, just set to 99.

643 ↗

(On Diff #469120)

The clause is not NoTaskwait, it's nowait, right? Why not HasNowaitClause?

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

1250

Same as above, new version only is sufficient.

2370

Unrelated and not even sure which one we should use nowadays. Please split this off.

openmp/runtime/src/kmp.h

4015

No default arguments here please.

openmp/runtime/src/kmp_taskdeps.cpp

905

Implement the old version by calling the new version with no_wait false. I assume you copied the original code, this should have indicated a problem (wrt. code duplication).

[koops]

Johannes Doerfert,
Most of your comments are related to avoiding having the kmpc_omp_taskwait_51( ) , in the runtime, and instead introduce a "nowait" related parameter in the kmpc_omp_taskwait( ) itself. I have explained earlier that kmpc_omp_taskwait_51( ) is a placeholder for introducing new code, at a later stage. The scheduling of tasks might change in the runtime due to "nowait" clause. I have done this with the sole purpose of having backward compatibility. If you feel that is not needed then I will remove the new code and modify the Code Generation also according to your comments.

Thank you for your detailed comments.

[jdoerfert]

In D131830#3874169, @koops wrote:

Johannes Doerfert,
Most of your comments are related to avoiding having the kmpc_omp_taskwait_51( ) , in the runtime, and instead introduce a "nowait" related parameter in the kmpc_omp_taskwait( ) itself. I have explained earlier that kmpc_omp_taskwait_51( ) is a placeholder for introducing new code, at a later stage. The scheduling of tasks might change in the runtime due to "nowait" clause. I have done this with the sole purpose of having backward compatibility. If you feel that is not needed then I will remove the new code and modify the Code Generation also according to your comments.

My comments are not about avoiding kmpc_omp_taskwait_51 in the runtime. On the contrary. My comments say: Always use kmpc_omp_taskwait_51 and eliminate the old kmpc_omp_taskwait. Do not generate calls to kmpc_omp_taskwait but to kmpc_omp_taskwait_51 instead. Implement kmpc_omp_taskwait with a call to kmpc_omp_taskwait_51. Doing both will remove most of the new code and we still get the same benefits.
I hope this is clearer now.

[koops]

Changes for the suggestion : "Always use kmpc_omp_taskwait_51 and eliminate the old kmpc_omp_taskwait".

[koops]

Can someone please review my latest changes?

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

4769

Use CGF.Int32Ty

5840

Use CGF.Int32Ty

[koops]

Changes for "Use CGF.Int32Ty" instead of CGF.IntTy.

[ABataev]

LG

[koops]

No Changes. Pulling in recent changes to trigger a rebuild.

[koops]

No changes. Pulling in recent changes to trigger a rebuild.

[ABataev]

Do you need to commit it?

[koops]

Yes, I don't have commit rights. I have tried to reproduce the pre-merge check failures on RedHat and Ubuntu but, the tests don't fail.

[ABataev]

Ok, I'll commit it.

[mstorsjo]

openmp/runtime/src/dllexports

349

This is missing a similar export of __kmpc_omp_taskwait_51, which causes a bunch of tests to fail if running them on Windows now. (Also, where do the numbers for these exports get assigned?)

[koops]

Addition of __kmpc_omp_taskwait_51 in dllexports.

[koops]

Martin Storsjö (or anybody who works on windows) I have uploaded a new patch which has __kmpc_omp_taskwait_51 in dllexports. Can you please check the patch on windows?

[koops]

git pull & reloading the earlier patch.

[mstorsjo]

In D131830#3985449, @koops wrote:

Martin Storsjö (or anybody who works on windows) I have uploaded a new patch which has __kmpc_omp_taskwait_51 in dllexports. Can you please check the patch on windows?

Thanks, now this seems to work fine for me! (There seems to be another unrelated regression at the same time though, regarding the symbol __kmpc_fork_call_if.)

[koops]

(There seems to be another unrelated regression at the same time though, regarding the symbol __kmpc_fork_call_if.)

I do not know about this regression. I have not touched this function.

[mstorsjo]

In D131830#3986436, @koops wrote:

(There seems to be another unrelated regression at the same time though, regarding the symbol __kmpc_fork_call_if.)

I do not know about this regression. I have not touched this function.

Yeah, that one should be separate, I guess it stems from D138495.

This updated patch should be ok with me to recommit.

[protze.joachim]

Why do we need __kmpc_omp_taskwait_51 ? Taskwait nowait only makes sense with dependencies. It should only result in additional nodes/edges in the task dependency graph without any code to execute. There is not taskwait (aka task barrier) in this case.

The taskwait construct has the restriction in the spec:

The nowait clause may only appear on a taskwait directive if the depend clause is present.

[nikic]

clang/test/OpenMP/taskwait_depend_nowait_codegen.cpp

2

Please do not add any new tests using -no-opaque-pointers.

[koops]

Taking care of :

1. "The nowait clause may only appear on a taskwait directive if the depend clause is present.".
2. "Please do not add any new tests using -no-opaque-pointers".
3. Added a new test to expect failure for point 1.

[ABataev]

LG

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_target_contains_not_only_teams : Error<
   "target construct with nested teams region contains statements outside of the teams construct">;
 def note_omp_nested_teams_construct_here : Note<
   "nested teams construct here">;
 def note_omp_nested_statement_here : Note<
   "%select{statement|directive}0 outside teams construct here">;
 def err_omp_single_copyprivate_with_nowait : Error<
   "the 'copyprivate' clause must not be used with the 'nowait' clause">;
+def err_omp_nowait_clause_without_depend: Error<
+  "directive '#pragma omp taskwait' cannot use 'nowait' clause without 'depend' clause">;
 def note_omp_nowait_clause_here : Note<
   "'nowait' clause is here">;
 def err_omp_single_decl_in_declare_simd_variant : Error<
   "single declaration is expected after 'declare %select{simd|variant}0' directive">;
 def err_omp_function_expected : Error<
   "'#pragma omp declare %select{simd|variant}0' can only be applied to functions">;
 def err_omp_wrong_cancel_region : Error<
   "one of 'for', 'parallel', 'sections' or 'taskgroup' is expected">;

clang/lib/CodeGen/CGOpenMPRuntime.h

struct OMPTaskDataTy final {
   llvm::PointerIntPair<llvm::Value *, 1, bool> Schedule;
   llvm::PointerIntPair<llvm::Value *, 1, bool> Priority;
   llvm::Value *Reductions = nullptr;
   unsigned NumberOfParts = 0;
   bool Tied = true;
   bool Nogroup = false;
   bool IsReductionWithTaskMod = false;
   bool IsWorksharingReduction = false;
+  bool HasNowaitClause = false;
 };
 
 /// Class intended to support codegen of all kind of the reduction clauses.
 class ReductionCodeGen {
 private:
   /// Data required for codegen of reduction clauses.
   struct ReductionData {
     /// Reference to the item shared between tasks to reduce into.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

if (!Data.Dependences.empty()) {
                           TaskArgs);
     }
     // Check if parent region is untied and build return for untied task;
     if (auto *Region =
             dyn_cast_or_null<CGOpenMPRegionInfo>(CGF.CapturedStmtInfo))
       Region->emitUntiedSwitch(CGF);
   };
 
-  llvm::Value *DepWaitTaskArgs[6];
+  llvm::Value *DepWaitTaskArgs[7];
   if (!Data.Dependences.empty()) {
     DepWaitTaskArgs[0] = UpLoc;
     DepWaitTaskArgs[1] = ThreadID;
     DepWaitTaskArgs[2] = NumOfElements;
     DepWaitTaskArgs[3] = DependenciesArray.getPointer();
     DepWaitTaskArgs[4] = CGF.Builder.getInt32(0);
     DepWaitTaskArgs[5] = llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
+    DepWaitTaskArgs[6] =
+        llvm::ConstantInt::get(CGF.Int32Ty, Data.HasNowaitClause);
   }
   auto &M = CGM.getModule();
   auto &&ElseCodeGen = [this, &M, &TaskArgs, ThreadID, NewTaskNewTaskTTy,
                         TaskEntry, &Data, &DepWaitTaskArgs,
                         Loc](CodeGenFunction &CGF, PrePostActionTy &) {
     CodeGenFunction::RunCleanupsScope LocalScope(CGF);
     // Build void __kmpc_omp_wait_deps(ident_t *, kmp_int32 gtid,
     // kmp_int32 ndeps, kmp_depend_info_t *dep_list, kmp_int32
     // ndeps_noalias, kmp_depend_info_t *noalias_dep_list); if dependence info
     // is specified.
     if (!Data.Dependences.empty())
-      CGF.EmitRuntimeCall(
-          OMPBuilder.getOrCreateRuntimeFunction(M, OMPRTL___kmpc_omp_wait_deps),
+      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
+                              M, OMPRTL___kmpc_omp_taskwait_deps_51),
                           DepWaitTaskArgs);
     // Call proxy_task_entry(gtid, new_task);
     auto &&CodeGen = [TaskEntry, ThreadID, NewTaskNewTaskTTy,
                       Loc](CodeGenFunction &CGF, PrePostActionTy &Action) {
       Action.Enter(CGF);
       llvm::Value *OutlinedFnArgs[] = {ThreadID, NewTaskNewTaskTTy};
       CGF.CGM.getOpenMPRuntime().emitOutlinedFunctionCall(CGF, Loc, TaskEntry,
                                                           OutlinedFnArgs);
     };
return Address(
           OMPBuilder.getOrCreateRuntimeFunction(
               CGM.getModule(), OMPRTL___kmpc_task_reduction_get_th_data),
           Args),
       CGF.Int8Ty, SharedLVal.getAlignment());
 }
 
 void CGOpenMPRuntime::emitTaskwaitCall(CodeGenFunction &CGF, SourceLocation Loc,
                                        const OMPTaskDataTy &Data) {
   if (!CGF.HaveInsertPoint())
     return;
 
   if (CGF.CGM.getLangOpts().OpenMPIRBuilder && Data.Dependences.empty()) {
     // TODO: Need to support taskwait with dependences in the OpenMPIRBuilder.
     OMPBuilder.createTaskwait(CGF.Builder);
   } else {
     llvm::Value *ThreadID = getThreadID(CGF, Loc);
     llvm::Value *UpLoc = emitUpdateLocation(CGF, Loc);
     auto &M = CGM.getModule();
     Address DependenciesArray = Address::invalid();
     llvm::Value *NumOfElements;
     std::tie(NumOfElements, DependenciesArray) =
         emitDependClause(CGF, Data.Dependences, Loc);
-    llvm::Value *DepWaitTaskArgs[6];
     if (!Data.Dependences.empty()) {
+      llvm::Value *DepWaitTaskArgs[7];
       DepWaitTaskArgs[0] = UpLoc;
       DepWaitTaskArgs[1] = ThreadID;
       DepWaitTaskArgs[2] = NumOfElements;
       DepWaitTaskArgs[3] = DependenciesArray.getPointer();
       DepWaitTaskArgs[4] = CGF.Builder.getInt32(0);
       DepWaitTaskArgs[5] = llvm::ConstantPointerNull::get(CGF.VoidPtrTy);
+      DepWaitTaskArgs[6] =
+          llvm::ConstantInt::get(CGF.Int32Ty, Data.HasNowaitClause);
 
       CodeGenFunction::RunCleanupsScope LocalScope(CGF);
 
-      // Build void __kmpc_omp_wait_deps(ident_t *, kmp_int32 gtid,
+      // Build void __kmpc_omp_taskwait_deps_51(ident_t *, kmp_int32 gtid,
       // kmp_int32 ndeps, kmp_depend_info_t *dep_list, kmp_int32
-      // ndeps_noalias, kmp_depend_info_t *noalias_dep_list); if dependence info
-      // is specified.
-      CGF.EmitRuntimeCall(
-          OMPBuilder.getOrCreateRuntimeFunction(M, OMPRTL___kmpc_omp_wait_deps),
+      // ndeps_noalias, kmp_depend_info_t *noalias_dep_list,
+      // kmp_int32 has_no_wait); if dependence info is specified.
+      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
+                              M, OMPRTL___kmpc_omp_taskwait_deps_51),
                           DepWaitTaskArgs);
 
     } else {
 
       // Build call kmp_int32 __kmpc_omp_taskwait(ident_t *loc, kmp_int32
       // global_tid);
       llvm::Value *Args[] = {UpLoc, ThreadID};
       // Ignore return result until untied tasks are supported.
       CGF.EmitRuntimeCall(
           OMPBuilder.getOrCreateRuntimeFunction(M, OMPRTL___kmpc_omp_taskwait),
           Args);
     }
   }
 
   if (auto *Region = dyn_cast_or_null<CGOpenMPRegionInfo>(CGF.CapturedStmtInfo))
     Region->emitUntiedSwitch(CGF);
 }
 
 void CGOpenMPRuntime::emitInlinedDirective(CodeGenFunction &CGF,
                                            OpenMPDirectiveKind InnerKind,
                                            const RegionCodeGenTy &CodeGen,
                                            bool HasCancel) {
   if (!CGF.HaveInsertPoint())
     return;
   InlinedOpenMPRegionRAII Region(CGF, CodeGen, InnerKind, HasCancel,
                                  InnerKind != OMPD_critical &&
                                      InnerKind != OMPD_master &&
                                      InnerKind != OMPD_masked);
   CGF.CapturedStmtInfo->EmitBody(CGF, /*S=*/nullptr);
 }
 
 namespace {
 enum RTCancelKind {
   CancelNoreq = 0,
   CancelParallel = 1,
   CancelLoop = 2,
   CancelSections = 3,

clang/lib/CodeGen/CGStmtOpenMP.cpp

 void CodeGenFunction::EmitOMPBarrierDirective(const OMPBarrierDirective &S) {
   CGM.getOpenMPRuntime().emitBarrierCall(*this, S.getBeginLoc(), OMPD_barrier);
 }
 
 void CodeGenFunction::EmitOMPTaskwaitDirective(const OMPTaskwaitDirective &S) {
   OMPTaskDataTy Data;
   // Build list of dependences
   buildDependences(S, Data);
+  Data.HasNowaitClause = S.hasClausesOfKind<OMPNowaitClause>();
   CGM.getOpenMPRuntime().emitTaskwaitCall(*this, S.getBeginLoc(), Data);
 }
 
 bool isSupportedByOpenMPIRBuilder(const OMPTaskgroupDirective &T) {
   return T.clauses().empty();
 }
 
 void CodeGenFunction::EmitOMPTaskgroupDirective(
     const OMPTaskgroupDirective &S) {

clang/lib/Sema/SemaOpenMP.cpp

if (!SeverityC || SeverityC->getSeverityKind() != OMPC_SEVERITY_warning)
       return StmtError();
   }
   return OMPErrorDirective::Create(Context, StartLoc, EndLoc, Clauses);
 }
 
 StmtResult Sema::ActOnOpenMPTaskwaitDirective(ArrayRef<OMPClause *> Clauses,
                                               SourceLocation StartLoc,
                                               SourceLocation EndLoc) {
+  const OMPNowaitClause *NowaitC =
+      OMPExecutableDirective::getSingleClause<OMPNowaitClause>(Clauses);
+  const OMPDependClause *DependC =
+      OMPExecutableDirective::getSingleClause<OMPDependClause>(Clauses);
+  if (NowaitC && !DependC) {
+    Diag(StartLoc, diag::err_omp_nowait_clause_without_depend);
+    return StmtError();
+  }
+
   return OMPTaskwaitDirective::Create(Context, StartLoc, EndLoc, Clauses);
 }
 
 StmtResult Sema::ActOnOpenMPTaskgroupDirective(ArrayRef<OMPClause *> Clauses,
                                                Stmt *AStmt,
                                                SourceLocation StartLoc,
                                                SourceLocation EndLoc) {
   if (!AStmt)

clang/test/OpenMP/target_depend_codegen.cpp

 // CHECK:       [[DEV:%.+]] = load i32, i32* [[DEVICE_CAP]],
 // CHECK:       store i32 [[DEV]], i32* [[GEP]],
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START:%.+]], i[[SZ]] 1
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 2
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP_START]], i[[SZ]] 3
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* [[DEP_START]] to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target device(global + a) depend(in                   \
                                              : global) depend(out \
                                                               : a, b, cn[4])
   {
   }
#pragma omp target device(global + a) nowait depend(inout                                          \
     static int local1;
     *plocal = global;
     local1 = global;
   }
 
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target if (0) firstprivate(global) depend(out \
                                                       : global)
   {
     global += 1;
   }

clang/test/OpenMP/target_enter_data_depend_codegen.cpp

void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 2
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY2]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
 
   // CK1: br label %[[IFEND:[^,]+]]
 
   // CK1: [[IFELSE]]
   // CK1: br label %[[IFEND]]
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 3
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY3]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target enter data map(alloc:lb) depend(out: lb, arg, la, gc)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 4
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] [[BC_ADR]], i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 4, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 1, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target enter data map(to:gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
   {++arg;}
 }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias noundef %1)

clang/test/OpenMP/target_exit_data_depend_codegen.cpp

void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 2
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY2]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
 
   // CK1: br label %[[IFEND:[^,]+]]
 
   // CK1: [[IFELSE]]
   // CK1: br label %[[IFEND]]
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 3
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY3]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target exit data map(from:lb) depend(out: lb, arg, la, gc)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 4
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] [[BC_ADR]], i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 4, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 1, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target exit data map(from:gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
   {++arg;}
 }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias noundef %1)

clang/test/OpenMP/target_parallel_depend_codegen.cpp

 // CHECK:       store i32 [[DEV]], i32* [[GEP]],
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target parallel device(global + a) depend(in                   \
                                                       : global) depend(out \
                                                                        : a, b, cn[4])
   {
   }
#pragma omp target parallel device(global + a) nowait depend(inout                       \
     *plocal = global;
     local1 = global;
   }
 
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target parallel if (0) firstprivate(global) depend(out \
                                                                : global)
   {
     global += 1;
   }

clang/test/OpenMP/target_parallel_for_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[IN:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[IN]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[IN]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target parallel for device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target parallel for if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_parallel_for_simd_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[IN:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[IN]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[IN]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target parallel for simd device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[IN]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[IN]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target parallel for simd if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_simd_depend_codegen.cpp

 // CHECK:       store i32 [[DEV]], i32* [[GEP]],
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target simd device(global + a) depend(in                   \
                                                   : global) depend(out \
                                                                    : a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target simd if (0) firstprivate(global) depend(out \
                                                            : global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }

clang/test/OpenMP/target_teams_depend_codegen.cpp

 // CHECK:       store i32 [[DEV]], i32* [[GEP]],
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target teams device(global + a) depend(in                   \
                                                    : global) depend(out \
                                                                     : a, b, cn[4])
   {
   }
#pragma omp target teams device(global + a) nowait depend(inout \
     *plocal = global;
     local1 = global;
   }
 
 // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @1, i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
 // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
 // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
 // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-// CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+// CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @1, i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
 // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
 // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @1, i32 [[GTID]], i8* [[TASK]])
 #pragma omp target teams if (0) firstprivate(global) depend(out \
                                                             : global)
   {
     global += 1;
   }

clang/test/OpenMP/target_teams_distribute_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_teams_distribute_parallel_for_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute parallel for device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute parallel for if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute parallel for simd device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute parallel for simd if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_teams_distribute_simd_depend_codegen.cpp

int foo(int n) {
   // CHECK:       store i32 [[DEV]], i32* [[GEP]],
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID:@.+]], i32 [[GTID:%.+]], i32 1, i[[SZ]] {{20|40}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY0:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY0:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 1
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 2
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 3
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 4, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY0]](i32 [[GTID]], [[TASK_TY0]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute simd device(global + a) depend(in: global) depend(out: a, b, cn[4])
   for (int i = 0; i < 10; ++i) {
   }
 
   // CHECK:       [[ADD:%.+]] = add nsw i32
for (int i = 0; i < *plocal; ++i) {
     *plocal = global;
     local1 = global;
   }
 
   // CHECK:       [[TASK:%.+]] = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i[[SZ]] {{48|24}}, i[[SZ]] 4, i32 (i32, i8*)* bitcast (i32 (i32, %{{.+}}*)* [[TASK_ENTRY2:@.+]] to i32 (i32, i8*)*))
   // CHECK:       [[BC_TASK:%.+]] = bitcast i8* [[TASK]] to [[TASK_TY2:%.+]]*
   // CHECK:       getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* %{{.+}}, i[[SZ]] 0
   // CHECK:       [[DEP:%.+]] = bitcast %struct.kmp_depend_info* %{{.+}} to i8*
-  // CHECK:       call void @__kmpc_omp_wait_deps(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null)
+  // CHECK:       call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* [[ID]], i32 [[GTID]], i32 1, i8* [[DEP]], i32 0, i8* null, i32 0)
   // CHECK:       call void @__kmpc_omp_task_begin_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   // CHECK:       call i32 [[TASK_ENTRY2]](i32 [[GTID]], [[TASK_TY2]]* [[BC_TASK]])
   // CHECK:       call void @__kmpc_omp_task_complete_if0(%struct.ident_t* [[ID]], i32 [[GTID]], i8* [[TASK]])
   #pragma omp target teams distribute simd if(0) firstprivate(global) depend(out:global)
   for (int i = 0; i < global; ++i) {
     global += 1;
   }
 

clang/test/OpenMP/target_update_depend_codegen.cpp

void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 2
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 3, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY2]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
 
   // CK1: br label %[[IFEND:[^,]+]]
 
   // CK1: [[IFELSE]]
   // CK1: br label %[[IFEND]]
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 3
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] ptrtoint ([100 x double]* @gc to i[[sz]]), i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 800, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 3, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 4, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY3]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target update from(lb) depend(out: lb, arg, la, gc)
   {++arg;}
 
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   {++arg;}
void foo(int arg) {
   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, %struct.kmp_depend_info* [[MAIN_DEP]], i[[sz]] 4
   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 0
   // CK1: store i[[sz]] [[BC_ADR]], i[[sz]]* [[DEP_ADR]],
   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 1
   // CK1: store i[[sz]] 4, i[[sz]]* [[DEP_SIZE]],
   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds %struct.kmp_depend_info, %struct.kmp_depend_info* [[DEP]], i32 0, i32 2
   // CK1: store i8 1, i8* [[DEP_ATTRS]]
   // CK1: [[BC:%.+]] = bitcast %struct.kmp_depend_info* [[MAIN_DEP]] to i8*
-  // CK1: call void @__kmpc_omp_wait_deps(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null)
+  // CK1: call void @__kmpc_omp_taskwait_deps_51(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 5, i8* [[BC]], i32 0, i8* null, i32 0)
   // CK1: call void @__kmpc_omp_task_begin_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, %struct.kmp_task_t_with_privates{{.+}}* [[RES_BC]])
   // CK1: call void @__kmpc_omp_task_complete_if0(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i8* [[RES]])
   #pragma omp target update to(gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
   {++arg;}
 }
 
 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, %struct.kmp_task_t_with_privates* noalias noundef %1)

clang/test/OpenMP/task_if_codegen.cpp

 // CHECK1-NEXT:    store i64 1, ptr [[DEP_COUNTER_ADDR]], align 8
 // CHECK1-NEXT:    [[TMP24:%.*]] = load i32, ptr @Arg, align 4
 // CHECK1-NEXT:    [[TOBOOL4:%.*]] = icmp ne i32 [[TMP24]], 0
 // CHECK1-NEXT:    br i1 [[TOBOOL4]], label [[OMP_IF_THEN5:%.*]], label [[OMP_IF_ELSE6:%.*]]
 // CHECK1:       omp_if.then5:
 // CHECK1-NEXT:    [[TMP25:%.*]] = call i32 @__kmpc_omp_task_with_deps(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]], i32 1, ptr [[TMP18]], i32 0, ptr null)
 // CHECK1-NEXT:    br label [[OMP_IF_END7:%.*]]
 // CHECK1:       omp_if.else6:
-// CHECK1-NEXT:    call void @__kmpc_omp_wait_deps(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP18]], i32 0, ptr null)
+// CHECK1-NEXT:    call void @__kmpc_omp_taskwait_deps_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP18]], i32 0, ptr null, i32 0)
 // CHECK1-NEXT:    call void @__kmpc_omp_task_begin_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]])
 // CHECK1-NEXT:    [[TMP26:%.*]] = call i32 @.omp_task_entry..9(i32 [[TMP0]], ptr [[TMP15]]) #[[ATTR3]]
 // CHECK1-NEXT:    call void @__kmpc_omp_task_complete_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]])
 // CHECK1-NEXT:    br label [[OMP_IF_END7]]
 // CHECK1:       omp_if.end7:
 // CHECK1-NEXT:    [[TMP27:%.*]] = load i32, ptr @Arg, align 4
 // CHECK1-NEXT:    [[CALL:%.*]] = call noundef i32 @_Z5tmainIiEiT_(i32 noundef [[TMP27]])
 // CHECK1-NEXT:    ret i32 [[CALL]]

 // CHECK1-NEXT:    store i64 1, ptr [[DEP_COUNTER_ADDR]], align 8
 // CHECK1-NEXT:    [[TMP25:%.*]] = load i32, ptr [[ARG_ADDR]], align 4
 // CHECK1-NEXT:    [[TOBOOL4:%.*]] = icmp ne i32 [[TMP25]], 0
 // CHECK1-NEXT:    br i1 [[TOBOOL4]], label [[OMP_IF_THEN5:%.*]], label [[OMP_IF_ELSE6:%.*]]
 // CHECK1:       omp_if.then5:
 // CHECK1-NEXT:    [[TMP26:%.*]] = call i32 @__kmpc_omp_task_with_deps(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]], i32 1, ptr [[TMP18]], i32 0, ptr null)
 // CHECK1-NEXT:    br label [[OMP_IF_END7:%.*]]
 // CHECK1:       omp_if.else6:
-// CHECK1-NEXT:    call void @__kmpc_omp_wait_deps(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP18]], i32 0, ptr null)
+// CHECK1-NEXT:    call void @__kmpc_omp_taskwait_deps_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP18]], i32 0, ptr null, i32 0)
 // CHECK1-NEXT:    call void @__kmpc_omp_task_begin_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]])
 // CHECK1-NEXT:    [[TMP27:%.*]] = call i32 @.omp_task_entry..17(i32 [[TMP0]], ptr [[TMP15]]) #[[ATTR3]]
 // CHECK1-NEXT:    call void @__kmpc_omp_task_complete_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP15]])
 // CHECK1-NEXT:    br label [[OMP_IF_END7]]
 // CHECK1:       omp_if.end7:
 // CHECK1-NEXT:    [[TMP28:%.*]] = call ptr @__kmpc_omp_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i64 40, i64 1, ptr @.omp_task_entry..19)
 // CHECK1-NEXT:    [[TMP30:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES_17:%.*]], ptr [[TMP28]], i32 0, i32 0
 // CHECK1-NEXT:    [[TMP31:%.*]] = getelementptr inbounds [1 x %struct.kmp_depend_info], ptr [[DOTDEP_ARR_ADDR9]], i64 0, i64 0
 // CHECK1-NEXT:    [[TMP32:%.*]] = ptrtoint ptr [[ARG_ADDR]] to i64
 // CHECK1-NEXT:    [[TMP33:%.*]] = getelementptr [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP31]], i64 0
 // CHECK1-NEXT:    [[TMP34:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP33]], i32 0, i32 0
 // CHECK1-NEXT:    store i64 [[TMP32]], ptr [[TMP34]], align 8
 // CHECK1-NEXT:    [[TMP35:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP33]], i32 0, i32 1
 // CHECK1-NEXT:    store i64 4, ptr [[TMP35]], align 8
 // CHECK1-NEXT:    [[TMP36:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP33]], i32 0, i32 2
 // CHECK1-NEXT:    store i8 3, ptr [[TMP36]], align 8
 // CHECK1-NEXT:    store i64 1, ptr [[DEP_COUNTER_ADDR10]], align 8
 // CHECK1-NEXT:    [[TMP38:%.*]] = load i32, ptr [[ARG_ADDR]], align 4
 // CHECK1-NEXT:    [[TOBOOL11:%.*]] = icmp ne i32 [[TMP38]], 0
 // CHECK1-NEXT:    br i1 [[TOBOOL11]], label [[OMP_IF_THEN12:%.*]], label [[OMP_IF_ELSE13:%.*]]
 // CHECK1:       omp_if.then12:
 // CHECK1-NEXT:    [[TMP39:%.*]] = call i32 @__kmpc_omp_task_with_deps(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP28]], i32 1, ptr [[TMP31]], i32 0, ptr null)
 // CHECK1-NEXT:    br label [[OMP_IF_END14:%.*]]
 // CHECK1:       omp_if.else13:
-// CHECK1-NEXT:    call void @__kmpc_omp_wait_deps(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP31]], i32 0, ptr null)
+// CHECK1-NEXT:    call void @__kmpc_omp_taskwait_deps_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP31]], i32 0, ptr null, i32 0)
 // CHECK1-NEXT:    call void @__kmpc_omp_task_begin_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP28]])
 // CHECK1-NEXT:    [[TMP40:%.*]] = call i32 @.omp_task_entry..19(i32 [[TMP0]], ptr [[TMP28]]) #[[ATTR3]]
 // CHECK1-NEXT:    call void @__kmpc_omp_task_complete_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP28]])
 // CHECK1-NEXT:    br label [[OMP_IF_END14]]
 // CHECK1:       omp_if.end14:
 // CHECK1-NEXT:    [[TMP41:%.*]] = call ptr @__kmpc_omp_task_alloc(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, i64 40, i64 1, ptr @.omp_task_entry..21)
 // CHECK1-NEXT:    [[TMP43:%.*]] = getelementptr inbounds [[STRUCT_KMP_TASK_T_WITH_PRIVATES_19:%.*]], ptr [[TMP41]], i32 0, i32 0
 // CHECK1-NEXT:    [[TMP44:%.*]] = getelementptr inbounds [1 x %struct.kmp_depend_info], ptr [[DOTDEP_ARR_ADDR16]], i64 0, i64 0
 // CHECK1-NEXT:    [[TMP45:%.*]] = ptrtoint ptr [[ARG_ADDR]] to i64
 // CHECK1-NEXT:    [[TMP46:%.*]] = getelementptr [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP44]], i64 0
 // CHECK1-NEXT:    [[TMP47:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP46]], i32 0, i32 0
 // CHECK1-NEXT:    store i64 [[TMP45]], ptr [[TMP47]], align 8
 // CHECK1-NEXT:    [[TMP48:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP46]], i32 0, i32 1
 // CHECK1-NEXT:    store i64 4, ptr [[TMP48]], align 8
 // CHECK1-NEXT:    [[TMP49:%.*]] = getelementptr inbounds [[STRUCT_KMP_DEPEND_INFO]], ptr [[TMP46]], i32 0, i32 2
 // CHECK1-NEXT:    store i8 3, ptr [[TMP49]], align 8
 // CHECK1-NEXT:    store i64 1, ptr [[DEP_COUNTER_ADDR17]], align 8
 // CHECK1-NEXT:    [[TMP51:%.*]] = load i32, ptr [[ARG_ADDR]], align 4
 // CHECK1-NEXT:    [[TOBOOL18:%.*]] = icmp ne i32 [[TMP51]], 0
 // CHECK1-NEXT:    br i1 [[TOBOOL18]], label [[OMP_IF_THEN19:%.*]], label [[OMP_IF_ELSE20:%.*]]
 // CHECK1:       omp_if.then19:
 // CHECK1-NEXT:    [[TMP52:%.*]] = call i32 @__kmpc_omp_task_with_deps(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP41]], i32 1, ptr [[TMP44]], i32 0, ptr null)
 // CHECK1-NEXT:    br label [[OMP_IF_END21:%.*]]
 // CHECK1:       omp_if.else20:
-// CHECK1-NEXT:    call void @__kmpc_omp_wait_deps(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP44]], i32 0, ptr null)
+// CHECK1-NEXT:    call void @__kmpc_omp_taskwait_deps_51(ptr @[[GLOB1]], i32 [[TMP0]], i32 1, ptr [[TMP44]], i32 0, ptr null, i32 0)
 // CHECK1-NEXT:    call void @__kmpc_omp_task_begin_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP41]])
 // CHECK1-NEXT:    [[TMP53:%.*]] = call i32 @.omp_task_entry..21(i32 [[TMP0]], ptr [[TMP41]]) #[[ATTR3]]
 // CHECK1-NEXT:    call void @__kmpc_omp_task_complete_if0(ptr @[[GLOB1]], i32 [[TMP0]], ptr [[TMP41]])
 // CHECK1-NEXT:    br label [[OMP_IF_END21]]
 // CHECK1:       omp_if.end21:
 // CHECK1-NEXT:    ret i32 0
 //
 //

clang/test/OpenMP/taskwait_ast_print.cpp

-// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
-// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
-// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void foo() {}
 
 template <class T>
 T tmain(T argc) {
   static T a;
 #pragma omp taskwait
 #pragma omp taskwait depend(in:a, argc)
   return a + argc;
 }
+
+template <class T>
+T ndmain(T argc) {
+  static T a;
+#pragma omp taskwait
+#pragma omp taskwait depend(in:a, argc) nowait
+  return a + argc;
+}
+
 // CHECK:      static T a;
 // CHECK-NEXT: #pragma omp taskwait{{$}}
 // CHECK-NEXT: #pragma omp taskwait depend(in : a,argc){{$}}
 // CHECK:      static int a;
 // CHECK-NEXT: #pragma omp taskwait
 // CHECK-NEXT: #pragma omp taskwait depend(in : a,argc){{$}}
 // CHECK:      static char a;
 // CHECK-NEXT: #pragma omp taskwait
 // CHECK-NEXT: #pragma omp taskwait depend(in : a,argc){{$}}
 
+// CHECK:      static T a;
+// CHECK-NEXT: #pragma omp taskwait{{$}}
+// CHECK-NEXT: #pragma omp taskwait depend(in : a,argc) nowait{{$}}
+// CHECK:      static int a;
+// CHECK-NEXT: #pragma omp taskwait
+// CHECK-NEXT: #pragma omp taskwait depend(in : a,argc) nowait{{$}}
+// CHECK:      static char a;
+// CHECK-NEXT: #pragma omp taskwait
+// CHECK-NEXT: #pragma omp taskwait depend(in : a,argc) nowait{{$}}
+
 int main(int argc, char **argv) {
   static int a;
 // CHECK: static int a;
 #pragma omp taskwait
 #pragma omp taskwait depend(out:a, argc)
   // CHECK-NEXT: #pragma omp taskwait
   // CHECK-NEXT: #pragma omp taskwait depend(out : a,argc)
-  return tmain(argc) + tmain(argv[0][0]) + a;
+  return tmain(argc) + tmain(argv[0][0]) + a + ndmain(argc) + ndmain(argv[0][0]);
 }
 
 #endif

clang/test/OpenMP/taskwait_depend_codegen.cpp

 
 // CHECK-LABEL: @main
 // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(ptr @{{.+}})
 // CHECK: [[ALLOC:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr @{{.+}}, i32 [[GTID]], i32 1, i64 40, i64 16, ptr @{{.+}})
 // CHECK: %{{.+}} = call i32 @__kmpc_omp_task_with_deps(ptr @{{.+}}, i32 [[GTID]], ptr [[ALLOC]], i32 1, ptr %{{[0-9]*}}, i32 0, ptr null)
 
 // CHECK-LABEL: tmain
 // CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(ptr @{{.+}})
-// CHECK: call void @__kmpc_omp_wait_deps(ptr @{{.+}}, i32 [[GTID]], i32 1, ptr %{{.}}, i32 0, ptr null)
+// CHECK: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 [[GTID]], i32 1, ptr %{{.}}, i32 0, ptr null, i32 0)
 
 
 #endif

clang/test/OpenMP/taskwait_depend_nowait_codegen.cpp

+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp -fopenmp-version=51 -x c++ -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-apple-darwin10 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -x c++ -triple x86_64-apple-darwin10 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp-simd -fopenmp-version=51 -x c++ -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-apple-darwin10 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple x86_64-apple-darwin10 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+void foo() {}
+
+template <class T>
+T nmain(T &argc) {
+  static T a;
+  #pragma omp taskwait depend(in:argc) nowait
+  return a + argc;
+}
+int main(int argc, char **argv) {
+  int n = 0;
+  n = argc;
+  return nmain(n);
+}
+
+// CHECK-LABEL: @main{{.*}}
+
+// CHECK-LABEL: nmain
+// CHECK: [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(ptr @{{.+}})
+// CHECK: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 [[GTID]], i32 1, ptr %{{.}}, i32 0, ptr null, i32 1)
+
+
+#endif

llvm/include/llvm/Frontend/OpenMP/OMP.td

def OMP_Error : Directive<"error"> {
   let allowedClauses = [
     VersionedClause<OMPC_At, 51>,
     VersionedClause<OMPC_Severity, 51>,
     VersionedClause<OMPC_Message, 51>
   ];
 }
 def OMP_TaskWait : Directive<"taskwait"> {
   let allowedClauses = [
-    VersionedClause<OMPC_Depend, 50>
+    VersionedClause<OMPC_Depend, 50>,
+    VersionedClause<OMPC_NoWait, 51>
   ];
 }
 def OMP_TaskGroup : Directive<"taskgroup"> {
   let allowedClauses = [
     VersionedClause<OMPC_TaskReduction, 50>,
     VersionedClause<OMPC_Allocate, 50>
   ];
 }

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

__OMP_RTL(__kmpc_task_reduction_get_th_data, false, VoidPtr, Int32, VoidPtr,
           VoidPtr)
 __OMP_RTL(__kmpc_task_reduction_init, false, VoidPtr, Int32, Int32, VoidPtr)
 __OMP_RTL(__kmpc_task_reduction_modifier_init, false, VoidPtr, VoidPtr, Int32,
           Int32, Int32, VoidPtr)
 __OMP_RTL(__kmpc_proxy_task_completed_ooo, false, Void, VoidPtr)
 
 __OMP_RTL(__kmpc_omp_wait_deps, false, Void, IdentPtr, Int32, Int32,
           /* kmp_depend_info_t */ VoidPtr, Int32, VoidPtr)
+__OMP_RTL(__kmpc_omp_taskwait_deps_51, false, Void, IdentPtr, Int32, Int32,
+          /* kmp_depend_info_t */ VoidPtr, Int32, VoidPtr, Int32)
 __OMP_RTL(__kmpc_cancellationpoint, false, Int32, IdentPtr, Int32, Int32)
 
 __OMP_RTL(__kmpc_fork_teams, true, Void, IdentPtr, Int32, ParallelTaskPtr)
 __OMP_RTL(__kmpc_push_num_teams, false, Void, IdentPtr, Int32, Int32, Int32)
 
 __OMP_RTL(__kmpc_copyprivate, false, Void, IdentPtr, Int32, SizeTy, VoidPtr,
           CopyFunctionPtr, Int32)
 __OMP_RTL(__kmpc_threadprivate_cached, false, VoidPtr, IdentPtr, Int32, VoidPtr,

 __OMP_RTL_ATTRS(__kmpc_task_reduction_modifier_init, DefaultAttrs,
                 ReturnPtrAttrs, ParamAttrs())
 __OMP_RTL_ATTRS(__kmpc_proxy_task_completed_ooo, DefaultAttrs, AttributeSet(),
                 ParamAttrs())
 
 __OMP_RTL_ATTRS(__kmpc_omp_wait_deps, BarrierAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), AttributeSet(),
                            ReadOnlyPtrAttrs))
+__OMP_RTL_ATTRS(__kmpc_omp_taskwait_deps_51, BarrierAttrs, AttributeSet(),
+                ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), AttributeSet(),
+                           ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_cancellationpoint, DefaultAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs))
 
 __OMP_RTL_ATTRS(__kmpc_fork_teams, ForkAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_push_num_teams, InaccessibleArgOnlyAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs))
 

llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp

 
 void OpenMPIRBuilder::createFlush(const LocationDescription &Loc) {
   if (!updateToLocation(Loc))
     return;
   emitFlush(Loc);
 }
 
 void OpenMPIRBuilder::emitTaskwaitImpl(const LocationDescription &Loc) {
   // Build call kmp_int32 __kmpc_omp_taskwait(ident_t *loc, kmp_int32
   // global_tid);
   uint32_t SrcLocStrSize;
   Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
   Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
   Value *Args[] = {Ident, getOrCreateThreadID(Ident)};
 
   // Ignore return result until untied tasks are supported.
   Builder.CreateCall(getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_taskwait),
                      Args);
 }
 
 void OpenMPIRBuilder::createTaskwait(const LocationDescription &Loc) {
   if (!updateToLocation(Loc))
     return;
   emitTaskwaitImpl(Loc);
 }
 
 void OpenMPIRBuilder::emitTaskyieldImpl(const LocationDescription &Loc) {
   // Build call __kmpc_omp_taskyield(loc, thread_id, 0);
   uint32_t SrcLocStrSize;
   Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
   Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
   Constant *I32Null = ConstantInt::getNullValue(Int32);
   Value *Args[] = {Ident, getOrCreateThreadID(Ident), I32Null};
 
   Builder.CreateCall(getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_taskyield),
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::applyWorkshareLoop(
   case OMPScheduleType::BaseAuto:
   case OMPScheduleType::BaseGreedy:
   case OMPScheduleType::BaseBalanced:
   case OMPScheduleType::BaseSteal:
   case OMPScheduleType::BaseGuidedSimd:
   case OMPScheduleType::BaseRuntimeSimd:
     assert(!ChunkSize &&
            "schedule type does not support user-defined chunk sizes");
-    [[fallthrough]];
+    LLVM_FALLTHROUGH;
   case OMPScheduleType::BaseDynamicChunked:
   case OMPScheduleType::BaseGuidedChunked:
   case OMPScheduleType::BaseGuidedIterativeChunked:
   case OMPScheduleType::BaseGuidedAnalyticalChunked:
   case OMPScheduleType::BaseStaticBalancedChunked:
     return applyDynamicWorkshareLoop(DL, CLI, AllocaIP, EffectiveScheduleType,
                                      NeedsBarrier, ChunkSize);
 

openmp/runtime/src/dllexports

 %ifndef stub
         __kmpc_push_proc_bind               237
         __kmpc_taskgroup                    238
         __kmpc_end_taskgroup                239
         __kmpc_push_num_teams               240
         __kmpc_fork_teams                   241
         __kmpc_omp_task_with_deps           242
         __kmpc_omp_wait_deps                243
+        __kmpc_omp_taskwait_deps_51         292
         __kmpc_cancel                       244
         __kmpc_cancellationpoint            245
         __kmpc_cancel_barrier               246
         __kmpc_dist_for_static_init_4       247
         __kmpc_dist_for_static_init_4u      248
         __kmpc_dist_for_static_init_8       249
         __kmpc_dist_for_static_init_8u      250
         __kmpc_dist_dispatch_init_4         251

openmp/runtime/src/kmp.h

KMP_EXPORT kmp_task_t *__kmpc_omp_target_task_alloc(
     size_t sizeof_shareds, kmp_routine_entry_t task_entry, kmp_int64 device_id);
 KMP_EXPORT void __kmpc_omp_task_begin_if0(ident_t *loc_ref, kmp_int32 gtid,
                                           kmp_task_t *task);
 KMP_EXPORT void __kmpc_omp_task_complete_if0(ident_t *loc_ref, kmp_int32 gtid,
                                              kmp_task_t *task);
 KMP_EXPORT kmp_int32 __kmpc_omp_task_parts(ident_t *loc_ref, kmp_int32 gtid,
                                            kmp_task_t *new_task);
 KMP_EXPORT kmp_int32 __kmpc_omp_taskwait(ident_t *loc_ref, kmp_int32 gtid);
-
 KMP_EXPORT kmp_int32 __kmpc_omp_taskyield(ident_t *loc_ref, kmp_int32 gtid,
                                           int end_part);
 
 #if TASK_UNUSED
 void __kmpc_omp_task_begin(ident_t *loc_ref, kmp_int32 gtid, kmp_task_t *task);
 void __kmpc_omp_task_complete(ident_t *loc_ref, kmp_int32 gtid,
                               kmp_task_t *task);
 #endif // TASK_UNUSED
 
 /* ------------------------------------------------------------------------ */
 
 KMP_EXPORT void __kmpc_taskgroup(ident_t *loc, int gtid);
 KMP_EXPORT void __kmpc_end_taskgroup(ident_t *loc, int gtid);
 
 KMP_EXPORT kmp_int32 __kmpc_omp_task_with_deps(
     ident_t *loc_ref, kmp_int32 gtid, kmp_task_t *new_task, kmp_int32 ndeps,
     kmp_depend_info_t *dep_list, kmp_int32 ndeps_noalias,
     kmp_depend_info_t *noalias_dep_list);
 KMP_EXPORT void __kmpc_omp_wait_deps(ident_t *loc_ref, kmp_int32 gtid,
                                      kmp_int32 ndeps,
                                      kmp_depend_info_t *dep_list,
                                      kmp_int32 ndeps_noalias,
                                      kmp_depend_info_t *noalias_dep_list);
+/* __kmpc_omp_taskwait_deps_51 : Function for OpenMP 5.1 nowait clause.
+ *                               Placeholder for taskwait with nowait clause.*/
+KMP_EXPORT void __kmpc_omp_taskwait_deps_51(ident_t *loc_ref, kmp_int32 gtid,
+                                            kmp_int32 ndeps,
+                                            kmp_depend_info_t *dep_list,
+                                            kmp_int32 ndeps_noalias,
+                                            kmp_depend_info_t *noalias_dep_list,
+                                            kmp_int32 has_no_wait);
+
 extern kmp_int32 __kmp_omp_task(kmp_int32 gtid, kmp_task_t *new_task,
                                 bool serialize_immediate);
 
 KMP_EXPORT kmp_int32 __kmpc_cancel(ident_t *loc_ref, kmp_int32 gtid,
                                    kmp_int32 cncl_kind);
 KMP_EXPORT kmp_int32 __kmpc_cancellationpoint(ident_t *loc_ref, kmp_int32 gtid,
                                               kmp_int32 cncl_kind);
 KMP_EXPORT kmp_int32 __kmpc_cancel_barrier(ident_t *loc_ref, kmp_int32 gtid);

openmp/runtime/src/kmp_taskdeps.cpp

 @param ndeps_noalias Number of depend items with no aliasing
 @param noalias_dep_list List of depend items with no aliasing
 
 Blocks the current task until all specifies dependences have been fulfilled.
 */
 void __kmpc_omp_wait_deps(ident_t *loc_ref, kmp_int32 gtid, kmp_int32 ndeps,
                           kmp_depend_info_t *dep_list, kmp_int32 ndeps_noalias,
                           kmp_depend_info_t *noalias_dep_list) {
-  KA_TRACE(10, ("__kmpc_omp_wait_deps(enter): T#%d loc=%p\n", gtid, loc_ref));
+  __kmpc_omp_taskwait_deps_51(loc_ref, gtid, ndeps, dep_list, ndeps_noalias,
+                              noalias_dep_list, false);
+}
 
+/* __kmpc_omp_taskwait_deps_51 : Function for OpenMP 5.1 nowait clause.
+                                 Placeholder for taskwait with nowait clause.
+                                 Earlier code of __kmpc_omp_wait_deps() is now
+                                 in this function.
+*/
+void __kmpc_omp_taskwait_deps_51(ident_t *loc_ref, kmp_int32 gtid,
+                                 kmp_int32 ndeps, kmp_depend_info_t *dep_list,
+                                 kmp_int32 ndeps_noalias,
+                                 kmp_depend_info_t *noalias_dep_list,
+                                 kmp_int32 has_no_wait) {
+  KA_TRACE(10, ("__kmpc_omp_taskwait_deps(enter): T#%d loc=%p nowait#%d\n",
+                gtid, loc_ref, has_no_wait));
   if (ndeps == 0 && ndeps_noalias == 0) {
-    KA_TRACE(10, ("__kmpc_omp_wait_deps(exit): T#%d has no dependences to "
+    KA_TRACE(10, ("__kmpc_omp_taskwait_deps(exit): T#%d has no dependences to "
                   "wait upon : loc=%p\n",
                   gtid, loc_ref));
     return;
   }
   __kmp_assert_valid_gtid(gtid);
   kmp_info_t *thread = __kmp_threads[gtid];
   kmp_taskdata_t *current_task = thread->th.th_current_task;
 
bool ignore = current_task->td_flags.team_serial ||
                 current_task->td_flags.final;
   ignore =
       ignore && thread->th.th_task_team != NULL &&
       thread->th.th_task_team->tt.tt_found_proxy_tasks == FALSE &&
       thread->th.th_task_team->tt.tt_hidden_helper_task_encountered == FALSE;
   ignore = ignore || current_task->td_dephash == NULL;
 
   if (ignore) {
-    KA_TRACE(10, ("__kmpc_omp_wait_deps(exit): T#%d has no blocking "
+    KA_TRACE(10, ("__kmpc_omp_taskwait_deps(exit): T#%d has no blocking "
                   "dependences : loc=%p\n",
                   gtid, loc_ref));
 #if OMPT_SUPPORT
     __ompt_taskwait_dep_finish(current_task, taskwait_task_data);
 #endif /* OMPT_SUPPORT */
     return;
   }
 
   kmp_depnode_t node = {0};
   __kmp_init_node(&node);
 
   if (!__kmp_check_deps(gtid, &node, NULL, &current_task->td_dephash,
                         DEP_BARRIER, ndeps, dep_list, ndeps_noalias,
                         noalias_dep_list)) {
-    KA_TRACE(10, ("__kmpc_omp_wait_deps(exit): T#%d has no blocking "
+    KA_TRACE(10, ("__kmpc_omp_taskwait_deps(exit): T#%d has no blocking "
                   "dependences : loc=%p\n",
                   gtid, loc_ref));
 #if OMPT_SUPPORT
     __ompt_taskwait_dep_finish(current_task, taskwait_task_data);
 #endif /* OMPT_SUPPORT */
     return;
   }
 
   int thread_finished = FALSE;
   kmp_flag_32<false, false> flag(
       (std::atomic<kmp_uint32> *)&node.dn.npredecessors, 0U);
   while (node.dn.npredecessors > 0) {
     flag.execute_tasks(thread, gtid, FALSE,
                        &thread_finished USE_ITT_BUILD_ARG(NULL),
                        __kmp_task_stealing_constraint);
   }
 
 #if OMPT_SUPPORT
   __ompt_taskwait_dep_finish(current_task, taskwait_task_data);
 #endif /* OMPT_SUPPORT */
-  KA_TRACE(10, ("__kmpc_omp_wait_deps(exit): T#%d finished waiting : loc=%p\n",
+  KA_TRACE(10, ("__kmpc_omp_taskwait_deps(exit): T#%d finished waiting : loc=%p\
+                \n",
                 gtid, loc_ref));
 }