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);
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);
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
