[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 created this revision.
carlo.bertolli added reviewers: ABataev, sfantao, Hahnfeld.
carlo.bertolli added subscribers: arpith-jacob, caomhin, cfe-commits.
carlo.bertolli set the repository for this revision to rL LLVM.

This patch fixes a bug in code generation of the correct OpenMP runtime library call in presence of target and teams, when target is separated by teams with multiple curly brackets. The current implementation will not be able to see the teams directive inside target and issue a call to tgt_target instead of the correct one tgt_target_teams.

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
@@ -4249,6 +4249,29 @@
       DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID);
 }
 
+/// \brief Search in a target region until either a teams directive is found
+/// or a different statement or construct. Return the teams directive object
+/// in the former case, nullptr otherwise.
+const static OMPTeamsDirective *
+hasEnclosingTeams(const Stmt *TargetBody) {
+  if(auto *TeamsDir = dyn_cast<OMPTeamsDirective>(TargetBody)) return TeamsDir;
+
+  auto *NextBlock = dyn_cast<CompoundStmt>(TargetBody);
+  auto *LastBlock = NextBlock;
+  // keep reading compound statements until something else is found
+  // or there is nothing more to read
+  while (NextBlock && NextBlock->body_front()) {
+    LastBlock = NextBlock;
+    NextBlock = dyn_cast<CompoundStmt>(NextBlock->body_front());
+  }
+
+  if (LastBlock && LastBlock->body_front())
+    return dyn_cast<OMPTeamsDirective>(LastBlock->body_front());
+
+  // no teams or no compund stmt in the target region body
+  return 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
@@ -4279,7 +4302,7 @@
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+   if (auto *TeamsDir = hasEnclosingTeams(CS.getCapturedStmt())) {
     if (auto *NTE = TeamsDir->getSingleClause<OMPNumTeamsClause>()) {
       CGOpenMPInnerExprInfo CGInfo(CGF, CS);
       CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
@@ -4327,7 +4350,7 @@
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+  if (auto *TeamsDir = hasEnclosingTeams(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.51640.patch
Type: text/x-patch
Size: 3151 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160325/3f2e9bb8/attachment.bin>