diff options
-rw-r--r-- | clang/lib/CodeGen/CGOpenMPRuntime.cpp | 14 | ||||
-rw-r--r-- | clang/test/OpenMP/teams_codegen.cpp | 10 |
2 files changed, 22 insertions, 2 deletions
diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 70057551d91..b280d53dd39 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -4556,6 +4556,14 @@ void CGOpenMPRuntime::emitTargetOutlinedFunctionHelper( 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(CGOpenMPRuntime &OMPRuntime, // 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(CGOpenMPRuntime &OMPRuntime, // 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); diff --git a/clang/test/OpenMP/teams_codegen.cpp b/clang/test/OpenMP/teams_codegen.cpp index 1d616c37815..c26b5862575 100644 --- a/clang/test/OpenMP/teams_codegen.cpp +++ b/clang/test/OpenMP/teams_codegen.cpp @@ -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:%[^,]+]], |