ffrankies updated this revision to Diff 342467.
ffrankies added a comment.
Addressed comment by @aaron.ballman and the Pre-update checks linter
- Removed calls to `std::move` for `llvm::Twine::str()` object in
`IdDependencyRecord` constructors
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D70094/new/
https://reviews.llvm.org/D70094
Files:
clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
clang-tools-extra/clang-tidy/altera/CMakeLists.txt
clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp
clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h
clang-tools-extra/docs/ReleaseNotes.rst
clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst
clang-tools-extra/docs/clang-tidy/checks/list.rst
clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp
Index: clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp
@@ -0,0 +1,86 @@
+// RUN: %check_clang_tidy %s altera-id-dependent-backward-branch %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h
+
+typedef struct ExampleStruct {
+ int IDDepField;
+} ExampleStruct;
+
+void error() {
+ // ==== Conditional Expressions ====
+ int accumulator = 0;
+ for (int i = 0; i < get_local_id(0); i++) {
+ // CHECK-NOTES: :[[@LINE-1]]:19: warning: backward branch (for loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ int j = 0;
+ while (j < get_local_id(0)) {
+ // CHECK-NOTES: :[[@LINE-1]]:10: warning: backward branch (while loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ do {
+ accumulator++;
+ } while (j < get_local_id(0));
+ // CHECK-NOTES: :[[@LINE-1]]:12: warning: backward branch (do loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch]
+
+ // ==== Assignments ====
+ int ThreadID = get_local_id(0);
+
+ while (j < ThreadID) {
+ // CHECK-NOTES: :[[@LINE-3]]:3: note: assignment of ID-dependent variable ThreadID
+ // CHECK-NOTES: :[[@LINE-2]]:10: warning: backward branch (while loop) is ID-dependent due to variable reference to 'ThreadID' and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ ExampleStruct Example;
+ Example.IDDepField = get_local_id(0);
+
+ // ==== Inferred Assignments ====
+ int ThreadID2 = ThreadID * get_local_size(0);
+
+ int ThreadID3 = Example.IDDepField; // OK: not used in any loops
+
+ ExampleStruct UnusedStruct = {
+ ThreadID * 2 // OK: not used in any loops
+ };
+
+ for (int i = 0; i < ThreadID2; i++) {
+ // CHECK-NOTES: :[[@LINE-9]]:3: note: inferred assignment of ID-dependent value from ID-dependent variable ThreadID
+ // CHECK-NOTES: :[[@LINE-2]]:19: warning: backward branch (for loop) is ID-dependent due to variable reference to 'ThreadID2' and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ do {
+ accumulator++;
+ } while (j < ThreadID);
+ // CHECK-NOTES: :[[@LINE-29]]:3: note: assignment of ID-dependent variable ThreadID
+ // CHECK-NOTES: :[[@LINE-2]]:12: warning: backward branch (do loop) is ID-dependent due to variable reference to 'ThreadID' and may cause performance degradation [altera-id-dependent-backward-branch]
+
+ for (int i = 0; i < Example.IDDepField; i++) {
+ // CHECK-NOTES: :[[@LINE-24]]:3: note: assignment of ID-dependent field IDDepField
+ // CHECK-NOTES: :[[@LINE-2]]:19: warning: backward branch (for loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ while (j < Example.IDDepField) {
+ // CHECK-NOTES: :[[@LINE-30]]:3: note: assignment of ID-dependent field IDDepField
+ // CHECK-NOTES: :[[@LINE-2]]:10: warning: backward branch (while loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch]
+ accumulator++;
+ }
+
+ do {
+ accumulator++;
+ } while (j < Example.IDDepField);
+ // CHECK-NOTES: :[[@LINE-38]]:3: note: assignment of ID-dependent field IDDepField
+ // CHECK-NOTES: :[[@LINE-2]]:12: warning: backward branch (do loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch]
+}
+
+void success() {
+ int accumulator = 0;
+
+ for (int i = 0; i < 1000; i++) {
+ if (i < get_local_id(0)) {
+ accumulator++;
+ }
+ }
+}
Index: clang-tools-extra/docs/clang-tidy/checks/list.rst
===================================================================
--- clang-tools-extra/docs/clang-tidy/checks/list.rst
+++ clang-tools-extra/docs/clang-tidy/checks/list.rst
@@ -30,6 +30,7 @@
`abseil-time-comparison <abseil-time-comparison.html>`_, "Yes"
`abseil-time-subtraction <abseil-time-subtraction.html>`_, "Yes"
`abseil-upgrade-duration-conversions <abseil-upgrade-duration-conversions.html>`_, "Yes"
+ `altera-id-dependent-backward-branch <altera-id-dependent-backward-branch.html>`_,
`altera-kernel-name-restriction <altera-kernel-name-restriction.html>`_,
`altera-single-work-item-barrier <altera-single-work-item-barrier.html>`_,
`altera-struct-pack-align <altera-struct-pack-align.html>`_, "Yes"
Index: clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst
===================================================================
--- /dev/null
+++ clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst
@@ -0,0 +1,28 @@
+.. title:: clang-tidy - altera-id-dependent-backward-branch
+
+altera-id-dependent-backward-branch
+===================================
+
+Finds ID-dependent variables and fields that are used within loops. This causes
+branches to occur inside the loops, and thus leads to performance degradation.
+
+.. code-block:: c++
+
+ // The following code will produce a warning because this ID-dependent
+ // variable is used in a loop condition statement.
+ int ThreadID = get_local_id(0);
+
+ // The following loop will produce a warning because the loop condition
+ // statement depends on an ID-dependent variable.
+ for (int i = 0; i < ThreadID; ++i) {
+ std::cout << i << std::endl;
+ }
+
+ // The following loop will not produce a warning, because the ID-dependent
+ // variable is not used in the loop condition statement.
+ for (int i = 0; i < 100; ++i) {
+ std::cout << ThreadID << std::endl;
+ }
+
+Based on the `Altera SDK for OpenCL: Best Practices Guide
+<https://www.altera.com/en_US/pdfs/literature/hb/opencl-sdk/aocl_optimization_guide.pdf>`_.
Index: clang-tools-extra/docs/ReleaseNotes.rst
===================================================================
--- clang-tools-extra/docs/ReleaseNotes.rst
+++ clang-tools-extra/docs/ReleaseNotes.rst
@@ -83,6 +83,13 @@
Finds ``pthread_setcanceltype`` function calls where a thread's cancellation
type is set to asynchronous.
+- New :doc:`altera-id-dependent-backward-branch
+ <clang-tidy/checks/altera-id-dependent-backward-branch>` check.
+
+ Finds ID-dependent variables and fields that are used within loops. This
+ causes branches to occur inside the loops, and thus leads to performance
+ degradation.
+
- New :doc:`altera-unroll-loops
<clang-tidy/checks/altera-unroll-loops>` check.
Index: clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h
@@ -0,0 +1,83 @@
+//===--- IdDependentBackwardBranchCheck.h - clang-tidy ----------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_IDDEPENDENTBACKWARDBRANCHCHECK_H
+#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_IDDEPENDENTBACKWARDBRANCHCHECK_H
+
+#include "../ClangTidyCheck.h"
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+/// Finds ID-dependent variables and fields used within loops, and warns of
+/// their usage. Using these variables in loops can lead to performance
+/// degradation.
+///
+/// For the user-facing documentation see:
+/// http://clang.llvm.org/extra/clang-tidy/checks/altera-id-dependent-backward-branch.html
+class IdDependentBackwardBranchCheck : public ClangTidyCheck {
+private:
+ enum LoopType { UnknownLoop = -1, DoLoop = 0, WhileLoop = 1, ForLoop = 2 };
+ // Stores information necessary for printing out source of error.
+ struct IdDependencyRecord {
+ IdDependencyRecord(const VarDecl *Declaration, SourceLocation Location,
+ const llvm::Twine &Message)
+ : VariableDeclaration(Declaration), Location(Location),
+ Message(Message.str()) {}
+ IdDependencyRecord(const FieldDecl *Declaration, SourceLocation Location,
+ const llvm::Twine &Message)
+ : FieldDeclaration(Declaration), Location(Location),
+ Message(Message.str()) {}
+ IdDependencyRecord() = default;
+ const VarDecl *VariableDeclaration = nullptr;
+ const FieldDecl *FieldDeclaration = nullptr;
+ SourceLocation Location;
+ std::string Message;
+ };
+ // Stores the locations where ID-dependent variables are created.
+ std::map<const VarDecl *, IdDependencyRecord> IdDepVarsMap;
+ // Stores the locations where ID-dependent fields are created.
+ std::map<const FieldDecl *, IdDependencyRecord> IdDepFieldsMap;
+ /// Returns an IdDependencyRecord if the Expression contains an ID-dependent
+ /// variable, returns a nullptr otherwise.
+ IdDependencyRecord *hasIdDepVar(const Expr *Expression);
+ /// Returns an IdDependencyRecord if the Expression contains an ID-dependent
+ /// field, returns a nullptr otherwise.
+ IdDependencyRecord *hasIdDepField(const Expr *Expression);
+ /// Stores the location an ID-dependent variable is created from a call to
+ /// an ID function in IdDepVarsMap.
+ void saveIdDepVar(const Stmt *Statement, const VarDecl *Variable);
+ /// Stores the location an ID-dependent field is created from a call to an ID
+ /// function in IdDepFieldsMap.
+ void saveIdDepField(const Stmt *Statement, const FieldDecl *Field);
+ /// Stores the location an ID-dependent variable is created from a reference
+ /// to another ID-dependent variable or field in IdDepVarsMap.
+ void saveIdDepVarFromReference(const DeclRefExpr *RefExpr,
+ const MemberExpr *MemExpr,
+ const VarDecl *PotentialVar);
+ /// Stores the location an ID-dependent field is created from a reference to
+ /// another ID-dependent variable or field in IdDepFieldsMap.
+ void saveIdDepFieldFromReference(const DeclRefExpr *RefExpr,
+ const MemberExpr *MemExpr,
+ const FieldDecl *PotentialField);
+ /// Returns the loop type.
+ LoopType getLoopType(const Stmt *Loop);
+
+public:
+ IdDependentBackwardBranchCheck(StringRef Name, ClangTidyContext *Context)
+ : ClangTidyCheck(Name, Context) {}
+ void registerMatchers(ast_matchers::MatchFinder *Finder) override;
+ void check(const ast_matchers::MatchFinder::MatchResult &Result) override;
+};
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang
+
+#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_IDDEPENDENTBACKWARDBRANCHCHECK_H
Index: clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp
@@ -0,0 +1,264 @@
+//===--- IdDependentBackwardBranchCheck.cpp - clang-tidy ------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "IdDependentBackwardBranchCheck.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+void IdDependentBackwardBranchCheck::registerMatchers(MatchFinder *Finder) {
+ // Prototype to identify all variables which hold a thread-variant ID.
+ // First Matcher just finds all the direct assignments of either ID call.
+ const auto ThreadID = expr(hasDescendant(callExpr(callee(functionDecl(
+ anyOf(hasName("get_global_id"), hasName("get_local_id")))))));
+
+ const auto RefVarOrField = forEachDescendant(
+ stmt(anyOf(declRefExpr(to(varDecl())).bind("assign_ref_var"),
+ memberExpr(member(fieldDecl())).bind("assign_ref_field"))));
+
+ Finder->addMatcher(
+ compoundStmt(
+ // Bind on actual get_local/global_id calls.
+ forEachDescendant(
+ stmt(
+ anyOf(declStmt(hasDescendant(varDecl(hasInitializer(ThreadID))
+ .bind("tid_dep_var"))),
+ binaryOperator(allOf(
+ isAssignmentOperator(), hasRHS(ThreadID),
+ hasLHS(anyOf(
+ declRefExpr(to(varDecl().bind("tid_dep_var"))),
+ memberExpr(member(
+ fieldDecl().bind("tid_dep_field")))))))))
+ .bind("straight_assignment"))),
+ this);
+
+ // Bind all VarDecls that include an initializer with a variable DeclRefExpr
+ // (in case it is ID-dependent).
+ Finder->addMatcher(
+ stmt(forEachDescendant(
+ varDecl(hasInitializer(RefVarOrField)).bind("pot_tid_var"))),
+ this);
+
+ // Bind all VarDecls that are assigned a value with a variable DeclRefExpr (in
+ // case it is ID-dependent).
+ Finder->addMatcher(
+ stmt(forEachDescendant(binaryOperator(
+ allOf(isAssignmentOperator(), hasRHS(RefVarOrField),
+ hasLHS(anyOf(
+ declRefExpr(to(varDecl().bind("pot_tid_var"))),
+ memberExpr(member(fieldDecl().bind("pot_tid_field"))))))))),
+ this);
+
+ // Second Matcher looks for branch statements inside of loops and bind on the
+ // condition expression IF it either calls an ID function or has a variable
+ // DeclRefExpr. DeclRefExprs are checked later to confirm whether the variable
+ // is ID-dependent.
+ const auto CondExpr =
+ expr(anyOf(hasDescendant(callExpr(callee(functionDecl(
+ anyOf(hasName("get_global_id"),
+ hasName("get_local_id")))))
+ .bind("id_call")),
+ hasDescendant(stmt(anyOf(declRefExpr(to(varDecl())),
+ memberExpr(member(fieldDecl())))))))
+ .bind("cond_expr");
+ Finder->addMatcher(stmt(anyOf(forStmt(hasCondition(CondExpr)),
+ doStmt(hasCondition(CondExpr)),
+ whileStmt(hasCondition(CondExpr))))
+ .bind("backward_branch"),
+ this);
+}
+
+IdDependentBackwardBranchCheck::IdDependencyRecord *
+IdDependentBackwardBranchCheck::hasIdDepVar(const Expr *Expression) {
+ if (const auto *Declaration = dyn_cast<DeclRefExpr>(Expression)) {
+ // It is a DeclRefExpr, so check if it's an ID-dependent variable.
+ const auto *CheckVariable = dyn_cast<VarDecl>(Declaration->getDecl());
+ auto FoundVariable = IdDepVarsMap.find(CheckVariable);
+ if (FoundVariable == IdDepVarsMap.end())
+ return nullptr;
+ return &(FoundVariable->second);
+ }
+ for (const auto *Child : Expression->children())
+ if (const auto *ChildExpression = dyn_cast<Expr>(Child))
+ if (IdDependencyRecord *Result = hasIdDepVar(ChildExpression))
+ return Result;
+ return nullptr;
+}
+
+IdDependentBackwardBranchCheck::IdDependencyRecord *
+IdDependentBackwardBranchCheck::hasIdDepField(const Expr *Expression) {
+ if (const auto *MemberExpression = dyn_cast<MemberExpr>(Expression)) {
+ const auto *CheckField =
+ dyn_cast<FieldDecl>(MemberExpression->getMemberDecl());
+ auto FoundField = IdDepFieldsMap.find(CheckField);
+ if (FoundField == IdDepFieldsMap.end())
+ return nullptr;
+ return &(FoundField->second);
+ }
+ for (const auto *Child : Expression->children())
+ if (const auto *ChildExpression = dyn_cast<Expr>(Child))
+ if (IdDependencyRecord *Result = hasIdDepField(ChildExpression))
+ return Result;
+ return nullptr;
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepVar(const Stmt *Statement,
+ const VarDecl *Variable) {
+ // Record that this variable is thread-dependent.
+ IdDepVarsMap[Variable] =
+ IdDependencyRecord(Variable, Variable->getBeginLoc(),
+ Twine("assignment of ID-dependent variable ") +
+ Variable->getNameAsString());
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepField(const Stmt *Statement,
+ const FieldDecl *Field) {
+ // Record that this field is thread-dependent.
+ IdDepFieldsMap[Field] = IdDependencyRecord(
+ Field, Statement->getBeginLoc(),
+ Twine("assignment of ID-dependent field ") + Field->getNameAsString());
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepVarFromReference(
+ const DeclRefExpr *RefExpr, const MemberExpr *MemExpr,
+ const VarDecl *PotentialVar) {
+ // If the variable is already in IdDepVarsMap, ignore it.
+ if (IdDepVarsMap.find(PotentialVar) != IdDepVarsMap.end())
+ return;
+ std::string Message;
+ llvm::raw_string_ostream StringStream(Message);
+ StringStream << "inferred assignment of ID-dependent value from "
+ "ID-dependent ";
+ if (RefExpr) {
+ const auto *RefVar = dyn_cast<VarDecl>(RefExpr->getDecl());
+ // If variable isn't ID-dependent, but RefVar is.
+ if (IdDepVarsMap.find(RefVar) != IdDepVarsMap.end())
+ StringStream << "variable " << RefVar->getNameAsString();
+ }
+ if (MemExpr) {
+ const auto *RefField = dyn_cast<FieldDecl>(MemExpr->getMemberDecl());
+ // If variable isn't ID-dependent, but RefField is.
+ if (IdDepFieldsMap.find(RefField) != IdDepFieldsMap.end())
+ StringStream << "member " << RefField->getNameAsString();
+ }
+ IdDepVarsMap[PotentialVar] =
+ IdDependencyRecord(PotentialVar, PotentialVar->getBeginLoc(), Message);
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepFieldFromReference(
+ const DeclRefExpr *RefExpr, const MemberExpr *MemExpr,
+ const FieldDecl *PotentialField) {
+ // If the field is already in IdDepFieldsMap, ignore it.
+ if (IdDepFieldsMap.find(PotentialField) != IdDepFieldsMap.end())
+ return;
+ std::string Message;
+ llvm::raw_string_ostream StringStream(Message);
+ StringStream << "inferred assignment of ID-dependent member from "
+ "ID-dependent ";
+ if (RefExpr) {
+ const auto *RefVar = dyn_cast<VarDecl>(RefExpr->getDecl());
+ // If field isn't ID-dependent, but RefVar is.
+ if (IdDepVarsMap.find(RefVar) != IdDepVarsMap.end())
+ StringStream << "variable " << RefVar->getNameAsString();
+ }
+ if (MemExpr) {
+ const auto *RefField = dyn_cast<FieldDecl>(MemExpr->getMemberDecl());
+ if (IdDepFieldsMap.find(RefField) != IdDepFieldsMap.end())
+ StringStream << "member " << RefField->getNameAsString();
+ }
+ IdDepFieldsMap[PotentialField] = IdDependencyRecord(
+ PotentialField, PotentialField->getBeginLoc(), Message);
+}
+
+IdDependentBackwardBranchCheck::LoopType
+IdDependentBackwardBranchCheck::getLoopType(const Stmt *Loop) {
+ switch (Loop->getStmtClass()) {
+ case Stmt::DoStmtClass:
+ return DoLoop;
+ case Stmt::WhileStmtClass:
+ return WhileLoop;
+ case Stmt::ForStmtClass:
+ return ForLoop;
+ default:
+ return UnknownLoop;
+ }
+}
+
+void IdDependentBackwardBranchCheck::check(
+ const MatchFinder::MatchResult &Result) {
+ // The first half of the callback only deals with identifying and storing
+ // ID-dependency information into the IdDepVars and IdDepFields maps.
+ const auto *Variable = Result.Nodes.getNodeAs<VarDecl>("tid_dep_var");
+ const auto *Field = Result.Nodes.getNodeAs<FieldDecl>("tid_dep_field");
+ const auto *Statement = Result.Nodes.getNodeAs<Stmt>("straight_assignment");
+ const auto *RefExpr = Result.Nodes.getNodeAs<DeclRefExpr>("assign_ref_var");
+ const auto *MemExpr = Result.Nodes.getNodeAs<MemberExpr>("assign_ref_field");
+ const auto *PotentialVar = Result.Nodes.getNodeAs<VarDecl>("pot_tid_var");
+ const auto *PotentialField =
+ Result.Nodes.getNodeAs<FieldDecl>("pot_tid_field");
+
+ // Save variables and fields assigned directly through ID function calls.
+ if (Statement && (Variable || Field)) {
+ if (Variable)
+ saveIdDepVar(Statement, Variable);
+ else if (Field)
+ saveIdDepField(Statement, Field);
+ }
+
+ // Save variables assigned to values of Id-dependent variables and fields.
+ if ((RefExpr || MemExpr) && PotentialVar)
+ saveIdDepVarFromReference(RefExpr, MemExpr, PotentialVar);
+
+ // Save fields assigned to values of ID-dependent variables and fields.
+ if ((RefExpr || MemExpr) && PotentialField)
+ saveIdDepFieldFromReference(RefExpr, MemExpr, PotentialField);
+
+ // The second part of the callback deals with checking if a branch inside a
+ // loop is thread dependent.
+ const auto *CondExpr = Result.Nodes.getNodeAs<Expr>("cond_expr");
+ const auto *IDCall = Result.Nodes.getNodeAs<CallExpr>("id_call");
+ const auto *Loop = Result.Nodes.getNodeAs<Stmt>("backward_branch");
+ if (!Loop)
+ return;
+ LoopType Type = getLoopType(Loop);
+ if (CondExpr) {
+ if (IDCall) { // Conditional expression calls an ID function directly.
+ diag(CondExpr->getBeginLoc(),
+ "backward branch (%select{do|while|for}0 loop) is ID-dependent due "
+ "to ID function call and may cause performance degradation")
+ << Type;
+ return;
+ }
+ // Conditional expression has DeclRefExpr(s), check ID-dependency.
+ IdDependencyRecord *IdDepVar = hasIdDepVar(CondExpr);
+ IdDependencyRecord *IdDepField = hasIdDepField(CondExpr);
+ if (IdDepVar) {
+ // Change one of these to a Note
+ diag(IdDepVar->Location, IdDepVar->Message, DiagnosticIDs::Note);
+ diag(CondExpr->getBeginLoc(),
+ "backward branch (%select{do|while|for}0 loop) is ID-dependent due "
+ "to variable reference to %1 and may cause performance degradation")
+ << Type << IdDepVar->VariableDeclaration;
+ } else if (IdDepField) {
+ diag(IdDepField->Location, IdDepField->Message, DiagnosticIDs::Note);
+ diag(CondExpr->getBeginLoc(),
+ "backward branch (%select{do|while|for}0 loop) is ID-dependent due "
+ "to member reference to %1 and may cause performance degradation")
+ << Type << IdDepField->FieldDeclaration;
+ }
+ }
+}
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang
Index: clang-tools-extra/clang-tidy/altera/CMakeLists.txt
===================================================================
--- clang-tools-extra/clang-tidy/altera/CMakeLists.txt
+++ clang-tools-extra/clang-tidy/altera/CMakeLists.txt
@@ -5,6 +5,7 @@
add_clang_library(clangTidyAlteraModule
AlteraTidyModule.cpp
+ IdDependentBackwardBranchCheck.cpp
KernelNameRestrictionCheck.cpp
SingleWorkItemBarrierCheck.cpp
StructPackAlignCheck.cpp
Index: clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
===================================================================
--- clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
+++ clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
@@ -9,6 +9,7 @@
#include "../ClangTidy.h"
#include "../ClangTidyModule.h"
#include "../ClangTidyModuleRegistry.h"
+#include "IdDependentBackwardBranchCheck.h"
#include "KernelNameRestrictionCheck.h"
#include "SingleWorkItemBarrierCheck.h"
#include "StructPackAlignCheck.h"
@@ -23,6 +24,8 @@
class AlteraModule : public ClangTidyModule {
public:
void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
+ CheckFactories.registerCheck<IdDependentBackwardBranchCheck>(
+ "altera-id-dependent-backward-branch");
CheckFactories.registerCheck<KernelNameRestrictionCheck>(
"altera-kernel-name-restriction");
CheckFactories.registerCheck<SingleWorkItemBarrierCheck>(
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
