[clang-tidy] new altera ID dependent backward branch check

[ffrankies]

Summary

This lint check is a part of the FLOCL (FPGA Linters for OpenCL) project out of the Synergy Lab at Virginia Tech.

FLOCL is a set of lint checks aimed at FPGA developers who write code in OpenCL.

The altera ID dependent backward branch lint check finds ID dependent variables and fields used within loops, and warns of their usage. Using these variables in loops can lead to performance degradation.

[Eugene.Zelenko]

clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

106 ↗

(On Diff #228743)

Please don't use auto when type is not explicitly spelled in same statement or iterator.

129 ↗

(On Diff #228743)

Please don't use auto when type is not explicitly spelled in same statement or iterator.

[ffrankies]

Implemented changes requested by @Eugene.Zelenko

Also changed for loop in hasIdDepVar and hasIdDepField to range-based for loops, and updated the license information in IdDependentBackwardBranchCheck.cpp

[Eugene.Zelenko]

clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

115 ↗

(On Diff #231309)

May be comments were shifted with code, but here const auto * could be used, because type is spelled in dyn_cast<>.

116 ↗

(On Diff #231309)

May be comments were shifted with code, but here const auto * could be used, because type is spelled in dyn_cast<>.

124 ↗

(On Diff #231309)

May be comments were shifted with code, but here const auto * could be used, because type is iterator.

125 ↗

(On Diff #231309)

May be comments were shifted with code, but here const auto * could be used, because type is spelled in dyn_cast<>.

165 ↗

(On Diff #231309)

const auto *. Same in other dyn_cast<>.

[ffrankies]

Addressed @Eugene.Zelenko's comments on use of auto

[ffrankies]

• Updated underlying repo to https://github.com/llvm/llvm-project
• Removed braces from one-line if-statements

[ffrankies]

• Rebased on top of latest changes in main branch
• The diagnostic that identifies the code location where an ID-dependent variable/field is assigned has been changed from a warning to a note
• Changes addressing code style

[Eugene.Zelenko]

clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst

10

Usually link to external documentation is placed at the end. See other checks documentation as example.

[ffrankies]

• Addressed comments from the automated pre-merge checks
• Moved link to external documentation to the end of clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst as requested by @Eugene.Zelenko

[Eugene.Zelenko]

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

201

Is loop ID is not enough? Why does comment with numerical code used? Same below.

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h

2

Please add space after clang-tidy and language code. See other checks as example.

37

Please use = default;

[ffrankies]

Addressed comments by @Eugene.Zelenko and the automated build bot

• Fixed header comments and include guard style
• Removed unnecessary comments in getLoopType()
• changed IDDependencyRecord() {} to IDDependencyRecord() = default;

[Eugene.Zelenko]

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

2

*- C++ -* is not necessary for .cpp files.

[ffrankies]

Removed *- C++ -* from IdDependentBackwardBranchCheck.cpp file-level comment.

[ffrankies]

@Eugene.Zelenko @aaron.ballman Are there any more changes that need to be made to this check or comments that need to be addressed?

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

201

Removed the comment with numerical code (I simply added it to avoid having to check the header file to see their numerical value). The LoopType is used in the diagnostics to select and emit the correct loop type as part of the diagnostic message, e.g.:

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;

I'm assuming the loop ID is a unique object identifier, so I don't think it'll serve the same purpose.

[aaron.ballman]

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

29–32

Can you use the isAssignmentOperator() AST matcher instead of spelling out the operator names manually?

55

102–103

122–123

133–135

You should use an llvm::Twine for this.

142–144

Same here as above.

155

I'd use llvm::raw_string_ostream here (basically, try to avoid STL iostream functionality as much as possible, it's really slow/bad).

200–205

Rather than repeatedly use isa<>, I'd consider using a switch over Loop->getStmtClass(). This also addresses the "no else after a return" coding style concern.

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h

27

The enumerators don't match the style guide -- can you rename them to UnknownLoop, DoLoop, etc so they look less like macros?

33

36

38–39

To ensure they get initialized to something sensible from the default ctor.

[ffrankies]

Addressed comments by @aaron.ballman

• Changed capitalization in enum
• Used std::move in IdDependencyRecord constructors
• Initialized VariableDeclaration and FieldDeclaration to nullptr
• Used isAssignmentOperator() instead of listing the assingment operators in the matchers
• Simplified code around if statement expressions
• Switched to llvm::Twine and llvm::raw_string_ostream when constructing warning and note messages
• Changed getLoopType() body to a switch statement instead of a series of if-else statements

[ffrankies]

Removed unused import statements from IdDependentBackwardBranch.cpp

[aaron.ballman]

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h

33

Ah, it looks like some changes conflicted with my suggestion -- when the function was taking a std::string, the move was needed, but now with a Twine, the move is an issue. You should remove the std::move.

36

Same here.

[ffrankies]

Addressed comment by @aaron.ballman and the Pre-update checks linter

• Removed calls to std::move for llvm::Twine::str() object in IdDependencyRecord constructors

[aaron.ballman]

LGTM! Btw, given that you've had several of these approved, would you be interested in getting commit access (https://llvm.org/docs/DeveloperPolicy.html#obtaining-commit-access)?

[ffrankies]

@aaron.ballman Thank you! As far as I'm aware, this is the last check that we are planning to submit, so if I do get commit access now it's likely to be unused. However, if that does change, then yes I would be interested in obtaining commit access. For now, can you please commit this on my behalf? My github username is ffrankies.

[aaron.ballman]

In D70094#2742760, @ffrankies wrote:

@aaron.ballman Thank you! As far as I'm aware, this is the last check that we are planning to submit, so if I do get commit access now it's likely to be unused. However, if that does change, then yes I would be interested in obtaining commit access. For now, can you please commit this on my behalf? My github username is ffrankies.

Ah, good to know! Thank you for the new check -- I've commit on your behalf in 83af66e18e3d3760d56ea7e3bdbff3428ae7730d

clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp

 //===--- AlteraTidyModule.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 "../ClangTidy.h"
 #include "../ClangTidyModule.h"
 #include "../ClangTidyModuleRegistry.h"
+#include "IdDependentBackwardBranchCheck.h"
 #include "KernelNameRestrictionCheck.h"
 #include "SingleWorkItemBarrierCheck.h"
 #include "StructPackAlignCheck.h"
 #include "UnrollLoopsCheck.h"
 
 using namespace clang::ast_matchers;
 
 namespace clang {
 namespace tidy {
 namespace altera {
 
 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>(
         "altera-single-work-item-barrier");
     CheckFactories.registerCheck<StructPackAlignCheck>(
         "altera-struct-pack-align");
     CheckFactories.registerCheck<UnrollLoopsCheck>("altera-unroll-loops");
   }

clang-tools-extra/clang-tidy/altera/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   FrontendOpenMP
   support
   )
 
 add_clang_library(clangTidyAlteraModule
   AlteraTidyModule.cpp
+  IdDependentBackwardBranchCheck.cpp
   KernelNameRestrictionCheck.cpp
   SingleWorkItemBarrierCheck.cpp
   StructPackAlignCheck.cpp
   UnrollLoopsCheck.cpp
 
   LINK_LIBS
   clangTidy
   clangTidyUtils

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h

+//===--- 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

clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

+//===--- 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

clang-tools-extra/docs/ReleaseNotes.rst

 ^^^^^^^^^^
 
 - New :doc:`concurrency-thread-canceltype-asynchronous
   <clang-tidy/checks/concurrency-thread-canceltype-asynchronous>` check.
 
   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.
 
   Finds inner loops that have not been unrolled, as well as fully unrolled
   loops with unknown loops bounds or a large number of iterations.
 
 - New :doc:`cppcoreguidelines-prefer-member-initializer
   <clang-tidy/checks/cppcoreguidelines-prefer-member-initializer>` check.

clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst

+.. 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>`_.

clang-tools-extra/docs/clang-tidy/checks/list.rst

.. csv-table::
    `abseil-no-namespace <abseil-no-namespace.html>`_,
    `abseil-redundant-strcat-calls <abseil-redundant-strcat-calls.html>`_, "Yes"
    `abseil-str-cat-append <abseil-str-cat-append.html>`_, "Yes"
    `abseil-string-find-startswith <abseil-string-find-startswith.html>`_, "Yes"
    `abseil-string-find-str-contains <abseil-string-find-str-contains.html>`_, "Yes"
    `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"
    `altera-unroll-loops <altera-unroll-loops.html>`_,
    `android-cloexec-accept <android-cloexec-accept.html>`_, "Yes"
    `android-cloexec-accept4 <android-cloexec-accept4.html>`_,
    `android-cloexec-creat <android-cloexec-creat.html>`_, "Yes"
    `android-cloexec-dup <android-cloexec-dup.html>`_, "Yes"

clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp

+// 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++;
+    }
+  }
+}