Commit 56252580 authored by Joel E. Denny's avatar Joel E. Denny
Browse files

WIP: [Clacc][OpenACC] Prototype `acc routine seq`

parent a59e21d4
......@@ -240,7 +240,7 @@ Run-Time Environment Variables
* Appearing within a `parallel` construct and any number of levels
of nesting within other `loop` directives are supported.
* Appearing outside a `parallel` construct (that is, an orphaned
loop) is not yet supported.
`loop` construct) is not yet supported.
* Use without clauses is supported.
* Supported partitionability clauses
* Implicit `independent`
......@@ -303,6 +303,25 @@ Run-Time Environment Variables
* A `reduction` clause implies a `copy` clause (overriding the
implicit `firstprivate` clause).
`routine` Directive
-------------------
* Lexical context
* Appearing at file scope is supported.
* Appearing within a function definition is not supported.
* Supported clauses
* `seq` (required)
* Associated declaration
* A lone function definition or prototype is supported.
* A declaration containing multiple declarators is not supported.
For example, `void foo(), bar();`.
* Function definition body
* Appearance of any OpenACC directive produces a compile-time
error diagnostic. Thus, orphaned `loop` constructs are not yet
supported.
* Declaration of a static local variable produces a compile-time
error diagnostic.
Subarrays
---------
......@@ -317,10 +336,11 @@ Subarrays
Device-Side Directives
----------------------
Nesting of an `update`, `data`, `parallel`, or `parallel loop`
directive inside a `parallel`, `loop`, or `parallel loop` construct is
not yet supported. We're not aware of any OpenACC implementation that
supports this yet.
Nesting of an `update`, `enter data`, `exit data`, `data`, `parallel`,
or `parallel loop` directive inside a `parallel`, `loop`, or `parallel
loop` construct or inside a function attributed with a `routine`
directive is not yet supported. We're not aware of any OpenACC
implementation that supports such cases yet.
OpenACC Runtime Library API and Preprocessor
--------------------------------------------
......
......@@ -28,9 +28,18 @@ namespace clang {
// AST classes for directives.
//===----------------------------------------------------------------------===//
/// This is a basic class for representing single OpenACC executable
/// directive.
/// This is a basic class for representing OpenACC executable directives and
/// constructs.
///
/// FIXME: The name should be changed to reflect that it's not just OpenACC
/// executable directives (like acc update). The base class ExecutableDirective
/// should become something more generic too then, but it has to also make sense
/// for OMPExecutableDirective. ActOnOpenACCExecutableDirective should be
/// renamed in Sema accordingly. Perhaps ACCExecutableDirectiveOrConstruct is
/// the right name to distinguish it from declarative directives. Perhaps the
/// base class should be ExecutableDirectiveOrConstruct. Or maybe
/// ACCDirectiveStmt and base class DirectiveStmt is a nice summary of a
/// directive that is more than declarative.
class ACCExecutableDirective : public ExecutableDirective {
friend class ASTStmtReader;
/// Kind of the directive.
......
......@@ -534,6 +534,9 @@ class Attr {
// Set to true if this attribute meaningful when applied to or inherited
// in a class template definition.
bit MeaningfulToClassTemplateDefinition = 0;
// Set to true if provides a custom implementation of the printPretty
// function.
bit HasCustomPrintPretty = 0;
// Set to true if this attribute can be used with '#pragma clang attribute'.
// By default, an attribute is supported by the '#pragma clang attribute'
// only when:
......@@ -3604,7 +3607,8 @@ def OMPDeclareTargetDecl : InheritableAttr {
EnumArgument<"DevType", "DevTypeTy",
[ "host", "nohost", "any" ],
[ "DT_Host", "DT_NoHost", "DT_Any" ]>,
UnsignedArgument<"Level">
UnsignedArgument<"Level">,
BoolArgument<"IsOpenACCTranslation">
];
let AdditionalMembers = [{
void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const;
......@@ -3614,6 +3618,7 @@ def OMPDeclareTargetDecl : InheritableAttr {
static llvm::Optional<DevTypeTy> getDeviceType(const ValueDecl *VD);
static llvm::Optional<SourceLocation> getLocation(const ValueDecl *VD);
}];
let HasCustomPrintPretty = 1;
}
def OMPAllocateDecl : InheritableAttr {
......@@ -3659,6 +3664,19 @@ def OMPDeclareVariant : InheritableAttr {
}];
}
def ACCRoutineDecl : InheritableAttr {
let Spellings = [Pragma<"acc", "routine">];
let SemaHandler = 0;
let Subjects = SubjectList<[Function]>;
let Args = [
EnumArgument<"Partitioning", "PartitioningTy",
[ "seq" ],
[ "Seq" ]>
];
let Documentation = [Undocumented];
let HasCustomPrintPretty = 1;
}
def Assumption : InheritableAttr {
let Spellings = [Clang<"assume">];
let Subjects = SubjectList<[Function, ObjCMethod]>;
......
......@@ -319,6 +319,15 @@ def err_rewrite_acc_end_in_macro : Error<
def err_rewrite_acc_end_in_pragma_op : Error<
"cannot rewrite OpenACC directive that has no associated statement and that "
"appears within a _Pragma operator">;
def err_rewrite_acc_routine_in_pragma_op : Error<
"cannot rewrite OpenACC routine directive that appears within a _Pragma "
"operator">;
def err_rewrite_acc_routine_function_start_in_macro : Error<
"cannot rewrite OpenACC routine directive whose associated function "
"declaration starts within a macro expansion">;
def err_rewrite_acc_routine_function_end_in_macro : Error<
"cannot rewrite OpenACC routine directive whose associated function "
"declaration ends within a macro expansion">;
} // end of RewriteOpenACC category
}
......@@ -10842,6 +10842,22 @@ def note_acc_disable_diag : Note<
def err_acc_no_self_host_device_clause : Error<
"expected at least one 'self', 'host', or 'device' clause for '#pragma acc "
"update'">;
// TODO: Once the other clauses are supported, this should say:
// "expected at least one 'gang', 'worker', 'vector', or 'seq' clause for "
// "'#pragma acc routine'"
def err_acc_no_gang_worker_vector_seq_clause : Error<
"expected 'seq' clause for '#pragma acc routine'">;
def err_acc_expected_function_after_directive : Error<
"'#pragma acc %0' must be followed by a lone function prototype or "
"definition">;
def err_acc_routine_unexpected_directive : Error<
"'#pragma acc %0' is not permitted within function '%1' because the latter "
"is attributed with '#pragma acc routine'">;
def err_acc_routine_static_local : Error<
"static local variable '%0' is not permitted within function '%1' because "
"the latter is attributed with '#pragma acc routine'">;
def note_acc_routine : Note<
"function '%0' attributed with '#pragma acc routine' here">;
} // end of OpenACC category
 
let CategoryName = "Related Result Type Issue" in {
......
......@@ -112,6 +112,9 @@
#ifndef OPENACC_PARALLEL_LOOP_CLAUSE
# define OPENACC_PARALLEL_LOOP_CLAUSE(Name)
#endif
#ifndef OPENACC_ROUTINE_CLAUSE
# define OPENACC_ROUTINE_CLAUSE(Name)
#endif
#ifndef OPENACC_UPDATE_PARENT
# define OPENACC_UPDATE_PARENT(Name)
#endif
......@@ -133,6 +136,9 @@
#ifndef OPENACC_PARALLEL_LOOP_PARENT
# define OPENACC_PARALLEL_LOOP_PARENT(Name)
#endif
#ifndef OPENACC_ROUTINE_PARENT
# define OPENACC_ROUTINE_PARENT(Name)
#endif
#define OPENACC_ALIASED_CLAUSE(Def, Name, Class) \
Def(Name, Class) \
......@@ -173,6 +179,7 @@ OPENACC_DIRECTIVE(data)
OPENACC_DIRECTIVE(parallel)
OPENACC_DIRECTIVE(loop)
OPENACC_DIRECTIVE_EXT(parallel_loop, "parallel loop")
OPENACC_DIRECTIVE(routine)
// Specifying and iterating OpenACC clauses, data attributes (DAs), data
// mapping attributes (DMAs), data sharing attributes (DSAs), and clause
......@@ -407,6 +414,9 @@ OPENACC_PARALLEL_LOOP_CLAUSE(worker)
OPENACC_PARALLEL_LOOP_CLAUSE(vector)
OPENACC_PARALLEL_LOOP_CLAUSE(collapse)
// Explicit clauses allowed for OpenACC directive 'routine'.
OPENACC_ROUTINE_CLAUSE(seq)
// Parent directives allowed for 'update'.
OPENACC_UPDATE_PARENT(unknown)
OPENACC_UPDATE_PARENT(data)
......@@ -440,6 +450,9 @@ OPENACC_LOOP_PARENT(parallel_loop)
OPENACC_PARALLEL_LOOP_PARENT(unknown)
OPENACC_PARALLEL_LOOP_PARENT(data)
// Parent directives allowed for 'routine'.
OPENACC_ROUTINE_PARENT(unknown)
#undef OPENACC_DIRECTIVE
#undef OPENACC_DIRECTIVE_EXT
#undef OPENACC_DMA
......@@ -473,6 +486,7 @@ OPENACC_PARALLEL_LOOP_PARENT(data)
#undef OPENACC_PARALLEL_CLAUSE
#undef OPENACC_LOOP_CLAUSE
#undef OPENACC_PARALLEL_LOOP_CLAUSE
#undef OPENACC_ROUTINE_CLAUSE
#undef OPENACC_UPDATE_PARENT
#undef OPENACC_ENTER_DATA_PARENT
#undef OPENACC_EXIT_DATA_PARENT
......@@ -480,6 +494,7 @@ OPENACC_PARALLEL_LOOP_PARENT(data)
#undef OPENACC_PARALLEL_PARENT
#undef OPENACC_LOOP_PARENT
#undef OPENACC_PARALLEL_LOOP_PARENT
#undef OPENACC_ROUTINE_PARENT
#undef OPENACC_ALIASED_CLAUSE
#undef OPENACC_CLAUSE_AND_DMA
#undef OPENACC_CLAUSE_AND_DSA_MAPPABLE
......
......@@ -3373,12 +3373,16 @@ public:
/// Parses declarative OpenACC directives.
DeclGroupPtrTy ParseOpenACCDeclarativeDirective();
/// Parses declarative or executable OpenACC directive.
/// Parses OpenACC executable directives or constructs.
///
/// \param StmtCtx The context in which we're parsing the directive.
StmtResult
ParseOpenACCDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx);
/// Parses clause of kind \a CKind for directive of a kind \a Kind.
StmtResult ParseOpenACCExecutableDirective(ParsedStmtContext StmtCtx);
/// Parses clauses for directive of kind \a Kind.
///
/// \param DKind Kind of current directive.
void ParseOpenACCClauses(OpenACCDirectiveKind DKind,
SmallVectorImpl<ACCClause *> &Clauses);
/// Parses clause of kind \a CKind for directive of kind \a Kind.
///
/// \param DKind Kind of current directive.
/// \param CKind Kind of current clause.
......
......@@ -11254,14 +11254,25 @@ private:
void DestroyOpenACCDirectiveStack();
 
public:
/// Called on start of new data attribute block.
void StartOpenACCDABlock(OpenACCDirectiveKind RealDKind, SourceLocation Loc);
/// Start analysis of clauses.
/// Called at start of an OpenACC directive before its clauses. Returns true
/// in the case of an error.
bool StartOpenACCDirectiveAndAssociate(OpenACCDirectiveKind RealDKind,
SourceLocation Loc);
/// Called at start of a clause.
void StartOpenACCClause(OpenACCClauseKind K);
/// End analysis of clauses.
/// Called at end of a clause.
void EndOpenACCClause();
/// Called on end of data attribute block.
void EndOpenACCDABlock();
/// Called at start of an OpenACC directive's associated statement. Returns
/// true in the case of an error.
bool StartOpenACCAssociatedStatement(OpenACCDirectiveKind DKind,
ArrayRef<ACCClause *> Clauses,
SourceLocation StartLoc);
/// Called at end of an OpenACC directive's associated statement and before
/// the \c ActOn function for the directive. Returns true in the case of an
/// error.
bool EndOpenACCAssociatedStatement();
/// Called after the \c ActOn function for the directive.
void EndOpenACCDirectiveAndAssociate();
 
/// If the current region is an OpenACC loop region, record any loop control
/// variables assigned but not declared in \p Init, the init of the attached
......@@ -11271,13 +11282,10 @@ public:
/// directive.
void ActOnOpenACCLoopBreakStatement(SourceLocation BreakLoc,
Scope *CurScope);
/// Check a function definition against any OpenACC restrictions (such as any
/// OpenACC routine directive).
void ActOnFunctionDefinitionForOpenACC(FunctionDecl *FD);
 
/// Start of OpenACC region.
bool ActOnOpenACCRegionStart(OpenACCDirectiveKind DKind,
ArrayRef<ACCClause *> Clauses,
SourceLocation StartLoc);
/// End of OpenACC region.
bool ActOnOpenACCRegionEnd();
StmtResult ActOnOpenACCExecutableDirective(
OpenACCDirectiveKind Kind, ArrayRef<ACCClause *> Clauses, Stmt *AStmt,
SourceLocation StartLoc, SourceLocation EndLoc);
......@@ -11314,6 +11322,10 @@ public:
StmtResult ActOnOpenACCParallelLoopDirective(
ArrayRef<ACCClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
SourceLocation EndLoc);
/// Called on well-formed '\#pragma acc routine'.
void ActOnOpenACCRoutineDirective(ArrayRef<ACCClause *> Clauses,
SourceLocation StartLoc,
SourceLocation EndLoc, DeclGroupRef Decl);
 
ACCClause *ActOnOpenACCSingleExprClause(OpenACCClauseKind Kind,
Expr *Expr,
......
......@@ -132,6 +132,32 @@ void OMPDeclareSimdDeclAttr::printPrettyPragma(
}
}
void OMPDeclareTargetDeclAttr::printPretty(raw_ostream &OS,
const PrintingPolicy &Policy) const {
// TODO: If this isn't an OpenACC translation, print exactly the way upstream
// prints it to avoid merge conflicts in tests. However, we need to come to
// terms with why inherited attributes should be printed as that seems to
// indicate they weren't specified in the original source. See related todo
// in DeclPrinter::VisitDeclContext.
OpenACCPrintKind PrintMode = OpenACCPrint_OMP;
if (getIsOpenACCTranslation())
PrintMode = isInherited() ? OpenACCPrint_ACC : Policy.OpenACCPrint;
switch (PrintMode) {
case OpenACCPrint_ACC_OMP:
OS << "// ";
LLVM_FALLTHROUGH;
case OpenACCPrint_OMP:
case OpenACCPrint_OMP_ACC:
case OpenACCPrint_OMP_HEAD:
OS << "#pragma omp declare target";
printPrettyPragma(OS, Policy);
OS << '\n';
break;
case OpenACCPrint_ACC:
break;
}
}
void OMPDeclareTargetDeclAttr::printPrettyPragma(
raw_ostream &OS, const PrintingPolicy &Policy) const {
// Use fake syntax because it is for testing and debugging purpose only.
......@@ -197,4 +223,23 @@ void OMPDeclareVariantAttr::printPrettyPragma(
OS << " match(" << traitInfos << ")";
}
void ACCRoutineDeclAttr::printPretty(raw_ostream &OS,
const PrintingPolicy &Policy) const {
if (isInherited() || isImplicit())
return;
switch (Policy.OpenACCPrint) {
case OpenACCPrint_OMP_ACC:
OS << "// ";
LLVM_FALLTHROUGH;
case OpenACCPrint_ACC:
case OpenACCPrint_ACC_OMP:
OS << "#pragma acc routine "
<< ConvertPartitioningTyToStr(getPartitioning()) << "\n";
break;
case OpenACCPrint_OMP:
case OpenACCPrint_OMP_HEAD:
break;
}
}
#include "clang/AST/AttrImpl.inc"
......@@ -488,8 +488,30 @@ void DeclPrinter::VisitDeclContext(DeclContext *DC, bool Indent) {
// Declare target attribute is special one, natural spelling for the pragma
// assumes "ending" construct so print it here.
if (D->hasAttr<OMPDeclareTargetDeclAttr>())
if (OMPDeclareTargetDeclAttr *Attr =
D->getAttr<OMPDeclareTargetDeclAttr>()) {
// TODO: If this isn't an OpenACC translation, print exactly the way
// upstream prints it to avoid merge conflicts in tests. However, we need
// to come to terms with why inherited attributes should be printed as
// that seems to indicate they weren't specified in the original source.
// See related todo in OMPDeclareTargetDeclAttr::printPretty.
OpenACCPrintKind PrintMode = OpenACCPrint_OMP;
if (Attr->getIsOpenACCTranslation())
PrintMode = Attr->isInherited() ? OpenACCPrint_ACC
: Policy.OpenACCPrint;
switch (PrintMode) {
case OpenACCPrint_ACC_OMP:
Out << "// ";
LLVM_FALLTHROUGH;
case OpenACCPrint_OMP:
case OpenACCPrint_OMP_ACC:
case OpenACCPrint_OMP_HEAD:
Out << "#pragma omp end declare target\n";
break;
case OpenACCPrint_ACC:
break;
}
}
}
if (!Decls.empty())
......
......@@ -198,6 +198,8 @@ bool clang::isAllowedDAForDirective(OpenACCDirectiveKind DKind,
break;
}
break;
case ACCD_routine:
break;
case ACCD_unknown:
llvm_unreachable("unexpected unknown directive");
}
......@@ -278,6 +280,8 @@ bool clang::isAllowedDAForDirective(OpenACCDirectiveKind DKind,
break;
}
break;
case ACCD_routine:
break;
case ACCD_unknown:
llvm_unreachable("unexpected unknown directive");
}
......@@ -353,6 +357,16 @@ bool clang::isAllowedClauseForDirective(OpenACCDirectiveKind DKind,
#define OPENACC_PARALLEL_LOOP_CLAUSE(Name) \
case ACCC_##Name: \
return true;
#include "clang/Basic/OpenACCKinds.def"
default:
break;
}
break;
case ACCD_routine:
switch (CKind) {
#define OPENACC_ROUTINE_CLAUSE(Name) \
case ACCC_##Name: \
return true;
#include "clang/Basic/OpenACCKinds.def"
default:
break;
......@@ -432,6 +446,16 @@ bool clang::isAllowedParentForDirective(OpenACCDirectiveKind DKind,
#define OPENACC_PARALLEL_LOOP_PARENT(Name) \
case ACCD_##Name: \
return true;
#include "clang/Basic/OpenACCKinds.def"
default:
break;
}
break;
case ACCD_routine:
switch (ParentDKind) {
#define OPENACC_ROUTINE_PARENT(Name) \
case ACCD_##Name: \
return true;
#include "clang/Basic/OpenACCKinds.def"
default:
break;
......@@ -451,6 +475,7 @@ int clang::getOpenACCEffectiveDirectives(OpenACCDirectiveKind DKind) {
case ACCD_data:
case ACCD_parallel:
case ACCD_loop:
case ACCD_routine:
return 1;
case ACCD_parallel_loop:
return 2;
......
......@@ -218,6 +218,7 @@ public:
std::string RewriteString;
llvm::raw_string_ostream RewriteStream(RewriteString);
// Generate new text.
switch (OpenACCPrint) {
case OpenACCPrint_ACC:
case OpenACCPrint_OMP_HEAD:
......@@ -374,6 +375,203 @@ public:
// associated statements separately.
return DirectiveOnly;
}
bool VisitDecl(Decl *D) {
// TODO: There's a lot of logic in here that could probably be shared with
// VisitACCExecutableDirective by encapsulating into separate functions.
SourceManager &SM = Context->getSourceManager();
const LangOptions &LO = Context->getLangOpts();
// Is this a function prototype or definition?
FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
if (!FD)
return true;
// Is there an OpenACC routine directive to rewrite?
ACCRoutineDeclAttr *ACCAttr = FD->getAttr<ACCRoutineDeclAttr>();
OMPDeclareTargetDeclAttr *OMPAttr = FD->getAttr<OMPDeclareTargetDeclAttr>();
if (!ACCAttr || !OMPAttr || !OMPAttr->getIsOpenACCTranslation())
return true;
assert(ACCAttr->isInherited() == OMPAttr->isInherited() &&
"expected OpenACC attribute and its OpenMP translation to either "
"both be inherited or neither");
if (ACCAttr->isInherited())
return true;
// Are we rewriting an OpenACC directive or adding OpenMP directives where
// there was no OpenACC directive?
bool RewriteDirective = !ACCAttr->isImplicit();
// Can we rewrite the directive?
//
// TODO: If the directive is in a _Pragma, we give up. Maybe we can try
// harder as we do in VisitACCExecutableDirective. See comments there for
// how all this works. To summarize: If the directive end location is not
// rewritable, apparently we have _Pragma form. Otherwise, we have #pragma
// form, whose begin and end locations are always rewritable.
SourceRange DirectiveRange = ACCAttr->getRange();
if (RewriteDirective) {
if (!Rewrite.isRewritable(DirectiveRange.getEnd())) {
Context->getDiagnostics().Report(
DirectiveRange.getEnd(),
diag::err_rewrite_acc_routine_in_pragma_op);
return true;
}
assert(Rewrite.isRewritable(DirectiveRange.getBegin()) &&
"expected start of directive to be rewritable because end is");
}
// Can we insert before the associated function prototype or definition?
SourceLocation FDBegin = FD->getBeginLoc();
if (!RewriteDirective && !Rewrite.isRewritable(FDBegin)) {
Context->getDiagnostics().Report(
FDBegin, diag::err_rewrite_acc_routine_function_start_in_macro);
return true;
}
// Can we insert after the associated function prototype or definition?
//
// Function definitions are easy as the end location is already correct.
SourceLocation FDEnd = FD->getEndLoc();
if (FD != FD->getDefinition()) {
#ifndef NDEBUG
Token Tok;
Lexer::getRawToken(SM.getSpellingLoc(FDEnd), Tok, SM, LO);
assert(Tok.getKind() != tok::semi &&
"expected FunctionDecl's end token not to be a semicolon");
#endif
Optional<Token> Next = Lexer::findNextToken(FDEnd, SM, LO);
// If Tok is expanded from a macro and is not the last token in the
// expansion, findNextToken refuses to look for the next token and returns
// None. Assume the final semicolon is the next token and is thus within
// the expansion, and so refuse to rewrite.
if (!Next.hasValue()) {
Context->getDiagnostics().Report(
FDEnd, diag::err_rewrite_acc_routine_function_end_in_macro);
return true;
}
// If Next is an identifier, assume it's a macro whose expansion's
// first token is the final semicolon, and so refuse to rewrite.
if (Next.getValue().getKind() == tok::raw_identifier) {
Context->getDiagnostics().Report(
Next.getValue().getLocation(),
diag::err_rewrite_acc_routine_function_end_in_macro);
return true;
}
assert(Next.getValue().getKind() == tok::semi &&
"expected semicolon after FunctionDecl");
FDEnd = Next.getValue().getLocation();
assert(SM.getCharacterData(FDEnd)[0] == ';' &&
"expected tok::semi to look like a semicolon");
assert(Rewrite.isRewritable(FDEnd) &&
"expected Lexer::findNextToken not to return token within macro "
"expansion");
} else {
#ifndef NDEBUG
Token Tok;
Lexer::getRawToken(SM.getSpellingLoc(FDEnd), Tok, SM, LO);
assert(Tok.getKind() == tok::r_brace &&
"expected FunctionDecl's end token to be a right brace");
assert(SM.getCharacterData(FDEnd)[0] == '}' &&
"expected tok::r_brace to look like a right brace");
#endif
if (!Rewrite.isRewritable(FDEnd)) {
Context->getDiagnostics().Report(
FDEnd, diag::err_rewrite_acc_routine_function_end_in_macro);
return true;
}
}
// What is the existing indentation?
StringRef IndentText;
if (RewriteDirective)
IndentText = Lexer::getIndentationForLine(DirectiveRange.getBegin(), SM);
else
IndentText = Lexer::getIndentationForLine(FDBegin, SM);
// Set up buffers for the new text.
std::string RewriteBeginString;
std::string RewriteEndString;
llvm::raw_string_ostream RewriteBeginStream(RewriteBeginString);
llvm::raw_string_ostream RewriteEndStream(RewriteEndString);
// Generate new text.
PrintingPolicy PolicyOMP(LO);
PolicyOMP.OpenACCPrint = OpenACCPrint_OMP;
if (!RewriteDirective) {
// If the rest of the line preceding the function declaration isn't
// blank, insert a newline before the inserted directive. Otherwise,
// don't or you'll create a blank line.
FileID File = SM.getFileID(FDBegin);
StringRef Buffer = SM.getBufferData(File);
const char *P = SM.getCharacterData(FDBegin);
if (P != Buffer.begin())
--P;
for (; P != Buffer.begin() && *P !='\n'; --P) {
if (!std::isspace(*P)) {
RewriteBeginStream << '\n' << IndentText;
break;
}
}
}
RewriteEndStream << '\n' << IndentText;
switch (OpenACCPrint) {
case OpenACCPrint_ACC:
case OpenACCPrint_OMP_HEAD:
llvm_unreachable("unexpected OpenACC print kind while rewriting input");
case OpenACCPrint_OMP:
OMPAttr->printPretty(RewriteBeginStream, PolicyOMP);
break;
case OpenACCPrint_OMP_ACC:
OMPAttr->printPretty(RewriteBeginStream, PolicyOMP);
if (RewriteDirective) {
RewriteBeginStream << IndentText << "// " << Rewrite.getRewrittenText(
CharSourceRange::getCharRange(DirectiveRange));
}
break;