[PATCH] D18474: [OPENMP] Enable correct generation of runtime call when target directive is separated from teams directive by multiple curly brackets

[Carlo Bertolli via cfe-commits]

carlo.bertolli updated this revision to Diff 53283.
carlo.bertolli added a comment.

[OPENMP] Remove template for intervening statement and only allow CompoundStmt.


Repository:
  rL LLVM

http://reviews.llvm.org/D18474

Files:
  lib/CodeGen/CGOpenMPRuntime.cpp
  test/OpenMP/teams_codegen.cpp

Index: test/OpenMP/teams_codegen.cpp
===================================================================
--- test/OpenMP/teams_codegen.cpp
+++ test/OpenMP/teams_codegen.cpp
@@ -29,6 +29,16 @@
     ++comp;
   }
 
+  // CK1: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
+  // CK1: call void @{{.+}}(i{{64|32}} %{{.+}})
+  #pragma omp target
+  {{{
+    #pragma omp teams
+    {
+      ++comp;
+    }
+  }}}
+  
   // CK1-DAG: call i32 @__tgt_target_teams(i32 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 [[NT:%[^,]+]], i32 0)
   // CK1-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]],
 
Index: lib/CodeGen/CGOpenMPRuntime.cpp
===================================================================
--- lib/CodeGen/CGOpenMPRuntime.cpp
+++ lib/CodeGen/CGOpenMPRuntime.cpp
@@ -4257,6 +4257,20 @@
       DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID);
 }
 
+/// \brief look inside a Body stmt for a LF statement discarding any
+/// intervening CompoundStmt's
+template <typename LF>
+const static LF *hasEnclosingOpenMPDirective(const Stmt *Body) {
+  const CompoundStmt *S = nullptr;
+  // keep iterating until we find an LF or not a CompoundStmt or a nullptr
+  while ((S = dyn_cast_or_null<CompoundStmt>(Body)) &&
+         (!dyn_cast_or_null<LF>(Body)))
+      Body = S->body_front();
+
+  return (Body) ? dyn_cast_or_null<LF>(Body) :
+      nullptr;
+}
+
 /// \brief Emit the num_teams clause of an enclosed teams directive at the
 /// target region scope. If there is no teams directive associated with the
 /// target directive, or if there is no num_teams clause associated with the
@@ -4287,7 +4301,8 @@
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+  if (auto *TeamsDir = hasEnclosingOpenMPDirective<OMPTeamsDirective>(
+          CS.getCapturedStmt())) {
     if (auto *NTE = TeamsDir->getSingleClause<OMPNumTeamsClause>()) {
       CGOpenMPInnerExprInfo CGInfo(CGF, CS);
       CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
@@ -4335,7 +4350,8 @@
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+  if (auto *TeamsDir = hasEnclosingOpenMPDirective<OMPTeamsDirective>(
+          CS.getCapturedStmt())) {
     if (auto *TLE = TeamsDir->getSingleClause<OMPThreadLimitClause>()) {
       CGOpenMPInnerExprInfo CGInfo(CGF, CS);
       CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D18474.53283.patch
Type: text/x-patch
Size: 2828 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160411/43905e5a/attachment-0001.bin>