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

[Carlo Bertolli via cfe-commits]

Author: cbertol
Date: Thu Apr 28 20:37:30 2016
New Revision: 267972

URL: http://llvm.org/viewvc/llvm-project?rev=267972&view=rev
Log:
[OPENMP] Enable correct generation of runtime call when target directive is separated from teams directive by multiple curly brackets

http://reviews.llvm.org/D18474

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.


Modified:
    cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
    cfe/trunk/test/OpenMP/teams_codegen.cpp

Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=267972&r1=267971&r2=267972&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Apr 28 20:37:30 2016
@@ -4556,6 +4556,14 @@ void CGOpenMPRuntime::emitTargetOutlined
       DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID);
 }
 
+/// discard all CompoundStmts intervening between two constructs
+static const Stmt *ignoreCompoundStmts(const Stmt *Body) {
+  while (auto *CS = dyn_cast_or_null<CompoundStmt>(Body))
+    Body = CS->body_front();
+
+  return Body;
+}
+
 /// \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
@@ -4586,7 +4594,8 @@ emitNumTeamsClauseForTargetDirective(CGO
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+  if (auto *TeamsDir = dyn_cast_or_null<OMPTeamsDirective>(
+          ignoreCompoundStmts(CS.getCapturedStmt()))) {
     if (auto *NTE = TeamsDir->getSingleClause<OMPNumTeamsClause>()) {
       CGOpenMPInnerExprInfo CGInfo(CGF, CS);
       CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);
@@ -4634,7 +4643,8 @@ emitThreadLimitClauseForTargetDirective(
 
   // FIXME: Accommodate other combined directives with teams when they become
   // available.
-  if (auto *TeamsDir = dyn_cast<OMPTeamsDirective>(CS.getCapturedStmt())) {
+  if (auto *TeamsDir = dyn_cast_or_null<OMPTeamsDirective>(
+          ignoreCompoundStmts(CS.getCapturedStmt()))) {
     if (auto *TLE = TeamsDir->getSingleClause<OMPThreadLimitClause>()) {
       CGOpenMPInnerExprInfo CGInfo(CGF, CS);
       CodeGenFunction::CGCapturedStmtRAII CapInfoRAII(CGF, &CGInfo);

Modified: cfe/trunk/test/OpenMP/teams_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_codegen.cpp?rev=267972&r1=267971&r2=267972&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/teams_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/teams_codegen.cpp Thu Apr 28 20:37:30 2016
@@ -29,6 +29,16 @@ int teams_argument_global_local(int a){
     ++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:%[^,]+]],