[OpenMP] Support for thread_limit-clause on the 'target teams' directive.

[arpith-jacob]

Summary

The thread_limit-clause on the combined directive applies to the
'teams' region of this construct. We modify the ThreadLimitClause
class to capture the clause expression within the 'target' region.

[ABataev]

LG

cfe/trunk/include/clang/AST/OpenMPClause.h

 /// directive.
 ///
 /// \code
 /// #pragma omp teams thread_limit(n)
 /// \endcode
 /// In this example directive '#pragma omp teams' has clause 'thread_limit'
 /// with single expression 'n'.
 ///
-class OMPThreadLimitClause : public OMPClause {
+class OMPThreadLimitClause : public OMPClause, public OMPClauseWithPreInit {
   friend class OMPClauseReader;
   /// \brief Location of '('.
   SourceLocation LParenLoc;
   /// \brief ThreadLimit number.
   Stmt *ThreadLimit;
   /// \brief Set the ThreadLimit number.
   ///
   /// \param E ThreadLimit number.
   ///
   void setThreadLimit(Expr *E) { ThreadLimit = E; }
 
 public:
   /// \brief Build 'thread_limit' clause.
   ///
   /// \param E Expression associated with this clause.
+  /// \param HelperE Helper Expression associated with this clause.
+  /// \param CaptureRegion Innermost OpenMP region where expressions in this
+  /// clause must be captured.
   /// \param StartLoc Starting location of the clause.
   /// \param LParenLoc Location of '('.
   /// \param EndLoc Ending location of the clause.
   ///
-  OMPThreadLimitClause(Expr *E, SourceLocation StartLoc,
-                       SourceLocation LParenLoc, SourceLocation EndLoc)
-      : OMPClause(OMPC_thread_limit, StartLoc, EndLoc), LParenLoc(LParenLoc),
-        ThreadLimit(E) {}
+  OMPThreadLimitClause(Expr *E, Stmt *HelperE,
+                       OpenMPDirectiveKind CaptureRegion,
+                       SourceLocation StartLoc, SourceLocation LParenLoc,
+                       SourceLocation EndLoc)
+      : OMPClause(OMPC_thread_limit, StartLoc, EndLoc),
+        OMPClauseWithPreInit(this), LParenLoc(LParenLoc), ThreadLimit(E) {
+    setPreInitStmt(HelperE, CaptureRegion);
+  }
 
   /// \brief Build an empty clause.
   ///
   OMPThreadLimitClause()
       : OMPClause(OMPC_thread_limit, SourceLocation(), SourceLocation()),
-        LParenLoc(SourceLocation()), ThreadLimit(nullptr) {}
+        OMPClauseWithPreInit(this), LParenLoc(SourceLocation()),
+        ThreadLimit(nullptr) {}
   /// \brief Sets the location of '('.
   void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; }
   /// \brief Returns the location of '('.
   SourceLocation getLParenLoc() const { return LParenLoc; }
   /// \brief Return ThreadLimit number.
   Expr *getThreadLimit() { return cast<Expr>(ThreadLimit); }
   /// \brief Return ThreadLimit number.
   Expr *getThreadLimit() const { return cast<Expr>(ThreadLimit); }

cfe/trunk/include/clang/AST/RecursiveASTVisitor.h

bool RecursiveASTVisitor<Derived>::VisitOMPNumTeamsClause(
   TRY_TO(VisitOMPClauseWithPreInit(C));
   TRY_TO(TraverseStmt(C->getNumTeams()));
   return true;
 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPThreadLimitClause(
     OMPThreadLimitClause *C) {
+  TRY_TO(VisitOMPClauseWithPreInit(C));
   TRY_TO(TraverseStmt(C->getThreadLimit()));
   return true;
 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPPriorityClause(
     OMPPriorityClause *C) {
   TRY_TO(TraverseStmt(C->getPriority()));

cfe/trunk/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_linear:
     return static_cast<const OMPLinearClause *>(C);
   case OMPC_if:
     return static_cast<const OMPIfClause *>(C);
   case OMPC_num_threads:
     return static_cast<const OMPNumThreadsClause *>(C);
   case OMPC_num_teams:
     return static_cast<const OMPNumTeamsClause *>(C);
+  case OMPC_thread_limit:
+    return static_cast<const OMPThreadLimitClause *>(C);
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_final:
   case OMPC_safelen:
   case OMPC_simdlen:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_shared:
const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_map:
-  case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:

cfe/trunk/lib/AST/StmtProfile.cpp

 }
 void OMPClauseProfiler::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
   VistOMPClauseWithPreInit(C);
   if (C->getNumTeams())
     Profiler->VisitStmt(C->getNumTeams());
 }
 void OMPClauseProfiler::VisitOMPThreadLimitClause(
     const OMPThreadLimitClause *C) {
+  VistOMPClauseWithPreInit(C);
   if (C->getThreadLimit())
     Profiler->VisitStmt(C->getThreadLimit());
 }
 void OMPClauseProfiler::VisitOMPPriorityClause(const OMPPriorityClause *C) {
   if (C->getPriority())
     Profiler->VisitStmt(C->getPriority());
 }
 void OMPClauseProfiler::VisitOMPGrainsizeClause(const OMPGrainsizeClause *C) {

cfe/trunk/lib/Sema/SemaOpenMP.cpp

case OMPC_num_teams:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
       llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
+  case OMPC_thread_limit:
+    switch (DKind) {
+    case OMPD_target_teams:
+      CaptureRegion = OMPD_target;
+      break;
+    case OMPD_cancel:
+    case OMPD_parallel:
+    case OMPD_parallel_sections:
+    case OMPD_parallel_for:
+    case OMPD_parallel_for_simd:
+    case OMPD_target:
+    case OMPD_target_simd:
+    case OMPD_target_parallel:
+    case OMPD_target_parallel_for:
+    case OMPD_target_parallel_for_simd:
+    case OMPD_target_teams_distribute:
+    case OMPD_target_teams_distribute_simd:
+    case OMPD_target_teams_distribute_parallel_for:
+    case OMPD_target_teams_distribute_parallel_for_simd:
+    case OMPD_teams_distribute_parallel_for:
+    case OMPD_teams_distribute_parallel_for_simd:
+    case OMPD_distribute_parallel_for:
+    case OMPD_distribute_parallel_for_simd:
+    case OMPD_task:
+    case OMPD_taskloop:
+    case OMPD_taskloop_simd:
+    case OMPD_target_data:
+    case OMPD_target_enter_data:
+    case OMPD_target_exit_data:
+    case OMPD_target_update:
+    case OMPD_teams:
+    case OMPD_teams_distribute:
+    case OMPD_teams_distribute_simd:
+      // Do not capture thread_limit-clause expressions.
+      break;
+    case OMPD_threadprivate:
+    case OMPD_taskyield:
+    case OMPD_barrier:
+    case OMPD_taskwait:
+    case OMPD_cancellation_point:
+    case OMPD_flush:
+    case OMPD_declare_reduction:
+    case OMPD_declare_simd:
+    case OMPD_declare_target:
+    case OMPD_end_declare_target:
+    case OMPD_simd:
+    case OMPD_for:
+    case OMPD_for_simd:
+    case OMPD_sections:
+    case OMPD_section:
+    case OMPD_single:
+    case OMPD_master:
+    case OMPD_critical:
+    case OMPD_taskgroup:
+    case OMPD_distribute:
+    case OMPD_ordered:
+    case OMPD_atomic:
+    case OMPD_distribute_simd:
+      llvm_unreachable("Unexpected OpenMP directive with thread_limit-clause");
+    case OMPD_unknown:
+      llvm_unreachable("Unknown OpenMP directive");
+    }
+    break;
   case OMPC_schedule:
   case OMPC_dist_schedule:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_reduction:
   case OMPC_linear:
   case OMPC_default:
   case OMPC_proc_bind:
static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_map:
-  case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_defaultmap:
   case OMPC_unknown:
   case OMPC_uniform:
return new (Context) OMPNumTeamsClause(ValExpr, HelperValStmt, CaptureRegion,
                                          StartLoc, LParenLoc, EndLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPThreadLimitClause(Expr *ThreadLimit,
                                               SourceLocation StartLoc,
                                               SourceLocation LParenLoc,
                                               SourceLocation EndLoc) {
   Expr *ValExpr = ThreadLimit;
+  Stmt *HelperValStmt = nullptr;
+  OpenMPDirectiveKind CaptureRegion = OMPD_unknown;
 
   // OpenMP [teams Constrcut, Restrictions]
   // The thread_limit expression must evaluate to a positive integer value.
   if (!IsNonNegativeIntegerValue(ValExpr, *this, OMPC_thread_limit,
                                  /*StrictlyPositive=*/true))
     return nullptr;
 
-  return new (Context)
-      OMPThreadLimitClause(ValExpr, StartLoc, LParenLoc, EndLoc);
+  OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
+  CaptureRegion = getOpenMPCaptureRegionForClause(DKind, OMPC_thread_limit);
+  if (CaptureRegion != OMPD_unknown) {
+    llvm::MapVector<Expr *, DeclRefExpr *> Captures;
+    ValExpr = tryBuildCapture(*this, ValExpr, Captures).get();
+    HelperValStmt = buildPreInits(Context, Captures);
+  }
+
+  return new (Context) OMPThreadLimitClause(
+      ValExpr, HelperValStmt, CaptureRegion, StartLoc, LParenLoc, EndLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPPriorityClause(Expr *Priority,
                                            SourceLocation StartLoc,
                                            SourceLocation LParenLoc,
                                            SourceLocation EndLoc) {
   Expr *ValExpr = Priority;
 

cfe/trunk/lib/Serialization/ASTReaderStmt.cpp

 
 void OMPClauseReader::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) {
   VisitOMPClauseWithPreInit(C);
   C->setNumTeams(Reader->Record.readSubExpr());
   C->setLParenLoc(Reader->ReadSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPThreadLimitClause(OMPThreadLimitClause *C) {
+  VisitOMPClauseWithPreInit(C);
   C->setThreadLimit(Reader->Record.readSubExpr());
   C->setLParenLoc(Reader->ReadSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPPriorityClause(OMPPriorityClause *C) {
   C->setPriority(Reader->Record.readSubExpr());
   C->setLParenLoc(Reader->ReadSourceLocation());
 }

cfe/trunk/lib/Serialization/ASTWriterStmt.cpp

 
 void OMPClauseWriter::VisitOMPNumTeamsClause(OMPNumTeamsClause *C) {
   VisitOMPClauseWithPreInit(C);
   Record.AddStmt(C->getNumTeams());
   Record.AddSourceLocation(C->getLParenLoc());
 }
 
 void OMPClauseWriter::VisitOMPThreadLimitClause(OMPThreadLimitClause *C) {
+  VisitOMPClauseWithPreInit(C);
   Record.AddStmt(C->getThreadLimit());
   Record.AddSourceLocation(C->getLParenLoc());
 }
 
 void OMPClauseWriter::VisitOMPPriorityClause(OMPPriorityClause *C) {
   Record.AddStmt(C->getPriority());
   Record.AddSourceLocation(C->getLParenLoc());
 }

cfe/trunk/test/OpenMP/target_teams_thread_limit_codegen.cpp

+// Test host codegen.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// CHECK-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
+// CHECK-DAG: [[STR:@.+]] = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CHECK-DAG: [[DEF_LOC:@.+]] = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* [[STR]], i32 0, i32 0) }
+
+// CHECK-DAG: [[S1:%.+]] = type { double }
+// CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 }
+// CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* }
+// CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* }
+
+// TCHECK: [[ENTTY:%.+]] = type { i8*, i8*, i{{32|64}}, i32, i32 }
+
+// We have 6 target regions
+
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+// CHECK-DAG: @{{.*}} = private constant i8 0
+
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+// TCHECK: @{{.+}} = constant [[ENTTY]]
+
+// Check if offloading descriptor is created.
+// CHECK: [[ENTBEGIN:@.+]] = external constant [[ENTTY]]
+// CHECK: [[ENTEND:@.+]] = external constant [[ENTTY]]
+// CHECK: [[DEVBEGIN:@.+]] = external constant i8
+// CHECK: [[DEVEND:@.+]] = external constant i8
+// CHECK: [[IMAGES:@.+]] = internal unnamed_addr constant [1 x [[DEVTY]]] [{{.+}} { i8* [[DEVBEGIN]], i8* [[DEVEND]], [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }]
+// CHECK: [[DESC:@.+]] = internal constant [[DSCTY]] { i32 1, [[DEVTY]]* getelementptr inbounds ([1 x [[DEVTY]]], [1 x [[DEVTY]]]* [[IMAGES]], i32 0, i32 0), [[ENTTY]]* [[ENTBEGIN]], [[ENTTY]]* [[ENTEND]] }
+
+// Check target registration is registered as a Ctor.
+// CHECK: appending global [1 x { i32, void ()*, i8* }] [{ i32, void ()*, i8* } { i32 0, void ()* bitcast (void (i8*)* [[REGFN:@.+]] to void ()*), i8* null }]
+
+
+template<typename tx>
+tx ftemplate(int n) {
+  tx a = 0;
+
+  #pragma omp target teams thread_limit(tx(20))
+  {
+  }
+
+  short b = 1;
+  #pragma omp target teams num_teams(b) thread_limit(1024)
+  {
+    a += b;
+  }
+
+  return a;
+}
+
+static
+int fstatic(int n) {
+
+  #pragma omp target teams num_teams(n) thread_limit(n*32)
+  {
+  }
+
+  #pragma omp target teams thread_limit(32+n)
+  {
+  }
+
+  return n+1;
+}
+
+struct S1 {
+  double a;
+
+  int r1(int n){
+    int b = 1;
+
+    #pragma omp target teams thread_limit(n-b)
+    {
+      this->a = (double)b + 1.5;
+    }
+
+    #pragma omp target teams thread_limit(1024)
+    {
+      this->a = 2.5;
+    }
+
+    return (int)a;
+  }
+};
+
+// CHECK: define {{.*}}@{{.*}}bar{{.*}}
+int bar(int n){
+  int a = 0;
+
+  S1 S;
+  // CHECK: call {{.*}}i32 [[FS1:@.+]]([[S1]]* {{.*}}, i32 {{.*}})
+  a += S.r1(n);
+
+  // CHECK: call {{.*}}i32 [[FSTATIC:@.+]](i32 {{.*}})
+  a += fstatic(n);
+
+  // CHECK: call {{.*}}i32 [[FTEMPLATE:@.+]](i32 {{.*}})
+  a += ftemplate<int>(n);
+
+  return a;
+}
+
+
+
+//
+// CHECK: define {{.*}}[[FS1]]([[S1]]* {{%.+}}, i32 {{[^%]*}}[[PARM:%.+]])
+//
+// CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
+// CHECK:       store i32 1, i32* [[B:%.+]], align
+// CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
+// CHECK:       [[BV:%.+]] = load i32, i32* [[B]], align
+// CHECK:       [[SUB:%.+]] = sub nsw i32 [[NV]], [[BV]]
+// CHECK:       store i32 [[SUB]], i32* [[CAPE_ADDR:%.+]], align
+// CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
+// CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
+// CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
+// CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
+// CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 0, i32 [[TL]])
+// CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT1:@.+]]([[S1]]* {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
+// CHECK:       br label {{%?}}[[END]]
+// CHECK:       [[END]]
+//
+//
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i32 0, i32 1024)
+// CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT2:@.+]]([[S1]]* {{[^,]+}})
+// CHECK:       br label {{%?}}[[END]]
+// CHECK:       [[END]]
+//
+
+
+
+
+
+
+//
+// CHECK: define {{.*}}[[FSTATIC]](i32 {{[^%]*}}[[PARM:%.+]])
+//
+// CHECK-DAG:   store i32 [[PARM]], i32* [[N_ADDR:%.+]], align
+// CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
+// CHECK:       store i32 [[NV]], i32* [[CAPE_ADDR1:%.+]], align
+// CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
+// CHECK:       [[MUL:%.+]] = mul nsw i32 [[NV]], 32
+// CHECK:       store i32 [[MUL]], i32* [[CAPE_ADDR2:%.+]], align
+// CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR1:%.+]] to i32*
+// CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
+// CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR1:%.+]], align
+// CHECK:       [[ARG1:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR1]], align
+// CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR2:%.+]] to i32*
+// CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
+// CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR2:%.+]], align
+// CHECK:       [[ARG2:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR2]], align
+// CHECK:       [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
+// CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i32 [[TEAMS]], i32 [[TL]])
+// CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT3:@.+]](i[[SZ]] [[ARG1]], i[[SZ]] [[ARG2]])
+// CHECK:       br label {{%?}}[[END]]
+// CHECK:       [[END]]
+//
+//
+//
+// CHECK:       [[NV:%.+]] = load i32, i32* [[N_ADDR]], align
+// CHECK:       [[ADD:%.+]] = add nsw i32 32, [[NV]]
+// CHECK:       store i32 [[ADD]], i32* [[CAPE_ADDR:%.+]], align
+// CHECK:       [[CEV:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i32*
+// CHECK-64:    store i32 [[CEV]], i32* [[CONV]], align
+// CHECK-32:    store i32 [[CEV]], i32* [[CAPEC_ADDR:%.+]], align
+// CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
+// CHECK:       [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, {{.*}}, i32 0, i32 [[TL]])
+// CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT4:@.+]](i[[SZ]] [[ARG]])
+// CHECK:       br label {{%?}}[[END]]
+// CHECK:       [[END]]
+//
+
+
+
+
+
+
+//
+// CHECK: define {{.*}}[[FTEMPLATE]]
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 0, {{.*}}, i32 0, i32 20)
+// CHECK-NEXT:  store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK-NEXT:  [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK-NEXT:  [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK-NEXT:  br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT5:@.+]]()
+// CHECK:       br label {{%?}}[[END]]
+//
+// CHECK:       [[END]]
+//
+//
+//
+// CHECK:       store i16 1, i16* [[B:%.+]], align
+// CHECK:       [[BV:%.+]] = load i16, i16* [[B]], align
+// CHECK:       store i16 [[BV]], i16* [[CAPE_ADDR:%.+]], align
+// CHECK:       [[CEV:%.+]] = load i16, i16* [[CAPE_ADDR]], align
+// CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPEC_ADDR:%.+]] to i16*
+// CHECK:       store i16 [[CEV]], i16* [[CONV]], align
+// CHECK:       [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align
+// CHECK:       [[T:%.+]] = load i16, i16* [[CAPE_ADDR]], align
+// CHECK:       [[TEAMS:%.+]] = sext i16 [[T]] to i32
+//
+// CHECK-DAG:   [[RET:%.+]] = call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i32 [[TEAMS]], i32 1024)
+// CHECK:       store i32 [[RET]], i32* [[RHV:%.+]], align
+// CHECK:       [[RET2:%.+]] = load i32, i32* [[RHV]], align
+// CHECK:       [[ERROR:%.+]] = icmp ne i32 [[RET2]], 0
+// CHECK:       br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]]
+//
+// CHECK:       [[FAIL]]
+// CHECK:       call void [[HVT6:@.+]](i[[SZ]] {{%.+}}, i[[SZ]] {{%.+}}, i[[SZ]] [[ARG]])
+// CHECK:       br label {{%?}}[[END]]
+// CHECK:       [[END]]
+//
+
+
+
+
+
+
+// Check that the offloading functions are emitted and that the parallel function
+// is appropriately guarded.
+
+// CHECK:       define internal void [[HVT1]]([[S1]]* {{%.+}}, i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
+// CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
+// CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
+// CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2,
+//
+//
+
+
+// CHECK:       define internal void [[HVT2]]([[S1]]* {{%.+}})
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 1024)
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 1,
+//
+//
+
+
+
+
+
+
+
+
+// CHECK:       define internal void [[HVT3]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]])
+// CHECK-DAG:   store i[[SZ]] [[PARM1]], i[[SZ]]* [[CAPE_ADDR1:%.+]], align
+// CHECK-DAG:   store i[[SZ]] [[PARM2]], i[[SZ]]* [[CAPE_ADDR2:%.+]], align
+// CHECK-64:    [[CONV1:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR1]] to i32*
+// CHECK-64:    [[CONV2:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR2]] to i32*
+// CHECK-64:    [[NT:%.+]] = load i32, i32* [[CONV1]], align
+// CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV2]], align
+// CHECK-32:    [[NT:%.+]] = load i32, i32* [[CAPE_ADDR1]], align
+// CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR2]], align
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 [[TL]])
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
+//
+//
+// CHECK:       define internal void [[HVT4]](i[[SZ]] [[PARM:%.+]])
+// CHECK-DAG:   store i[[SZ]] [[PARM]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
+// CHECK-64:    [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i32*
+// CHECK-64:    [[TL:%.+]] = load i32, i32* [[CONV]], align
+// CHECK-32:    [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 [[TL]])
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
+//
+//
+
+
+
+
+
+// CHECK:       define internal void [[HVT5]](
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 0, i32 20)
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 0,
+//
+//
+
+
+// CHECK:       define internal void [[HVT6]](i[[SZ]] [[PARM1:%.+]], i[[SZ]] [[PARM2:%.+]], i[[SZ]] [[PARM3:%.+]])
+// CHECK-DAG:   store i[[SZ]] [[PARM3]], i[[SZ]]* [[CAPE_ADDR:%.+]], align
+// CHECK:       [[CONV:%.+]] = bitcast i[[SZ]]* [[CAPE_ADDR]] to i16*
+// CHECK:       [[T:%.+]] = load i16, i16* [[CONV]], align
+// CHECK:       [[NT:%.+]] = sext i16 [[T]] to i32
+// CHECK:       call i32 @__kmpc_push_num_teams(%ident_t* {{[^,]+}}, i32 {{[^,]+}}, i32 [[NT]], i32 1024)
+// CHECK:       call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC]], i32 2,
+//
+//
+
+
+
+#endif

cfe/trunk/tools/libclang/CIndex.cpp

 }
 
 void OMPClauseEnqueue::VisitOMPNumTeamsClause(const OMPNumTeamsClause *C) {
   VisitOMPClauseWithPreInit(C);
   Visitor->AddStmt(C->getNumTeams());
 }
 
 void OMPClauseEnqueue::VisitOMPThreadLimitClause(const OMPThreadLimitClause *C) {
+  VisitOMPClauseWithPreInit(C);
   Visitor->AddStmt(C->getThreadLimit());
 }
 
 void OMPClauseEnqueue::VisitOMPPriorityClause(const OMPPriorityClause *C) {
   Visitor->AddStmt(C->getPriority());
 }
 
 void OMPClauseEnqueue::VisitOMPGrainsizeClause(const OMPGrainsizeClause *C) {