[clang-tidy] new altera unroll loops 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 unroll loops check finds inner loops that have not been unrolled, as well as fully-unrolled loops that should be partially unrolled due to unknown loop bounds or a large number of loop iterations.

Based on the Altera SDK for OpenCL: Best Practices Guide

[Eugene.Zelenko]

I think will be good idea to separate module code in own review or refer to previous one of previous reviews as dependency.

clang-tidy/altera/UnrollLoopsCheck.cpp

31 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence. See modernize-use-auto.

69 ↗

(On Diff #236271)

May be merge declaration with initialization and use const auto *?

96 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence.

113 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence.

117 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence.

121 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence.

134 ↗

(On Diff #236271)

Could be const auto *, because type is spelled in same sentence.

docs/ReleaseNotes.rst

79 ↗

(On Diff #236271)

Please synchronize with first statement in documentation.

docs/clang-tidy/checks/altera-unroll-loops.rst

68 ↗

(On Diff #236271)

Please document default value.

[ffrankies]

Addressed comments by @Eugene.Zelenko

• Used const auto * when type is mentioned in same sentence
• Fixed formatting in documentation
• Documented default value of MaxLoopIterations option

[Eugene.Zelenko]

Please rebase from master.

[ffrankies]

@Eugene.Zelenko It turns out we were using an old mirror of the clang-tools-extra repository, that no longer seems to be updated. I switched to the llvm-project repo, and updated the documentation accordingly (some of the formatting is different now). Let me know if I missed anything - if not, I will be switching the other checks over to the llvm-project repo as well over the coming weeks.

[Eugene.Zelenko]

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

41

Please elide braces.

66

Please elide braces.

93

Please elide braces.

130

Please elide braces.

138

Please elide braces.

141

Please elide braces.

152

Please elide braces.

156

Please elide braces.

clang-tools-extra/docs/ReleaseNotes.rst

85

Please synchronize with first statement in documentation.

[ffrankies]

Addressed comments by @Eugene.Zelenko

• Removed braces from one-lien if statements
• Release notes on the altera unroll loops check now match the first line of documentation.

[ffrankies]

• Rebased with master branch
• Diagnostics are re-worded so they're not complete sentences

Check ready for further review

[ffrankies]

• removed unnecessary comment

[njames93]

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

23

Dont use all caps variable names, preferred style is CamelCase.

25–27

36

This loops like it should be a switch

69

Shouldn't this by dyn_cast too.

70

To reduce indentations can this be an early exit

80

Default statements are usually bad as the prevent compiler diagnostics if the enum is ever updated.

94–95

108–121

129

dyn_cast again

146

Elide inner parens

149–151

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

32–34

Should probably define this out of line in the cpp file.

[ffrankies]

• Addressed comments from @njames93
• Rebased with latest master branch to make sure there are no merge issues with the latest committed altera check

[aaron.ballman]

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

41

44

52–53

59

86

This looks like dead code (it's still inside the switch statement).

112

Should we also handle CXXForRangeStmt? I'm thinking of cases like: for (int i : {1, 2, 3, 4}) {} or other cases where there's a known bounds.

146

If the check won't work too well if the value is decremented rather than incremented, or if it uses an induction variable that increases by something other than 1 or -1, we should probably document the limitations and consider whether we want to note that scenario in the code and bail out (rather than produce incorrect diagnostics, assuming we do).

clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst

5

clang-tools-extra/test/clang-tidy/checkers/altera-unroll-loops.cpp

360

Can you also add some test cases where the loop is decrementing rather than incrementing, and some tests where the increment is by more than 1?

[ffrankies]

• Added support for CXXForRangeStmt loops
• Added support for different for loop increments (++, --, +=, -=, *=, \=)
  • Depending on the exit condition, the calculations for the number of Iterations may be off by 1, but that should not be an issue performance-wise since the threshold for a "large" number of iterations is likely to be arbitrary.
  • If any other increment is used in the increment part of the for loop, we recommend partial unrolling.
• Changed the way while and do..while loops are handled: because the loop variable is changed within the loop body and it is untrivial to determine what is happening to it, we now emit a Note diagnostic (instead of warning) recommending partial unrolling if a while or do..while loop is fully unrolled.
• Added tests for the above
• Documented the above caveats in clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst

[Eugene.Zelenko]

clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst

16

Please use double back-ticks for language constructs. Same in rest of note text.

[ffrankies]

• Surrounded language constructs with double back ticks in clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst.
• Removed trailing whitespace.
• Added single quotes around #pragma unrolls in diagnostics.
• Added full stops to the comments that didn't have them in clang-tools-extra/clang-tidy/altera/UnrollLoopsCheck.cpp

[ffrankies]

• Removed unnecessary braces.
• Simplified boolean return statements.

[ffrankies]

Rebased with the latest main branch to fix patch application errors.

[ffrankies]

• Ran git-clang-format HEAD^
• Addressed comments automatically added by the Lint: Pre-merge Checks

[aaron.ballman]

Mostly just nits for the check, otherwise this LGTM.

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

66

84–85

Pretty sure you don't need to use a qualified name here.

117–121

124

177–178

The cast isn't necessary (IntergerLiteral is already an Expr). You should probably assert that CXXLoopBound isn't null here though.

181–182

197–198

[ffrankies]

Implemented changes requested by @aaron.ballman

[ffrankies]

@aaron.ballman If there are no more changes, can you please commit this on my behalf? My github username is ffrankies.

[aaron.ballman]

Thanks! I've commit on your behalf in 5a87f81fe9aee996dfe3a84dd833f0a48e093e7f

[ffrankies]

Thank you very much!

clang-tools-extra/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-tools-extra/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-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 "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<UnrollLoopsCheck>(
+        "altera-unroll-loops");
+  }
+};
+
+} // 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-tools-extra/clang-tidy/altera/CMakeLists.txt

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

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

+//===--- UnrollLoopsCheck.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_UNROLL_LOOPS_CHECK_H
+#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_UNROLL_LOOPS_CHECK_H
+
+#include "../ClangTidy.h"
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+/// Finds for, while and do..while loop statements that have not been unrolled.
+/// Unrolling these loops could increase the performance of the OpenCL kernel.
+/// Also attempts to find for, while and do..while loops that have been fully
+/// unrolled, but either have unknown bounds or a large number of iterations.
+/// The compiler will not unroll these loops, so partial unrolling is
+/// recommended in these cases.
+///
+/// For the user-facing documentation see:
+/// http://clang.llvm.org/extra/clang-tidy/checks/altera-unroll-loops.html
+class UnrollLoopsCheck : public ClangTidyCheck {
+  const unsigned MaxLoopIterations;
+
+public:
+  UnrollLoopsCheck(StringRef Name, ClangTidyContext *Context)
+      : ClangTidyCheck(Name, Context),
+        MaxLoopIterations(Options.get("MaxLoopIterations", 100U)) {}
+  void registerMatchers(ast_matchers::MatchFinder *Finder) override;
+  void check(const ast_matchers::MatchFinder::MatchResult &Result) override;
+
+private:
+  /// The kind of unrolling, if any, applied to a given loop.
+  enum UnrollType {
+    // This loop has no #pragma unroll directive associated with it.
+    NotUnrolled,
+    // This loop has a #pragma unroll directive associated with it.
+    FullyUnrolled,
+    // This loop has a #pragma unroll <num> directive associated with it.
+    PartiallyUnrolled
+  };
+  /// Returns true if the given loop statement has a large number of iterations,
+  /// as determined by the integer value in the loop's condition expression,
+  /// if one exists.
+  bool hasLargeNumIterations(const Stmt *Statement, const ASTContext *Context);
+  /// Checks one hand side of the binary operator to ascertain if the upper
+  /// bound on the number of loops is greater than max_loop_iterations or not.
+  /// If the expression is not evaluatable or not an integer, returns false.
+  bool exprHasLargeNumIterations(const Expr *Expression,
+                                 const ASTContext *Context);
+  /// Returns the type of unrolling, if any, associated with the given
+  /// statement.
+  enum UnrollType unrollType(const Stmt *Statement, ASTContext *Context);
+  /// Returns the condition expression within a given for statement. If there is
+  /// none, or if the Statement is not a loop, then returns a NULL pointer.
+  const Expr *getCondExpr(const Stmt *Statement);
+  /// Returns True if the loop statement has known bounds.
+  bool hasKnownBounds(const Stmt *Statement, const ASTContext *Context);
+  void storeOptions(ClangTidyOptions::OptionMap &Opts);
+};
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang
+
+#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_UNROLL_LOOPS_CHECK_H

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

+//===--- UnrollLoopsCheck.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 "UnrollLoopsCheck.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/AST/ASTTypeTraits.h"
+#include "clang/AST/ParentMapContext.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+void UnrollLoopsCheck::registerMatchers(MatchFinder *Finder) {
+  const auto ANYLOOP = anyOf(forStmt(), whileStmt(), doStmt());
+  Finder->addMatcher(stmt(allOf(ANYLOOP, // Match all loop types,
+                                unless(hasDescendant(forStmt())),
+                                unless(hasDescendant(whileStmt())),
+                                unless(hasDescendant(doStmt()))))
+                         .bind("loop"),
+                     this);
+}
+
+void UnrollLoopsCheck::check(const MatchFinder::MatchResult &Result) {
+  const auto *MatchedLoop = Result.Nodes.getNodeAs<Stmt>("loop");
+  const ASTContext *Context = Result.Context;
+  UnrollType Unroll = unrollType(MatchedLoop, Result.Context);
+  if (Unroll == NotUnrolled) {
+    diag(MatchedLoop->getBeginLoc(),
+         "The performance of the kernel could be improved by unrolling this "
+         "loop with a #pragma unroll directive");
+    return;
+  }
+  if (Unroll == PartiallyUnrolled) {
+    return;
+  }
+  if (Unroll == FullyUnrolled) {
+    if (hasKnownBounds(MatchedLoop, Context)) {
+      if (hasLargeNumIterations(MatchedLoop, Context)) {
+        diag(MatchedLoop->getBeginLoc(),
+             "This loop likely has a large number of iterations and thus "
+             "cannot be fully unrolled. To partially unroll this loop, use the "
+             "#pragma unroll <num> directive");
+        return;
+      }
+      return;
+    }
+    diag(MatchedLoop->getBeginLoc(),
+         "Full unrolling was requested, but loop bounds are not known. To "
+         "partially unroll this loop, use the #pragma unroll <num> directive");
+  }
+}
+
+enum UnrollLoopsCheck::UnrollType
+UnrollLoopsCheck::unrollType(const Stmt *Statement, ASTContext *Context) {
+  const clang::DynTypedNodeList Parents = Context->getParents<Stmt>(*Statement);
+  for (const ast_type_traits::DynTypedNode &Parent : Parents) {
+    const auto *ParentStmt = Parent.get<AttributedStmt>();
+    if (!ParentStmt) {
+      continue;
+    }
+    for (const Attr *Attribute : ParentStmt->getAttrs()) {
+      const auto *LoopHint = static_cast<const LoopHintAttr *>(Attribute);
+      if (LoopHint) {
+        switch (LoopHint->getState()) {
+        case LoopHintAttr::Numeric:
+          return PartiallyUnrolled;
+        case LoopHintAttr::Disable:
+          return NotUnrolled;
+        case LoopHintAttr::Full:
+          return FullyUnrolled;
+        case LoopHintAttr::Enable:
+          return FullyUnrolled;
+        default:
+          return NotUnrolled;
+        }
+      }
+    }
+  }
+  return NotUnrolled;
+}
+
+bool UnrollLoopsCheck::hasKnownBounds(const Stmt *Statement,
+                                      const ASTContext *Context) {
+  const Expr *Conditional = getCondExpr(Statement);
+  if (!Conditional) {
+    return false;
+  }
+  if (isa<BinaryOperator>(Conditional)) {
+    const auto *BinaryOp = static_cast<const BinaryOperator *>(Conditional);
+    const Expr *LHS = BinaryOp->getLHS();
+    const Expr *RHS = BinaryOp->getRHS();
+    if (LHS->isEvaluatable(*Context) == RHS->isEvaluatable(*Context)) {
+      return false; // Both sides are value dependent or constant, so loop
+                    // bounds are not known.
+    }
+    return true; // At least 1 side isn't value dependent, so we know the loop
+                 // bounds.
+  }
+  return false; // If it's not a binary operator, we don't know the loop bounds.
+}
+
+const Expr *UnrollLoopsCheck::getCondExpr(const Stmt *Statement) {
+  const Expr *Conditional;
+  if (isa<ForStmt>(Statement)) {
+    const auto *ForLoop = static_cast<const ForStmt *>(Statement);
+    Conditional = ForLoop->getCond();
+  }
+  if (isa<WhileStmt>(Statement)) {
+    const auto *WhileLoop = static_cast<const WhileStmt *>(Statement);
+    Conditional = WhileLoop->getCond();
+  }
+  if (isa<DoStmt>(Statement)) {
+    const auto *DoWhileLoop = static_cast<const DoStmt *>(Statement);
+    Conditional = DoWhileLoop->getCond();
+  }
+  return Conditional;
+}
+
+bool UnrollLoopsCheck::hasLargeNumIterations(const Stmt *Statement,
+                                             const ASTContext *Context) {
+  const Expr *Conditional = getCondExpr(Statement);
+  if (!Conditional) {
+    return false;
+  }
+  if (isa<BinaryOperator>(Conditional)) {
+    const auto *BinaryOp = static_cast<const BinaryOperator *>(Conditional);
+    const Expr *LHS = BinaryOp->getLHS();
+    const Expr *RHS = BinaryOp->getRHS();
+    Expr::EvalResult result;
+    if (LHS->isEvaluatable(*Context) && !(RHS->isEvaluatable(*Context))) {
+      return exprHasLargeNumIterations(LHS, Context);
+    }
+    if (RHS->isEvaluatable(*Context) && !(LHS->isEvaluatable(*Context))) {
+      return exprHasLargeNumIterations(RHS, Context);
+    }
+  }
+  return false; // Cannot check number of iteration, return false to be safe.
+}
+
+bool UnrollLoopsCheck::exprHasLargeNumIterations(const Expr *Expression,
+                                                 const ASTContext *Context) {
+  Expr::EvalResult Result;
+  if (Expression->EvaluateAsRValue(Result, *Context)) {
+    if (!(Result.Val.isInt())) {
+      return false; // Cannot check number of iterations, return false to be
+                    // safe.
+    }
+    if (Result.Val.getInt() > MaxLoopIterations) {
+      return true; // Assumes values go from 0 to Val in increments of 1.
+    }
+    return false; // Number of iterations likely less than MaxLoopIterations.
+  }
+  // Cannot evaluate Expression as an r-value, so cannot check number of
+  // iterations.
+  return false;
+}
+
+void UnrollLoopsCheck::storeOptions(ClangTidyOptions::OptionMap &Opts) {
+  Options.store(Opts, "MaxLoopIterations", MaxLoopIterations);
+}
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang

clang-tools-extra/docs/ReleaseNotes.rst

 Improvements to clang-rename
 ----------------------------
 
 The improvements are...
 
 Improvements to clang-tidy
 --------------------------
 
+New modules
+^^^^^^^^^^^
+
+- 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 checks
 ^^^^^^^^^^
 
+- 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:`bugprone-misplaced-pointer-arithmetic-in-alloc
   <clang-tidy/checks/bugprone-misplaced-pointer-arithmetic-in-alloc>` check.
 
   Finds cases where an integer expression is added to or subtracted from the
   result of a memory allocation function (``malloc()``, ``calloc()``,
   ``realloc()``, ``alloca()``) instead of its argument.
 
 - New :doc:`bugprone-reserved-identifier

clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst

+.. title:: clang-tidy - altera-unroll-loops
+
+altera-unroll-loops
+=================
+
+Checks for inner loops that have not been unrolled. Unrolling
+these loops could improve the performance of OpenCL kernels.
+
+Also checks for fully unrolled loops with a large number of 
+iterations, or unknown loop bounds. These loops cannot be fully
+unrolled, and should be partially unrolled.
+
+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++
+
+   for (int i = 0; i < 10; i++) {  // ok: outer loops should not be unrolled
+      int j = 0;
+      do {  // warning: this inner do..while loop should be unrolled
+         j++;
+      } while (j < 15);
+
+      int k = 0;
+      #pragma unroll
+      while (k < 20) {  // ok: this inner loop is already unrolled 
+         k++;
+      }
+   }
+
+   #pragma unroll
+   for (int i = 0; i < 1000; ++i) {  // warning: this loop is too large and cannot be fully unrolled
+      printf("%d", i);
+   }
+
+   #pragma unroll 5
+   for (int i = 0; i < 1000; ++i) {  // ok: this loop is large, but is partially unrolled
+      printf("%d", i);
+   }
+
+   std::vector<int> someVector (100, 0);
+   int i = 0;
+   #pragma unroll
+   while (i < someVector.size()) {  // warning: this loop has unknown bounds and cannot be fully unrolled
+      someVector[i]++;
+   }
+
+   #pragma unroll
+   while (true) {  // warning: this loop has unknown bounds and cannot be fully unrolled
+      printf("In loop");
+   }
+
+   #pragma unroll
+   while (i < 5) {  // warning: this loop has unknown bounds and cannot be fully unrolled
+      printf("In loop");
+   }
+
+   #pragma unroll 5
+   while (i < someVector.size()) {  // ok: this loop has unknown bounds, but is partially unrolled
+      someVector[i]++;
+   }
+
+Options
+-------
+
+.. option:: MaxLoopIterations
+
+   Defines the maximum number of loop iterations that a fully unrolled loop
+   can have. By default, it is set to `100`.
+
+   In practice, this refers to the integer value of the upper bound
+   within the loop statement's condition expression.

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

.. csv-table::
    `abseil-no-internal-dependencies <abseil-no-internal-dependencies.html>`_,
    `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-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-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"
    `android-cloexec-epoll-create <android-cloexec-epoll-create.html>`_,
    `android-cloexec-epoll-create1 <android-cloexec-epoll-create1.html>`_,
    `android-cloexec-fopen <android-cloexec-fopen.html>`_,
    `android-cloexec-inotify-init <android-cloexec-inotify-init.html>`_,

clang-tools-extra/docs/clang-tidy/index.rst

 .. _checks-groups-table:
 
 There are currently the following groups of checks:
 
 ====================== =========================================================
 Name prefix            Description
 ====================== =========================================================
 ``abseil-``            Checks related to Abseil library.
+``altera-``            Checks related to OpenCL programming for FPGAs.
 ``android-``           Checks related to Android.
 ``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.

clang-tools-extra/test/clang-tidy/checkers/altera-unroll-loops.cpp

+// RUN: %check_clang_tidy %s altera-unroll-loops %t -- -config="{CheckOptions: [{key: "altera-unroll-loops.MaxLoopIterations", value: 50}]}" -header-filter=.* "--" --include opencl-c.h -cl-std=CL1.2 -c
+
+// Inner loops should be unrolled
+__kernel void nested_simple_loops(__global int *A) {
+    for (int i = 0; i < 1000; ++i) {
+        for (int j = 0; j < 2000; ++j) { 
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[0] += i + j;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0;
+        while (j < 2000) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[1] += i + j;
+            j++;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0; 
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[2] += i + j;
+            j++;
+        } while (j < 2000);
+    }
+
+    int i = 0;
+    while (i < 1000) {
+        for (int j = 0; j < 2000; ++j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[3] += i + j;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (i < 1000) {
+        int j = 0;
+        while (j < 2000) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[4] += i + j;
+            j++;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (i < 1000) {
+        int j = 0;
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[5] += i + j;
+            j++;
+        } while (j < 2000);
+        i++;
+    }
+
+    i = 0;
+    do {
+        for (int j = 0; j < 2000; ++j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[6] += i + j;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        while (j < 2000) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[7] += i + j;
+            j++;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+            A[8] += i + j;
+            j++;
+        } while (j < 2000);
+        i++;
+    } while (i < 1000);
+
+    for (int i = 0; i < 100; ++i) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+        printf("Hello");
+    }
+
+    i = 0;
+    while (i < 100) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+        i++;
+    }
+
+    i = 0;
+    do {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: The performance of the kernel could be improved by unrolling this loop with a #pragma unroll directive [altera-unroll-loops]
+        i++;
+    } while (i < 100);
+}
+
+// These loops are all correctly unrolled
+__kernel void unrolled_nested_simple_loops(__global int *A) {
+    for (int i = 0; i < 1000; ++i) {
+        #pragma unroll
+        for (int j = 0; j < 50; ++j) {
+            A[0] += i + j;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0;
+        #pragma unroll
+        while (j < 50) {
+            A[1] += i + j;
+            j++;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0; 
+        #pragma unroll
+        do {
+            A[2] += i + j;
+            j++;
+        } while (j < 50);
+    }
+
+    int i = 0;
+    while (i < 1000) {
+        #pragma unroll
+        for (int j = 0; j < 50; ++j) {
+            A[3] += i + j;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (i < 1000) {
+        int j = 0;
+        #pragma unroll
+        while (50 > j) {
+            A[4] += i + j;
+            j++;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (1000 > i) {
+        int j = 0;
+        #pragma unroll
+        do {
+            A[5] += i + j;
+            j++;
+        } while (j < 50);
+        i++;
+    }
+
+    i = 0;
+    do {
+        #pragma unroll
+        for (int j = 0; j < 50; ++j) {
+            A[6] += i + j;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        #pragma unroll
+        while (j < 50) {
+            A[7] += i + j;
+            j++;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        #pragma unroll
+        do {
+            A[8] += i + j;
+            j++;
+        } while (j < 50);
+        i++;
+    } while (i < 1000);
+}
+
+__kernel void unrolled_nested_simple_loops_large_num_iterations(__global int *A) {
+    for (int i = 0; i < 1000; ++i) {
+        #pragma unroll
+        for (int j = 0; j < 51; ++j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[0] += i + j;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0;
+        #pragma unroll
+        while (j < 51) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[1] += i + j;
+            j++;
+        }
+    }
+
+    for (int i = 0; i < 1000; ++i) {
+        int j = 0; 
+        #pragma unroll
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[2] += i + j;
+            j++;
+        } while (j < 51);
+    }
+
+    int i = 0;
+    while (i < 1000) {
+        #pragma unroll
+        for (int j = 0; j < 51; ++j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[3] += i + j;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (i < 1000) {
+        int j = 0;
+        #pragma unroll
+        while (51 > j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[4] += i + j;
+            j++;
+        }
+        i++;
+    }
+
+    i = 0;
+    while (1000 > i) {
+        int j = 0;
+        #pragma unroll
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[5] += i + j;
+            j++;
+        } while (j < 51);
+        i++;
+    }
+
+    i = 0;
+    do {
+        #pragma unroll
+        for (int j = 0; j < 51; ++j) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[6] += i + j;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        #pragma unroll
+        while (j < 51) {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[7] += i + j;
+            j++;
+        }
+        i++;
+    } while (i < 1000);
+
+    i = 0;
+    do {
+        int j = 0;
+        #pragma unroll
+        do {
+// CHECK-MESSAGES: :[[@LINE-1]]:9: warning: This loop likely has a large number of iterations and thus cannot be fully unrolled. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+            A[8] += i + j;
+            j++;
+        } while (j < 51);
+        i++;
+    } while (i < 1000);
+}
+
+__kernel void fully_unrolled_unknown_bounds(int vectorSize) {
+    int someVector[101];
+
+    // There is no loop condition
+    #pragma unroll
+    for (;;) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    }
+
+    // Infinite loop
+    #pragma unroll
+    while (true) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    }
+
+    #pragma unroll
+    do {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    } while (true);
+
+    #pragma unroll
+    while (1 < 5) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    }
+
+    #pragma unroll
+    do {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    } while (1 < 5);
+
+    #pragma unroll
+    for (int i = 0; 1 < 5; ++i) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        someVector[i]++;
+    }
+
+    // Both sides are value-dependent
+    #pragma unroll
+    for (int i = 0; i < vectorSize; ++i) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    }
+
+    int j = 0;
+    #pragma unroll
+    while (j < vectorSize) {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    }
+    
+    #pragma unroll
+    do {
+// CHECK-MESSAGES: :[[@LINE-1]]:5: warning: Full unrolling was requested, but loop bounds are not known. To partially unroll this loop, use the #pragma unroll <num> directive [altera-unroll-loops]
+        printf("%d", someVector);
+    } while (j < vectorSize);
+}
+// There are no fix-its for this check