[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

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

129

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

116

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

117

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

125

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

126

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

166

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

9 ↗

(On Diff #334979)

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

200 ↗

(On Diff #335195)

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

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

1 ↗

(On Diff #335195)

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

36 ↗

(On Diff #335195)

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

1 ↗

(On Diff #336512)

*- 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

200 ↗

(On Diff #335195)

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

28–31 ↗

(On Diff #338160)

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

54 ↗

(On Diff #338160)

101–102 ↗

(On Diff #338160)

121–122 ↗

(On Diff #338160)

132–134 ↗

(On Diff #338160)

You should use an llvm::Twine for this.

141–143 ↗

(On Diff #338160)

Same here as above.

154 ↗

(On Diff #338160)

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

199–204 ↗

(On Diff #338160)

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

26 ↗

(On Diff #338160)

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

32 ↗

(On Diff #338160)

35 ↗

(On Diff #338160)

37–38 ↗

(On Diff #338160)

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

32 ↗

(On Diff #338160)

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.

35 ↗

(On Diff #338160)

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-tidy/CMakeLists.txt

target_link_libraries(clangTidy PRIVATE
     clangStaticAnalyzerFrontend
   )
 endif()
 
 # Checks.
 # If you add a check, also add it to ClangTidyForceLinker.h in this directory.
 add_subdirectory(android)
 add_subdirectory(abseil)
+add_subdirectory(altera)
 add_subdirectory(boost)
 add_subdirectory(bugprone)
 add_subdirectory(cert)
 add_subdirectory(cppcoreguidelines)
 add_subdirectory(darwin)
 add_subdirectory(fuchsia)
 add_subdirectory(google)
 add_subdirectory(hicpp)
 add_subdirectory(linuxkernel)
 add_subdirectory(llvm)
 add_subdirectory(misc)
 add_subdirectory(modernize)
 if(CLANG_ENABLE_STATIC_ANALYZER)
   add_subdirectory(mpi)
 endif()
 add_subdirectory(objc)
 add_subdirectory(openmp)
 add_subdirectory(performance)
 add_subdirectory(portability)
 add_subdirectory(readability)
 add_subdirectory(zircon)
 set(ALL_CLANG_TIDY_CHECKS
   clangTidyAndroidModule
   clangTidyAbseilModule
+  clangTidyAlteraModule
   clangTidyBoostModule
   clangTidyBugproneModule
   clangTidyCERTModule
   clangTidyCppCoreGuidelinesModule
   clangTidyDarwinModule
   clangTidyFuchsiaModule
   clangTidyGoogleModule
   clangTidyHICPPModule

clang-tidy/ClangTidyForceLinker.h

 static int LLVM_ATTRIBUTE_UNUSED CERTModuleAnchorDestination =
     CERTModuleAnchorSource;
 
 // This anchor is used to force the linker to link the AbseilModule.
 extern volatile int AbseilModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED AbseilModuleAnchorDestination =
     AbseilModuleAnchorSource;
 
+// This anchor is used to force the linker to link the AlteraModule.
+extern volatile int AlteraModuleAnchorSource;
+static int LLVM_ATTRIBUTE_UNUSED AlteraModuleAnchorDestination =
+    AlteraModuleAnchorSource;
+
 // This anchor is used to force the linker to link the BoostModule.
 extern volatile int BoostModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED BoostModuleAnchorDestination =
     BoostModuleAnchorSource;
 
 // This anchor is used to force the linker to link the BugproneModule.
 extern volatile int BugproneModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED BugproneModuleAnchorDestination =

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"
+
+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");
+  }
+};
+
+} // namespace altera
+
+// Register the AlteraTidyModule using this statically initialized variable.
+static ClangTidyModuleRegistry::Add<altera::AlteraModule>
+    X("altera-module", "Adds Altera FPGA OpenCL lint checks.");
+
+// This anchor is used to force the linker to link in the generated object file
+// and thus register the AlteraModule.
+volatile int AlteraModuleAnchorSource = 0;
+
+} // namespace tidy
+} // namespace clang

clang-tidy/altera/CMakeLists.txt

+set(LLVM_LINK_COMPONENTS support)
+
+add_clang_library(clangTidyAlteraModule
+  AlteraTidyModule.cpp
+  IdDependentBackwardBranchCheck.cpp
+
+  LINK_LIBS
+  clangAnalysis
+  clangAST
+  clangASTMatchers
+  clangBasic
+  clangLex
+  clangTidy
+  clangTidyUtils
+  )

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_ID_DEPENDENT_BACKWARD_BRANCH_H
+#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_ID_DEPENDENT_BACKWARD_BRANCH_H
+
+#include "../ClangTidy.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 { UNK_LOOP = -1, DO_LOOP = 0, WHILE_LOOP = 1, FOR_LOOP = 2 };
+  // Stores information necessary for printing out source of error.
+  struct IdDependencyRecord {
+    IdDependencyRecord(const VarDecl *Declaration, SourceLocation Location,
+                       std::string Message)
+        : VariableDeclaration(Declaration), Location(Location),
+          Message(Message) {}
+    IdDependencyRecord(const FieldDecl *Declaration, SourceLocation Location,
+                       std::string Message)
+        : FieldDeclaration(Declaration), Location(Location), Message(Message) {}
+    IdDependencyRecord() {}
+    const VarDecl *VariableDeclaration;
+    const FieldDecl *FieldDeclaration;
+    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_ID_DEPENDENT_BACKWARD_BRANCH_H

clang-tidy/altera/IdDependentBackwardBranchCheck.cpp

+//===--- IdDependentBackwardBranchCheck.cpp - clang-tidy ------------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+
+#include "IdDependentBackwardBranchCheck.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+#include <sstream>
+#include <vector>
+
+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 TID_RHS = expr(hasDescendant(callExpr(callee(functionDecl(
+      anyOf(hasName("get_global_id"), hasName("get_local_id")))))));
+
+  const auto ANY_ASSIGN = anyOf(
+      hasOperatorName("="), hasOperatorName("*="), hasOperatorName("/="),
+      hasOperatorName("%="), hasOperatorName("+="), hasOperatorName("-="),
+      hasOperatorName("<<="), hasOperatorName(">>="), hasOperatorName("&="),
+      hasOperatorName("^="), hasOperatorName("|="));
+
+  Finder->addMatcher(
+      compoundStmt(
+          // Bind on actual get_local/global_id calls
+          forEachDescendant(
+              stmt(anyOf(declStmt(hasDescendant(varDecl(hasInitializer(TID_RHS))
+                                                    .bind("tid_dep_var"))),
+                         binaryOperator(allOf(
+                             ANY_ASSIGN, hasRHS(TID_RHS),
+                             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
+  // (incase it is ID-dependent)
+  Finder->addMatcher(
+      stmt(forEachDescendant(
+          varDecl(
+              hasInitializer(forEachDescendant(stmt(anyOf(
+                  declRefExpr(to(varDecl())).bind("assign_ref_var"),
+                  memberExpr(member(fieldDecl())).bind("assign_ref_field"))))))
+              .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(
+          ANY_ASSIGN,
+          hasRHS(forEachDescendant(stmt(anyOf(
+              declRefExpr(to(varDecl())).bind("assign_ref_var"),
+              memberExpr(member(fieldDecl())).bind("assign_ref_field"))))),
+          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 COND_EXPR =
+      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(COND_EXPR)),
+                                doStmt(hasCondition(COND_EXPR)),
+                                whileStmt(hasCondition(COND_EXPR))))
+                         .bind("backward_branch"),
+                     this);
+}
+
+IdDependentBackwardBranchCheck::IdDependencyRecord *
+IdDependentBackwardBranchCheck::hasIdDepVar(const Expr *Expression) {
+  if (const DeclRefExpr *expr = dyn_cast<DeclRefExpr>(Expression)) {
+    // It is a DeclRefExpr, so check if it's an ID-dependent variable
+    const VarDecl *CheckVariable = dyn_cast<VarDecl>(expr->getDecl());
+    auto FoundVariable = IdDepVarsMap.find(CheckVariable);
+    if (FoundVariable == IdDepVarsMap.end()) {
+      return nullptr;
+    }
+    return &(FoundVariable->second);
+  }
+  for (auto i = Expression->child_begin(), e = Expression->child_end(); i != e;
+       ++i) {
+    if (auto *ChildExpression = dyn_cast<Expr>(*i)) {
+      auto Result = hasIdDepVar(ChildExpression);
+      if (Result) {
+        return Result;
+      }
+    }
+  }
+  return nullptr;
+}
+
+IdDependentBackwardBranchCheck::IdDependencyRecord *
+IdDependentBackwardBranchCheck::hasIdDepField(const Expr *Expression) {
+  if (const MemberExpr *MemberExpression = dyn_cast<MemberExpr>(Expression)) {
+    const FieldDecl *CheckField =
+        dyn_cast<FieldDecl>(MemberExpression->getMemberDecl());
+    auto FoundField = IdDepFieldsMap.find(CheckField);
+    if (FoundField == IdDepFieldsMap.end()) {
+      return nullptr;
+    }
+    return &(FoundField->second);
+  }
+  for (auto I = Expression->child_begin(), E = Expression->child_end(); I != E;
+       ++I) {
+    if (auto *ChildExpression = dyn_cast<Expr>(*I)) {
+      auto Result = hasIdDepField(ChildExpression);
+      if (Result) {
+        return Result;
+      }
+    }
+  }
+  return nullptr;
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepVar(const Stmt *Statement,
+                                                  const VarDecl *Variable) {
+  // Record that this variable is thread-dependent
+  std::ostringstream StringStream;
+  StringStream << "assignment of ID-dependent variable "
+               << Variable->getNameAsString();
+  IdDepVarsMap[Variable] =
+      IdDependencyRecord(Variable, Variable->getBeginLoc(), StringStream.str());
+}
+
+void IdDependentBackwardBranchCheck::saveIdDepField(const Stmt *Statement,
+                                                    const FieldDecl *Field) {
+  std::ostringstream StringStream;
+  StringStream << "assignment of ID-dependent field "
+               << Field->getNameAsString();
+  IdDepFieldsMap[Field] =
+      IdDependencyRecord(Field, Statement->getBeginLoc(), StringStream.str());
+}
+
+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::ostringstream StringStream;
+  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(), StringStream.str());
+}
+
+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::ostringstream StringStream;
+  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(), StringStream.str());
+}
+
+IdDependentBackwardBranchCheck::LoopType
+IdDependentBackwardBranchCheck::getLoopType(const Stmt *Loop) {
+  if (const auto DoLoop = dyn_cast<DoStmt>(Loop)) {
+    return DO_LOOP; // loop_type = 0;
+  } else if (const auto WhileLoop = dyn_cast<WhileStmt>(Loop)) {
+    return WHILE_LOOP; // loop_type = 1;
+  } else if (const auto ForLoop = dyn_cast<ForStmt>(Loop)) {
+    return FOR_LOOP; // loop_type = 2;
+  }
+  return UNK_LOOP;
+}
+
+void IdDependentBackwardBranchCheck::check(
+    const MatchFinder::MatchResult &Result) {
+  // The first half of the callback only deals with identifying and propagating
+  // ID-dependency information into the IdDepVars vector
+  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) { // If 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;
+    } else { // Conditional expression has DeclRefExpr(s), check ID-dependency
+      IdDependencyRecord *IdDepVar = hasIdDepVar(CondExpr);
+      IdDependencyRecord *IdDepField = hasIdDepField(CondExpr);
+      if (IdDepVar) {
+        diag(IdDepVar->Location, IdDepVar->Message);
+        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);
+        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

docs/ReleaseNotes.rst

 Improvements to clang-rename
 ----------------------------
 
 The improvements are...
 
 Improvements to clang-tidy
 --------------------------
 
+- New :doc:`altera <clang-tidy/modules/altera>` module.
+
+  Includes checks related to OpenCL for FPGA coding guidelines, 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>`_.
+
+- 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.  
+
 - New :doc:`bugprone-dynamic-static-initializers
   <clang-tidy/checks/bugprone-dynamic-static-initializers>` check.
 
   Finds instances where variables with static storage are initialized
   dynamically in header files.
 
 - New :doc:`bugprone-infinite-loop
   <clang-tidy/checks/bugprone-infinite-loop>` check.

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

docs/clang-tidy/checks/list.rst

.. toctree::
    abseil-no-internal-dependencies
    abseil-no-namespace
    abseil-redundant-strcat-calls
    abseil-str-cat-append
    abseil-string-find-startswith
    abseil-time-comparison
    abseil-time-subtraction
    abseil-upgrade-duration-conversions
+   altera-id-dependent-backward-branch
    android-cloexec-accept
    android-cloexec-accept4
    android-cloexec-creat
    android-cloexec-dup
    android-cloexec-epoll-create
    android-cloexec-epoll-create1
    android-cloexec-fopen
    android-cloexec-inotify-init

docs/clang-tidy/index.rst

 
 There are currently the following groups of checks:
 
 ====================== =========================================================
 Name prefix            Description
 ====================== =========================================================
 ``abseil-``            Checks related to Abseil library.
 ``android-``           Checks related to Android.
+``altera-``            Checks related to OpenCL programming for FPGAs.
 ``boost-``             Checks related to Boost library.
 ``bugprone-``          Checks that target bugprone code constructs.
 ``cert-``              Checks related to CERT Secure Coding Guidelines.
 ``cppcoreguidelines-`` Checks related to C++ Core Guidelines.
 ``clang-analyzer-``    Clang Static Analyzer checks.
 ``fuchsia-``           Checks related to Fuchsia coding conventions.
 ``google-``            Checks related to Google coding conventions.
 ``hicpp-``             Checks related to High Integrity C++ Coding Standard.

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() {
+  // ==== Assignments ====
+  int ThreadID = get_local_id(0); 
+// CHECK-NOTES: :[[@LINE-1]]:3: warning: assignment of ID-dependent variable ThreadID [altera-id-dependent-backward-branch]
+
+  ExampleStruct Example;
+  Example.IDDepField = get_local_id(0);
+// CHECK-NOTES: :[[@LINE-1]]:3: warning: assignment of ID-dependent field IDDepField [altera-id-dependent-backward-branch]
+
+  // ==== Inferred Assignments ====
+  int ThreadID2 = ThreadID * get_local_size(0);
+// CHECK-NOTES: :[[@LINE-1]]:3: warning: inferred assignment of ID-dependent value from ID-dependent variable ThreadID [altera-id-dependent-backward-branch]
+  
+  int ThreadID3 = Example.IDDepField;  // OK: not used in any loops
+
+  ExampleStruct UnusedStruct = {
+    ThreadID * 2  // OK: not used in any loops
+  };
+
+  // ==== 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]
+  
+  for (int i = 0; i < ThreadID2; i++) {
+    // CHECK-NOTES: :[[@LINE-1]]: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++;
+  }
+
+  while (j < ThreadID) {
+    // CHECK-NOTES: :[[@LINE-1]]: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++;
+  }
+
+  do {
+    accumulator++;
+  } while (j < ThreadID);
+  // CHECK-NOTES: :[[@LINE-1]]: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-1]]: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-1]]: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-1]]: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++;
+    }
+  }
+}