[OpenMP] Codegen support for thread_limit on target directive

[sandeepkosuri]

Summary

• This patch adds support for thread_limit clause on target directive according to OpenMP 51 [2.14.5]
• The idea is to create an outer task for target region, when there is a thread_limit clause, and manipulate the thread_limit of task instead. This way, thread_limit will be applied to all the relevant constructs enclosed by the target region.

[sandeepkosuri]

Updated target_codegen.cpp test case to incorporate my changes

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9640

What if D is combined target directive, i.e. D.getDirectiveKind() is something like OMPD_target_teams, etc.?

clang/lib/CodeGen/CGStmtOpenMP.cpp

5143–5148

Avoid double call of S.getSingleClause<OMPThreadLimitClause>(), store in local variable call result.

clang/test/OpenMP/target_codegen.cpp

849–850

It requires extra resource consumption, can you try to avoid creating outer task, if possible?

openmp/runtime/src/kmp_ftn_entry.h

813

No need for else here

[sandeepkosuri]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9640

I will fix that, thanks for noticing.

clang/lib/CodeGen/CGStmtOpenMP.cpp

5143–5148

sure.

clang/test/OpenMP/target_codegen.cpp

849–850

I tried different ideas for making thread_limit work on target.

I tried to reuse the existing implementation by replacing the directive to target teams(1) thread_limit(x) at parsing , sema and IR stages. I couldn't successfully implement any of them. So, I tried adding num_threads for all the parallel directives within target, and there were corner cases like parallel directives in a function which is called in target region, which were becoming tedious to handle.

This method seem to encompass the idea of thread limit on target pretty well and also works... So I proceeded with this idea.

openmp/runtime/src/kmp_ftn_entry.h

813

oops, I will fix that

[jdoerfert]

openmp/runtime/test/target/target_thread_limit.cpp

29

This doesn't check much. You need to verify a 5th or print the team size. Same for the rest

36

Verify that the user can use the omp_set_ functions and see consistent behavior.

[sandeepkosuri]

• Added support for thread_limit clause on relevant combined directives which begin with target as per @ABataev 's comments.
• Added additional LIT test cases to check codegen of the thread_limit on the newly supported directives.
• Updated the runtime LIT as per @jdoerfert 's comments.

[sandeepkosuri]

Explicitly mentioned -fopenmp-version=51 in LIT test cases

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9866

I think you don't need needsTaskBasedThreadLimit call here, the emitTargetCall function itself can be called only for target-based directives

clang/lib/CodeGen/CGStmtOpenMP.cpp

5143

Same regarding needsTaskBasedThreadLimit(S.getDirectiveKind()) , the function EmitOMPTargetTaskBasedDirective is called only for target-based directives

[sandeepkosuri]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9866

emitTargetCall() is called for all the target based directives, even target - team based directives, which already have a thread limit implementation in place. So, I need needsTaskBasedThreadLimit to select applicable directives only.

clang/lib/CodeGen/CGStmtOpenMP.cpp

5143

Similarly here as well.

[tianshilei1992]

Is this patch to support thread_limit on target directive on the host?

[sandeepkosuri]

In D152054#4560353, @tianshilei1992 wrote:

Is this patch to support thread_limit on target directive on the host?

Yes @tianshilei1992 , It is for host only

[ABataev]

LG

[sandeepkosuri]

• Updated SemaOpenMP.cpp to support thread_limit clause on the newly allowed directives.

• This update is to fix the newly added LIT tests' failures (which were occurring only on debug build)

[ABataev]

Please, use the script to generate the checks for the newly added tests

[sandeepkosuri]

Used the python script update_cc_test_checks.py to generate the checks for the newly added tests.

[ABataev]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

5

Add PCH serialization/desrialization checks in your tests

[sandeepkosuri]

Added PCH options to the RUN lines in LIT tests

[ABataev]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

31

Why removed these checks?

clang/test/OpenMP/target_parallel_for_tl_codegen.cpp

31

Same, why it was removed?

[sandeepkosuri]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

31

I did not remove any check lines in this function.
But I removed checks in omp_task_entry function that were not related to my changes, to avoid failures. I only wanted to check whether __kmpc_set_thread_limit() is called.

Same for all the other test cases.

[ABataev]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

31

Better to restore it to be able to use the script in future without many changes

[sandeepkosuri]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

31

But a few check lines are failing on windows, while passing on debian.

[ABataev]

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

31

It must be investigated, it is bad idea just to remove these checks. Most probably related to the order of the expressions emission.

[sandeepkosuri]

Edited the LIT test cases to use more script generated check lines.

[ABataev]

LG

[sandeepkosuri]

Used CHECK-DAG s to avoid LIT test failures on Windows system

[sandeepkosuri]

Made LIT test cases more robust to check lines ordering problem

[sandeepkosuri]

made new LIT test cases target specific to linux

[ABataev]

LG

[mstorsjo]

This new test is failing on Windows, due to __kmpc_set_thread_limit not being exported - see e.g. https://github.com/mstorsjo/llvm-mingw/actions/runs/5994183421/job/16264501555. Can someone add it to openmp/runtime/src/dllexports?

[mstorsjo]

In D152054#4620927, @mstorsjo wrote:

This new test is failing on Windows, due to __kmpc_set_thread_limit not being exported - see e.g. https://github.com/mstorsjo/llvm-mingw/actions/runs/5994183421/job/16264501555. Can someone add it to openmp/runtime/src/dllexports?

CC @vadikp-intel @natgla about this. What's the procedure for allocating new ordinal numbers for new exported functions here?

[vadikp-intel]

Windows importing is now done by name, and new exports do not need to have an ordinal specified for them i.e. you can add a line with just the API name to dllexports.

[mstorsjo]

In D152054#4622642, @vadikp-intel wrote:

Windows importing is now done by name, and new exports do not need to have an ordinal specified for them i.e. you can add a line with just the API name to dllexports.

Oh, right, thanks. Will do.

[mstorsjo]

openmp/runtime/test/target/target_thread_limit.cpp

29

This test fails when running (on Windows) on GitHub Actions runners - see https://github.com/mstorsjo/llvm-mingw/actions/runs/6019088705/job/16342540379.

I believe that this bit of the test has got a hidden assumption that it is running in an environment with 4 or more cores. By setting #pragma omp target thread_limit(tl) (with tl=4) and running a line in parallel with #pragma omp parallel, it expects that we'll get 4 printouts - while in practice, we'll get anywhere between 1 and 4 printouts depending on the number of cores.

Is there something that can be done to make this test work in such an environment too?

[mstorsjo]

openmp/runtime/test/target/target_thread_limit.cpp

29

Can someone involved in this patch take on fixing it so that it works on machines with fewer than 4 cores? I'm not sure what's the most appropriate path forward here, as it breaks clearly in such configs (even if it might not be hit by one of the official llvm buildbots, but it shows up as breakage in my nightly builds every day now) - reverting seems a bit harsh. I guess I could just rip out this part of the test?

[sandeepkosuri]

openmp/runtime/test/target/target_thread_limit.cpp

29

@mstorsjo , I noticed that you have committed this https://github.com/llvm/llvm-project/commit/c2019c416c8d7ec50aec6ac6b82c9aa4e99b0f6f

Does this solve your problem ?

[mstorsjo]

openmp/runtime/test/target/target_thread_limit.cpp

29

Yes, that commit fixed the issue.

[kaz7]

I cannot write inline comment, so I'm leaving message here.

openmp/runtime/test/target/target_thread_limit.cpp:

// checking consecutive target regions with different thread_limits
#pragma omp target thread_limit(3)

{
  printf("\nsecond target: thread_limit = %d", omp_get_thread_limit());

// OMP51: second target: thread_limit = 3

Our VE architecture supports only OpenMP runtime. It doesn't support libomptarget. If I run check-openmp on our machine, this omp_get_thread_limit() returns default thread limit 2147483647 (=0x7fffffff). I guess it is OK because this pragma specifies only target and our VE doensn't support target. Other pragmas are containing not only target but also other keyword like parallel, so I guess others are running well. I'm not clear about omptarget, so my assumptions here may be wrong, though.

My question is what is the best way to correct the behavior of this test? I appreciate any comments or suggestions. Thank you!

[tianshilei1992]

In D152054#4642725, @kaz7 wrote:

I run check-openmp on our machine, this omp_get_thread_limit() returns default thread limit 2147483647 (=0x7fffffff).

That is something wrong because this patch is about host instead of offloading.

[kaz7]

In D152054#4642793, @tianshilei1992 wrote:

In D152054#4642725, @kaz7 wrote:

I run check-openmp on our machine, this omp_get_thread_limit() returns default thread limit 2147483647 (=0x7fffffff).

That is something wrong because this patch is about host instead of offloading.

Thank you for your comment. I try to inspect this.

[kaz7]

I figure out the reason of strange behavior I mentioned last night. However, I'm still not sure what is the good way to solve this problem. Can someone check the behavior with -O2 option please? Thanks!

openmp/runtime/test/target/target_thread_limit.cpp

74

Last night, I cannot use inline comment, but now I can. So, I'm writing down my comments here.

The problem I mentioned last night was caused by optimization. This works fine with optimization level by default, but this doesn't work with -O2 even on x86_64. I'm using -O2 for tests, so I come across this problem. I'll write command lines to reproduce my problem below. As you see, there is a warning message at -O2. So, this behavior may be expected, though. In addition, I'm using f8efa65 to produce following results.

$ cd build
$ ninja
...
$ ./bin/clang++  -fopenmp -I runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src -I ../llvm-project/openmp/runtime/test -L runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src -fno-omit-frame-pointer -I ../llvm-project/openmp/runtime/test/ompt -std=c++17 ../llvm-project/openmp/runtime/test/target/target_thread_limit.cpp -o ok -lm -latomic -fopenmp-version=51
$ LD_LIBRARY_PATH=runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src ./ok | grep "second target: thread_limit"
second target: thread_limit = 3
$ ./bin/clang++ -fopenmp -I runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src -I ../llvm-project/openmp/runtime/test -L runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src -fno-omit-frame-pointer -I ../llvm-project/openmp/runtime/test/ompt -std=c++17 ../llvm-project/openmp/runtime/test/target/target_thread_limit.cpp -o bad -lm -latomic -fopenmp-version=51 -O2
warning: loop not vectorized: the optimizer was unable to perform the requested transformation; the transformation might be disabled or specified as part of an unsupported transformation ordering [-Wpass-failed=transform-warning]
1 warning generated.
$ LD_LIBRARY_PATH=runtimes/runtimes-x86_64-unknown-linux-gnu-bins/openmp/runtime/src ./bad | grep "second target: thread_limit"
second target: thread_limit = 2147483647

[sandeepkosuri]

@kaz7, it seems that the thread_limit is being set properly, but the omp_get_thread_limit() is giving a wrong output when you enable anything more than -O1. I will fix it as soon as I can. Meanwhile, if you absolutely want the test case to work right now, remove the printf causing the issue or do not run that test case with a higher optimization level than -O1.

[kaz7]

In D152054#4648318, @sandeepkosuri wrote:

@kaz7, it seems that the thread_limit is being set properly, but the omp_get_thread_limit() is giving a wrong output when you enable anything more than -O1. I will fix it as soon as I can. Meanwhile, if you absolutely want the test case to work right now, remove the printf causing the issue or do not run that test case with a higher optimization level than -O1.

Thank you for investigating that. I'm fine with -O0 result. I just came across this issue and informed that. Thank you for your efforts!

clang/include/clang/Basic/OpenMPKinds.h

void getOpenMPCaptureRegions(
     OpenMPDirectiveKind DKind);
 
 /// Checks if the specified directive is a combined construct for which
 /// the first construct is a parallel construct.
 /// \param DKind Specified directive.
 /// \return true - if the above condition is met for this directive
 /// otherwise - false.
 bool isOpenMPCombinedParallelADirective(OpenMPDirectiveKind DKind);
+
+/// Checks if the specified target directive, combined or not, needs task based
+/// thread_limit
+/// \param DKind Specified directive.
+/// \return true - if the above condition is met for this directive
+/// otherwise - false.
+bool needsTaskBasedThreadLimit(OpenMPDirectiveKind DKind);
 }
 
 #endif
 

clang/lib/Basic/OpenMPKinds.cpp

 bool clang::isOpenMPCombinedParallelADirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_parallel_for || DKind == OMPD_parallel_for_simd ||
          DKind == OMPD_parallel_master ||
          DKind == OMPD_parallel_master_taskloop ||
          DKind == OMPD_parallel_master_taskloop_simd ||
          DKind == OMPD_parallel_sections;
 }
 
+bool clang::needsTaskBasedThreadLimit(OpenMPDirectiveKind DKind) {
+  return DKind == OMPD_target || DKind == OMPD_target_parallel ||
+         DKind == OMPD_target_parallel_for ||
+         DKind == OMPD_target_parallel_for_simd || DKind == OMPD_target_simd ||
+         DKind == OMPD_target_parallel_loop;
+}
+
 void clang::getOpenMPCaptureRegions(
     SmallVectorImpl<OpenMPDirectiveKind> &CaptureRegions,
     OpenMPDirectiveKind DKind) {
   assert(unsigned(DKind) < llvm::omp::Directive_enumSize);
   switch (DKind) {
   case OMPD_metadirective:
     CaptureRegions.push_back(OMPD_metadirective);
     break;

clang/lib/CodeGen/CGOpenMPRuntime.h

public:
   /// Emits call to void __kmpc_push_num_teams(ident_t *loc, kmp_int32
   /// global_tid, kmp_int32 num_teams, kmp_int32 thread_limit) to generate code
   /// for num_teams clause.
   /// \param NumTeams An integer expression of teams.
   /// \param ThreadLimit An integer expression of threads.
   virtual void emitNumTeamsClause(CodeGenFunction &CGF, const Expr *NumTeams,
                                   const Expr *ThreadLimit, SourceLocation Loc);
 
+  /// Emits call to void __kmpc_set_thread_limit(ident_t *loc, kmp_int32
+  /// global_tid, kmp_int32 thread_limit) to generate code for
+  /// thread_limit clause on target directive
+  /// \param ThreadLimit An integer expression of threads.
+  virtual void emitThreadLimitClause(CodeGenFunction &CGF,
+                                     const Expr *ThreadLimit,
+                                     SourceLocation Loc);
+
   /// Struct that keeps all the relevant information that should be kept
   /// throughout a 'target data' region.
   class TargetDataInfo : public llvm::OpenMPIRBuilder::TargetDataInfo {
   public:
     explicit TargetDataInfo() : llvm::OpenMPIRBuilder::TargetDataInfo() {}
     explicit TargetDataInfo(bool RequiresDevicePointerInfo,
                             bool SeparateBeginEndCalls)
         : llvm::OpenMPIRBuilder::TargetDataInfo(RequiresDevicePointerInfo,

clang/lib/CodeGen/CGOpenMPRuntime.cpp

 
 llvm::Value *emitDynCGGroupMem(const OMPExecutableDirective &D,
                                CodeGenFunction &CGF) {
   llvm::Value *DynCGroupMem = CGF.Builder.getInt32(0);
 
   if (auto *DynMemClause = D.getSingleClause<OMPXDynCGroupMemClause>()) {
     CodeGenFunction::RunCleanupsScope DynCGroupMemScope(CGF);
     llvm::Value *DynCGroupMemVal = CGF.EmitScalarExpr(
         DynMemClause->getSize(), /*IgnoreResultAssign=*/true);
     DynCGroupMem = CGF.Builder.CreateIntCast(DynCGroupMemVal, CGF.Int32Ty,
                                              /*isSigned=*/false);
   }
   return DynCGroupMem;
 }
 
 static void emitTargetCallKernelLaunch(
     CGOpenMPRuntime *OMPRuntime, llvm::Function *OutlinedFn,
void CGOpenMPRuntime::emitTargetCall(
   if (!CGF.HaveInsertPoint())
     return;
 
   const bool OffloadingMandatory = !CGM.getLangOpts().OpenMPIsTargetDevice &&
                                    CGM.getLangOpts().OpenMPOffloadMandatory;
 
   assert((OffloadingMandatory || OutlinedFn) && "Invalid outlined function!");
 
-  const bool RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>() ||
-                                 D.hasClausesOfKind<OMPNowaitClause>() ||
-                                 D.hasClausesOfKind<OMPInReductionClause>();
+  const bool RequiresOuterTask =
+      D.hasClausesOfKind<OMPDependClause>() ||
+      D.hasClausesOfKind<OMPInReductionClause>() ||
+      (CGM.getLangOpts().OpenMP >= 51 &&
+       needsTaskBasedThreadLimit(D.getDirectiveKind()) &&
+       D.hasClausesOfKind<OMPThreadLimitClause>());
   llvm::SmallVector<llvm::Value *, 16> CapturedVars;
   const CapturedStmt &CS = *D.getCapturedStmt(OMPD_target);
   auto &&ArgsCodegen = [&CS, &CapturedVars](CodeGenFunction &CGF,
                                             PrePostActionTy &) {
     CGF.GenerateOpenMPCapturedVars(CS, CapturedVars);
   };
   emitInlinedDirective(CGF, OMPD_unknown, ArgsCodegen);
 
void CGOpenMPRuntime::emitNumTeamsClause(CodeGenFunction &CGF,
   // Build call __kmpc_push_num_teamss(&loc, global_tid, num_teams, thread_limit)
   llvm::Value *PushNumTeamsArgs[] = {RTLoc, getThreadID(CGF, Loc), NumTeamsVal,
                                      ThreadLimitVal};
   CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                           CGM.getModule(), OMPRTL___kmpc_push_num_teams),
                       PushNumTeamsArgs);
 }
 
+void CGOpenMPRuntime::emitThreadLimitClause(CodeGenFunction &CGF,
+                                            const Expr *ThreadLimit,
+                                            SourceLocation Loc) {
+  llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
+  llvm::Value *ThreadLimitVal =
+      ThreadLimit
+          ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(ThreadLimit),
+                                      CGF.CGM.Int32Ty, /* isSigned = */ true)
+          : CGF.Builder.getInt32(0);
+
+  // Build call __kmpc_set_thread_limit(&loc, global_tid, thread_limit)
+  llvm::Value *ThreadLimitArgs[] = {RTLoc, getThreadID(CGF, Loc),
+                                    ThreadLimitVal};
+  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
+                          CGM.getModule(), OMPRTL___kmpc_set_thread_limit),
+                      ThreadLimitArgs);
+}
+
 void CGOpenMPRuntime::emitTargetDataCalls(
     CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr *IfCond,
     const Expr *Device, const RegionCodeGenTy &CodeGen,
     CGOpenMPRuntime::TargetDataInfo &Info) {
   if (!CGF.HaveInsertPoint())
     return;
 
   // Action used to replace the default codegen action and turn privatization

clang/lib/CodeGen/CGStmtOpenMP.cpp

if (InputInfo.NumberOfTargetItems > 0) {
       // If MVD is nullptr, the mapper array is not privatized
       if (MVD)
         InputInfo.MappersArray = CGF.Builder.CreateConstArrayGEP(
             CGF.GetAddrOfLocalVar(MVD), /*Index=*/0);
     }
 
     Action.Enter(CGF);
     OMPLexicalScope LexScope(CGF, S, OMPD_task, /*EmitPreInitStmt=*/false);
+    auto *TL = S.getSingleClause<OMPThreadLimitClause>();
+    if (CGF.CGM.getLangOpts().OpenMP >= 51 &&
+        needsTaskBasedThreadLimit(S.getDirectiveKind()) && TL) {
+      // Emit __kmpc_set_thread_limit() to set the thread_limit for the task
+      // enclosing this target region. This will indirectly set the thread_limit
+      // for every applicable construct within target region.
+      CGF.CGM.getOpenMPRuntime().emitThreadLimitClause(
+          CGF, TL->getThreadLimit(), S.getBeginLoc());
+    }
     BodyGen(CGF);
   };
   llvm::Function *OutlinedFn = CGM.getOpenMPRuntime().emitTaskOutlinedFunction(
       S, *I, *PartId, *TaskT, S.getDirectiveKind(), CodeGen, /*Tied=*/true,
       Data.NumberOfParts);
   llvm::APInt TrueOrFalse(32, S.hasClausesOfKind<OMPNowaitClause>() ? 1 : 0);
   IntegerLiteral IfCond(getContext(), TrueOrFalse,
                         getContext().getIntTypeForBitwidth(32, /*Signed=*/0),

clang/test/OpenMP/target_codegen.cpp

 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 0,
 
 #pragma omp target thread_limit(TargetTL)
 {}
 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
 // OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]
 // OMP51: load {{.*}} [[CEA]]
 // OMP51: [[CE:%.*]] = load {{.*}} [[CEA]]
-// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1, i32 [[CE]],
+// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
+// OMP51: call i32 [[OMP_TASK_ENTRY]]
 
 #pragma omp target thread_limit(TargetTL)
 #pragma omp teams
 {}
 // OMP51: [[TL:%.*]] = load {{.*}} %TargetTL.addr
 // OMP51: store {{.*}} [[TL]], {{.*}} [[CEA:%.*]]
 // OMP51: load {{.*}} [[CEA]]
-// OMP51: [[CE:%.*]] = load {{.*}} [[CEA]]
-// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[CE]],
+// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
+// OMP51: call i32 [[OMP_TASK_ENTRY]]
 
 #pragma omp target
 #pragma omp teams thread_limit(TeamsTL)
 {}
 // OMP51: load {{.*}} %TeamsTL.addr
 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
 // OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[TeamsL]],
 
 #pragma omp target thread_limit(TargetTL)
 #pragma omp teams thread_limit(TeamsTL)
 {}
 // OMP51: load {{.*}} %TeamsTL.addr
 // OMP51: [[TeamsL:%.*]] = load {{.*}} %TeamsTL.addr
-// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0, i32 [[TeamsL]],
+// OMP51: call ptr @__kmpc_omp_task_alloc({{.*@.omp_task_entry.*}})
+// OMP51: call i32 [[OMP_TASK_ENTRY]]
 
 }
 #endif
+// Check that the offloading functions are called after setting thread_limit in the task entry functions
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
+// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 -1,
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
+// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 %{{.+}})
+// OMP51: call i32 @__tgt_target_kernel({{.*}}, i64 -1, i32 0,
+
 
 // CHECK:     define internal void @.omp_offloading.requires_reg()
 // CHECK:     call void @__tgt_register_requires(i64 1)
 // CHECK:     ret void
 
 int main () {
   S2 bar;
   bar.zee();
 }
 
 #endif

clang/test/OpenMP/target_parallel_for_simd_tl_codegen.cpp

+// This file is to test thread_limit clause on target parallel for simd directive
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck --check-prefix=OMP51 %s
+
+// expected-no-diagnostics
+
+int thread_limit_target_parallel_for_simd() {
+
+// Check that the offloading function is called after setting thread_limit in the task entry function
+#pragma omp target parallel for simd thread_limit(2)
+    for(int i=0; i<2; i++) {}
+
+// OMP51: define {{.*}}thread_limit_target_parallel_for_simd
+// OMP51: call i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr %1)
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 2)
+// OMP51: call void {{.*omp_offloading.*thread_limit_target_parallel_for_simd.*}}
+  return 0;
+}
 No newline at end of file

clang/test/OpenMP/target_parallel_for_tl_codegen.cpp

+// This file is to test thread_limit clause on target parallel for directive
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck --check-prefix=OMP51 %s
+
+// expected-no-diagnostics
+
+int thread_limit_target_parallel_for() {
+
+// Check that the offloading function is called after setting thread_limit in the task entry function
+#pragma omp target parallel for thread_limit(2)
+    for(int i=0; i<2; i++) {}
+
+// OMP51: define {{.*}}thread_limit_target_parallel_for
+// OMP51: call i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr %1)
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 2)
+// OMP51: call void {{.*omp_offloading.*thread_limit_target_parallel_for.*}}
+  return 0;
+}
 No newline at end of file

clang/test/OpenMP/target_parallel_generic_loop_tl_codegen.cpp

+// This file is to test thread_limit clause on target prallel loop directive
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck --check-prefix=OMP51 %s
+
+// expected-no-diagnostics
+
+int thread_limit_target_parallel_loop() {
+
+// Check that the offloading function is called after setting thread_limit in the task entry function
+#pragma omp target parallel loop thread_limit(2)
+    for(int i=0; i<2; i++) {}
+
+// OMP51: define {{.*}}thread_limit_target_parallel_loop
+// OMP51: call i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr %1)
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 2)
+// OMP51: call void {{.*omp_offloading.*thread_limit_target_parallel_loop.*}}
+  return 0;
+}
 No newline at end of file

clang/test/OpenMP/target_parallel_tl_codegen.cpp

+// This file is to test thread_limit clause on target parallel directive
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck --check-prefix=OMP51 %s
+
+// expected-no-diagnostics
+
+int thread_limit_target_parallel() {
+
+// Check that the offloading function is called after setting thread_limit in the task entry function
+#pragma omp target parallel thread_limit(2)
+{}
+
+// OMP51: define {{.*}}thread_limit_target_parallel
+// OMP51: call i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr %1)
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 2)
+// OMP51: call void {{.*omp_offloading.*thread_limit_target_parallel.*}}
+  return 0;
+}
 No newline at end of file

clang/test/OpenMP/target_simd_tl_codegen.cpp

+// This file is to test thread_limit clause on target simd directive
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck --check-prefix=OMP51 %s
+
+// expected-no-diagnostics
+
+int thread_limit_target_simd() {
+
+// Check that the offloading function is called after setting thread_limit in the task entry function
+#pragma omp target simd thread_limit(2)
+    for(int i=0; i<2; i++) {}
+
+// OMP51: define {{.*}}thread_limit_target_simd
+// OMP51: call i32 [[OMP_TASK_ENTRY:@.+]](i32 {{.*}}%0, ptr %1)
+
+// OMP51: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
+// OMP51: call void @__kmpc_set_thread_limit(ptr @{{.+}}, i32 %{{.+}}, i32 2)
+// OMP51: call void {{.*omp_offloading.*thread_limit_target_simd.*}}
+  return 0;
+}
 No newline at end of file

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

def OMP_TargetParallel : Directive<"target parallel"> {
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_NumThreads>,
     VersionedClause<OMPC_ProcBind>,
     VersionedClause<OMPC_OMPX_DynCGroupMem>,
+    VersionedClause<OMPC_ThreadLimit, 51>,
   ];
 }
 def OMP_TargetParallelFor : Directive<"target parallel for"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
let allowedClauses = [
     VersionedClause<OMPC_IsDevicePtr>,
     VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_OMPX_DynCGroupMem>,
+    VersionedClause<OMPC_ThreadLimit, 51>,
   ];
 }
 def OMP_TargetParallelDo : Directive<"target parallel do"> {
   let allowedClauses = [
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_LastPrivate>,
let allowedClauses = [
     VersionedClause<OMPC_HasDeviceAddr, 51>,
     VersionedClause<OMPC_Allocate>,
     VersionedClause<OMPC_NonTemporal, 50>,
     VersionedClause<OMPC_Order, 50>,
     VersionedClause<OMPC_UsesAllocators, 50>
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_OMPX_DynCGroupMem>,
+    VersionedClause<OMPC_ThreadLimit, 51>,
   ];
 }
 def OMP_TargetParallelDoSimd : Directive<"target parallel do simd"> {
   let allowedClauses = [
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_Map>,
     VersionedClause<OMPC_Private>,
let allowedOnceClauses = [
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_NumThreads>,
     VersionedClause<OMPC_ProcBind>,
     VersionedClause<OMPC_Device>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_Schedule>,
     VersionedClause<OMPC_OMPX_DynCGroupMem>,
-    VersionedClause<OMPC_Order, 50>
+    VersionedClause<OMPC_Order, 50>,
+    VersionedClause<OMPC_ThreadLimit, 51>,
   ];
 }
 def OMP_TeamsDistribute : Directive<"teams distribute"> {
   let allowedClauses = [
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_Shared>,
let allowedOnceClauses = [
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_DefaultMap>,
     VersionedClause<OMPC_If>,
     VersionedClause<OMPC_NoWait>,
     VersionedClause<OMPC_NumThreads>,
     VersionedClause<OMPC_Order>,
     VersionedClause<OMPC_ProcBind>,
     VersionedClause<OMPC_OMPX_DynCGroupMem>,
+    VersionedClause<OMPC_ThreadLimit, 51>,
   ];
 }
 def OMP_Metadirective : Directive<"metadirective"> {
   let allowedClauses = [VersionedClause<OMPC_When>];
   let allowedOnceClauses = [VersionedClause<OMPC_Default>];
 }
 def OMP_Unknown : Directive<"unknown"> {
   let isDefault = true;
 }

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

 __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_set_thread_limit, false, Void, IdentPtr, Int32, Int32)
 
 __OMP_RTL(__kmpc_copyprivate, false, Void, IdentPtr, Int32, SizeTy, VoidPtr,
           CopyFunctionPtr, Int32)
 __OMP_RTL(__kmpc_threadprivate_cached, false, VoidPtr, IdentPtr, Int32, VoidPtr,
           SizeTy, VoidPtrPtrPtr)
 __OMP_RTL(__kmpc_threadprivate_register, false, Void, IdentPtr, VoidPtr,
           KmpcCtorPtr, KmpcCopyCtorPtr, KmpcDtorPtr)
 
__OMP_RTL_ATTRS(__kmpc_omp_taskwait_deps_51, BarrierAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, SExt, ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_cancellationpoint, DefaultAttrs, SExt,
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, SExt))
 
 __OMP_RTL_ATTRS(__kmpc_fork_teams, ForkAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, ReadOnlyPtrAttrs))
 __OMP_RTL_ATTRS(__kmpc_push_num_teams, InaccessibleArgOnlyAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, SExt, SExt))
+__OMP_RTL_ATTRS(__kmpc_set_thread_limit, InaccessibleArgOnlyAttrs, AttributeSet(),
+                ParamAttrs(ReadOnlyPtrAttrs, SExt, SExt))
 
 __OMP_RTL_ATTRS(__kmpc_copyprivate, DefaultAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, SizeTyExt,
                            ReadOnlyPtrAttrs, AttributeSet(), SExt))
 __OMP_RTL_ATTRS(__kmpc_threadprivate_cached, DefaultAttrs, ReturnPtrAttrs,
                 ParamAttrs(ReadOnlyPtrAttrs, SExt, AttributeSet(), SizeTyExt))
 __OMP_RTL_ATTRS(__kmpc_threadprivate_register, DefaultAttrs, AttributeSet(),
                 ParamAttrs(ReadOnlyPtrAttrs, AttributeSet(), ReadOnlyPtrAttrs,

openmp/runtime/src/kmp.h

kmp_int8
       bt_set; /* internal control for whether blocktime is explicitly set */
   int blocktime; /* internal control for blocktime */
 #if KMP_USE_MONITOR
   int bt_intervals; /* internal control for blocktime intervals */
 #endif
   int nproc; /* internal control for #threads for next parallel region (per
                 thread) */
   int thread_limit; /* internal control for thread-limit-var */
+  int task_thread_limit; /* internal control for thread-limit-var of a task*/
   int max_active_levels; /* internal control for max_active_levels */
   kmp_r_sched_t
       sched; /* internal control for runtime schedule {sched,chunk} pair */
   kmp_proc_bind_t proc_bind; /* internal control for affinity  */
   kmp_int32 default_device; /* internal control for default device */
   struct kmp_internal_control *next;
 } kmp_internal_control_t;
 

 extern int __kmp_xproc; /* number of processors in the system */
 extern int __kmp_avail_proc; /* number of processors available to the process */
 extern size_t __kmp_sys_min_stksize; /* system-defined minimum stack size */
 extern int __kmp_sys_max_nth; /* system-imposed maximum number of threads */
 // maximum total number of concurrently-existing threads on device
 extern int __kmp_max_nth;
 // maximum total number of concurrently-existing threads in a contention group
 extern int __kmp_cg_max_nth;
+extern int __kmp_task_max_nth; // max threads used in a task
 extern int __kmp_teams_max_nth; // max threads used in a teams construct
 extern int __kmp_threads_capacity; /* capacity of the arrays __kmp_threads and
                                       __kmp_root */
 extern int __kmp_dflt_team_nth; /* default number of threads in a parallel
                                    region a la OMP_NUM_THREADS */
 extern int __kmp_dflt_team_nth_ub; /* upper bound on "" determined at serial
                                       initialization */
 extern int __kmp_tp_capacity; /* capacity of __kmp_threads if threadprivate is

 KMP_EXPORT void __kmpc_push_num_threads(ident_t *loc, kmp_int32 global_tid,
                                         kmp_int32 num_threads);
 
 KMP_EXPORT void __kmpc_push_proc_bind(ident_t *loc, kmp_int32 global_tid,
                                       int proc_bind);
 KMP_EXPORT void __kmpc_push_num_teams(ident_t *loc, kmp_int32 global_tid,
                                       kmp_int32 num_teams,
                                       kmp_int32 num_threads);
+KMP_EXPORT void __kmpc_set_thread_limit(ident_t *loc, kmp_int32 global_tid,
+                                        kmp_int32 thread_limit);
 /* Function for OpenMP 5.1 num_teams clause */
 KMP_EXPORT void __kmpc_push_num_teams_51(ident_t *loc, kmp_int32 global_tid,
                                          kmp_int32 num_teams_lb,
                                          kmp_int32 num_teams_ub,
                                          kmp_int32 num_threads);
 KMP_EXPORT void __kmpc_fork_teams(ident_t *loc, kmp_int32 argc,
                                   kmpc_micro microtask, ...);
 struct kmp_dim { // loop bounds info casted to kmp_int64

openmp/runtime/src/kmp_csupport.cpp

void __kmpc_push_num_teams(ident_t *loc, kmp_int32 global_tid,
   __kmp_assert_valid_gtid(global_tid);
   __kmp_push_num_teams(loc, global_tid, num_teams, num_threads);
 }
 
 /*!
 @ingroup PARALLEL
 @param loc source location information
 @param global_tid global thread number
+@param thread_limit limit on number of threads which can be created within the
+current task
+
+Set the thread_limit for the current task
+This call is there to support `thread_limit` clause on the `target` construct
+*/
+void __kmpc_set_thread_limit(ident_t *loc, kmp_int32 global_tid,
+                             kmp_int32 thread_limit) {
+  __kmp_assert_valid_gtid(global_tid);
+  kmp_info_t *thread = __kmp_threads[global_tid];
+  if (thread_limit > 0)
+    thread->th.th_current_task->td_icvs.task_thread_limit = thread_limit;
+}
+
+/*!
+@ingroup PARALLEL
+@param loc source location information
+@param global_tid global thread number
 @param num_teams_lb lower bound on number of teams requested for the teams
 construct
 @param num_teams_ub upper bound on number of teams requested for the teams
 construct
 @param num_threads number of threads per team requested for the teams construct
 
 Set the number of teams to be used by the teams construct. The number of initial
 teams cretaed will be greater than or equal to the lower bound and less than or

openmp/runtime/src/kmp_ftn_entry.h

#else
   int gtid;
   kmp_info_t *thread;
   if (!__kmp_init_serial) {
     __kmp_serial_initialize();
   }
 
   gtid = __kmp_entry_gtid();
   thread = __kmp_threads[gtid];
+  // If thread_limit for the target task is defined, return that instead of the
+  // regular task thread_limit
+  if (int thread_limit = thread->th.th_current_task->td_icvs.task_thread_limit)
+    return thread_limit;
   return thread->th.th_current_task->td_icvs.thread_limit;
 #endif
 }
 
 int FTN_STDCALL KMP_EXPAND_NAME(FTN_IN_FINAL)(void) {
 #ifdef KMP_STUB
   return 0; // TO DO: clarify whether it returns 1 or 0?
 #else
   if (!TCR_4(__kmp_init_parallel)) {
     return 0;
   }
   return __kmp_entry_thread()->th.th_current_task->td_flags.final;
 #endif

openmp/runtime/src/kmp_global.cpp

 int __kmp_generate_warnings = kmp_warnings_low;
 int __kmp_reserve_warn = 0;
 int __kmp_xproc = 0;
 int __kmp_avail_proc = 0;
 size_t __kmp_sys_min_stksize = KMP_MIN_STKSIZE;
 int __kmp_sys_max_nth = KMP_MAX_NTH;
 int __kmp_max_nth = 0;
 int __kmp_cg_max_nth = 0;
+int __kmp_task_max_nth = 0;
 int __kmp_teams_max_nth = 0;
 int __kmp_threads_capacity = 0;
 int __kmp_dflt_team_nth = 0;
 int __kmp_dflt_team_nth_ub = 0;
 int __kmp_tp_capacity = 0;
 int __kmp_tp_cached = 0;
 int __kmp_dispatch_num_buffers = KMP_DFLT_DISP_NUM_BUFF;
 int __kmp_dflt_max_active_levels = 1; // Nesting off by default

openmp/runtime/src/kmp_runtime.cpp

int __kmp_fork_call(ident_t *loc, int gtid,
   int master_this_cons;
   kmp_team_t *team;
   kmp_team_t *parent_team;
   kmp_info_t *master_th;
   kmp_root_t *root;
   int nthreads;
   int master_active;
   int master_set_numthreads;
+  int task_thread_limit = 0;
   int level;
   int active_level;
   int teams_level;
 #if KMP_NESTED_HOT_TEAMS
   kmp_hot_team_ptr_t **p_hot_teams;
 #endif
   { // KMP_TIME_BLOCK
     KMP_TIME_DEVELOPER_PARTITIONED_BLOCK(KMP_fork_call);
#endif
     master_th = __kmp_threads[gtid];
 
     parent_team = master_th->th.th_team;
     master_tid = master_th->th.th_info.ds.ds_tid;
     master_this_cons = master_th->th.th_local.this_construct;
     root = master_th->th.th_root;
     master_active = root->r.r_active;
     master_set_numthreads = master_th->th.th_set_nproc;
+    task_thread_limit =
+        master_th->th.th_current_task->td_icvs.task_thread_limit;
 
 #if OMPT_SUPPORT
     ompt_data_t ompt_parallel_data = ompt_data_none;
     ompt_data_t *parent_task_data;
     ompt_frame_t *ompt_frame;
     void *return_address = NULL;
 
     if (ompt_enabled.enabled) {
if ((!enter_teams &&
         (__kmp_library == library_serial)) {
       KC_TRACE(10, ("__kmp_fork_call: T#%d serializing team\n", gtid));
       nthreads = 1;
     } else {
       nthreads = master_set_numthreads
                      ? master_set_numthreads
                      // TODO: get nproc directly from current task
                      : get__nproc_2(parent_team, master_tid);
+      // Use the thread_limit set for the current target task if exists, else go
+      // with the deduced nthreads
+      nthreads = task_thread_limit > 0 && task_thread_limit < nthreads
+                     ? task_thread_limit
+                     : nthreads;
       // Check if we need to take forkjoin lock? (no need for serialized
       // parallel out of teams construct).
       if (nthreads > 1) {
         /* determine how many new threads we can use */
         __kmp_acquire_bootstrap_lock(&__kmp_forkjoin_lock);
         /* AC: If we execute teams from parallel region (on host), then teams
            should be created but each can only have 1 thread if nesting is
            disabled. If teams called from serial region, then teams and their

 #if KMP_USE_MONITOR
     __kmp_bt_intervals, // int bt_intervals; //internal control for blocktime
 // intervals
 #endif
     __kmp_dflt_team_nth, // int nproc; //internal control for # of threads for
     // next parallel region (per thread)
     // (use a max ub on value if __kmp_parallel_initialize not called yet)
     __kmp_cg_max_nth, // int thread_limit;
+    __kmp_task_max_nth, // int task_thread_limit; // to set the thread_limit
+    // on task. This is used in the case of target thread_limit
     __kmp_dflt_max_active_levels, // int max_active_levels; //internal control
     // for max_active_levels
     r_sched, // kmp_r_sched_t sched; //internal control for runtime schedule
     // {sched,chunk} pair
     __kmp_nested_proc_bind.bind_types[0],
     __kmp_default_device,
     NULL // struct kmp_internal_control *next;
   };

openmp/runtime/test/target/target_thread_limit.cpp

+// RUN: %libomp-cxx-compile -fopenmp-version=51
+// RUN: %libomp-run | FileCheck %s --check-prefix OMP51
+
+#include <stdio.h>
+#include <omp.h>
+
+void foo() {
+#pragma omp parallel num_threads(10)
+  { printf("\ntarget: foo(): parallel num_threads(10)"); }
+}
+
+int main(void) {
+
+  int tl = 4;
+  printf("\nmain: thread_limit = %d", omp_get_thread_limit());
+  // OMP51: main: thread_limit = {{[0-9]+}}
+
+#pragma omp target thread_limit(tl)
+  {
+    printf("\ntarget: thread_limit = %d", omp_get_thread_limit());
+// OMP51: target: thread_limit = 4
+// check whether thread_limit is honoured
+#pragma omp parallel
+    { printf("\ntarget: parallel"); }
+// OMP51: target: parallel
+// OMP51: target: parallel
+// OMP51: target: parallel
+// OMP51: target: parallel
+// OMP51-NOT: target: parallel
+
+// check whether num_threads is honoured
+#pragma omp parallel num_threads(2)
+    { printf("\ntarget: parallel num_threads(2)"); }
+// OMP51: target: parallel num_threads(2)
+// OMP51: target: parallel num_threads(2)
+// OMP51-NOT: target: parallel num_threads(2)
+
+// check whether thread_limit is honoured when there is a conflicting
+// num_threads
+#pragma omp parallel num_threads(10)
+    { printf("\ntarget: parallel num_threads(10)"); }
+    // OMP51: target: parallel num_threads(10)
+    // OMP51: target: parallel num_threads(10)
+    // OMP51: target: parallel num_threads(10)
+    // OMP51: target: parallel num_threads(10)
+    // OMP51-NOT: target: parallel num_threads(10)
+
+    // check whether threads are limited across functions
+    foo();
+    // OMP51: target: foo(): parallel num_threads(10)
+    // OMP51: target: foo(): parallel num_threads(10)
+    // OMP51: target: foo(): parallel num_threads(10)
+    // OMP51: target: foo(): parallel num_threads(10)
+    // OMP51-NOT: target: foo(): parallel num_threads(10)
+
+    // check if user can set num_threads at runtime
+    omp_set_num_threads(2);
+#pragma omp parallel
+    { printf("\ntarget: parallel with omp_set_num_thread(2)"); }
+    // OMP51: target: parallel with omp_set_num_thread(2)
+    // OMP51: target: parallel with omp_set_num_thread(2)
+    // OMP51-NOT: target: parallel with omp_set_num_thread(2)
+
+    // make sure thread_limit is unaffected by omp_set_num_threads
+    printf("\ntarget: thread_limit = %d", omp_get_thread_limit());
+    // OMP51: target: thread_limit = 4
+  }
+
+// checking consecutive target regions with different thread_limits
+#pragma omp target thread_limit(3)
+  {
+    printf("\nsecond target: thread_limit = %d", omp_get_thread_limit());
+// OMP51: second target: thread_limit = 3
+#pragma omp parallel
+    { printf("\nsecond target: parallel"); }
+    // OMP51: second target: parallel
+    // OMP51: second target: parallel
+    // OMP51: second target: parallel
+    // OMP51-NOT: second target: parallel
+  }
+
+  // confirm that thread_limit's effects are limited to target region
+  printf("\nmain: thread_limit = %d", omp_get_thread_limit());
+  // OMP51: main: thread_limit = {{[0-9]+}}
+#pragma omp parallel num_threads(10)
+  { printf("\nmain: parallel num_threads(10)"); }
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51: main: parallel num_threads(10)
+  // OMP51-NOT: main: parallel num_threads(10)
+
+// check combined target directives which support thread_limit
+// target parallel
+#pragma omp target parallel thread_limit(2)
+  printf("\ntarget parallel thread_limit(2)");
+  // OMP51: target parallel thread_limit(2)
+  // OMP51: target parallel thread_limit(2)
+  // OMP51-NOT: target parallel thread_limit(2)
+
+#pragma omp target parallel num_threads(2) thread_limit(3)
+  printf("\ntarget parallel num_threads(2) thread_limit(3)");
+  // OMP51: target parallel num_threads(2) thread_limit(3)
+  // OMP51: target parallel num_threads(2) thread_limit(3)
+  // OMP51-NOT: target parallel num_threads(2) thread_limit(3)
+
+#pragma omp target parallel num_threads(3) thread_limit(2)
+  printf("\ntarget parallel num_threads(3) thread_limit(2)");
+  // OMP51: target parallel num_threads(3) thread_limit(2)
+  // OMP51: target parallel num_threads(3) thread_limit(2)
+  // OMP51-NOT: target parallel num_threads(3) thread_limit(2)
+
+// target parallel for
+#pragma omp target parallel for thread_limit(2)
+  for (int i = 0; i < 5; ++i)
+    printf("\ntarget parallel for thread_limit(2) : thread num = %d",
+           omp_get_thread_num());
+    // OMP51: target parallel for thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for thread_limit(2) : thread num = {{0|1}}
+    // OMP51-NOT: target parallel for thread_limit(3) : thread num = {{0|1}}
+
+// target parallel for simd
+#pragma omp target parallel for simd thread_limit(2)
+  for (int i = 0; i < 5; ++i)
+    printf("\ntarget parallel for simd thread_limit(2) : thread num = %d",
+           omp_get_thread_num());
+    // OMP51: target parallel for simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target parallel for simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51-NOT: target parallel for simd thread_limit(2) : thread num =
+    // {{0|1}}
+
+// target simd
+#pragma omp target simd thread_limit(2)
+  for (int i = 0; i < 5; ++i)
+    printf("\ntarget simd thread_limit(2) : thread num = %d",
+           omp_get_thread_num());
+    // OMP51: target simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51: target simd thread_limit(2) : thread num = {{0|1}}
+    // OMP51-NOT: target simd thread_limit(2) : thread num = {{0|1}}
+
+// target parallel loop
+#pragma omp target parallel loop thread_limit(2)
+  for (int i = 0; i < 5; ++i)
+    printf("\ntarget parallel loop thread_limit(2) : thread num = %d",
+           omp_get_thread_num());
+  // # OMP51: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  // # OMP51: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  // # OMP51: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  // # OMP51: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  // # OMP51: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  // # OMP51-NOT: target parallel loop thread_limit(2) : thread num = {{0|1}}
+  return 0;
+}