[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

42

Please elide braces.

67

Please elide braces.

94

Please elide braces.

131

Please elide braces.

139

Please elide braces.

142

Please elide braces.

153

Please elide braces.

157

Please elide braces.

clang-tools-extra/docs/ReleaseNotes.rst

96

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/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 "KernelNameRestrictionCheck.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<KernelNameRestrictionCheck>(
         "altera-kernel-name-restriction");
     CheckFactories.registerCheck<StructPackAlignCheck>(
         "altera-struct-pack-align");
+    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
   FrontendOpenMP
   support
   )
 
 add_clang_library(clangTidyAlteraModule
   AlteraTidyModule.cpp
   KernelNameRestrictionCheck.cpp
   StructPackAlignCheck.cpp
+  UnrollLoopsCheck.cpp
 
   LINK_LIBS
   clangTidy
   clangTidyUtils
   )
 
 clang_target_link_libraries(clangTidyAlteraModule
   PRIVATE
   clangAnalysis
   clangAST
   clangASTMatchers
   clangBasic
   clangLex
   )

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 "../ClangTidyCheck.h"
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+/// Finds inner loops that have not been unrolled, as well as fully unrolled
+/// loops with unknown loop bounds or a large number of iterations.
+/// 
+/// Unrolling inner loops could improve the performance of OpenCL kernels.
+/// However, if they have unknown loop bounds or a large number of iterations,
+/// they cannot be fully unrolled, and should be partially unrolled.
+///
+/// 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(),
+         "kernel performance 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(),
+             "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 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;
+    // Instead of looping over attrs, get the LoopHintAttr directly like in
+    // single-work-item-barrier
+    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))
+      // Both sides are value dependent or constant, so loop bounds are unknown.
+      return false;
+    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

- New ``concurrency`` module.
   coroutines, etc.).
 
 New checks
 ^^^^^^^^^^
 
 - New :doc:`altera-kernel-name-restriction
   <clang-tidy/checks/altera-kernel-name-restriction>` check.
 
   Finds kernel files and include directives whose filename is `kernel.cl`,
   `Verilog.cl`, or `VHDL.cl`.
 
 - New :doc:`altera-struct-pack-align
   <clang-tidy/checks/altera-struct-pack-align>` check.
 
   Finds structs that are inefficiently packed or aligned, and recommends
   packing and/or aligning of said structs as needed.
 
+- 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.
 
   Finds member initializations in the constructor body which can be placed into
   the initialization list instead.
 
 - New :doc:`bugprone-misplaced-pointer-arithmetic-in-alloc
   <clang-tidy/checks/bugprone-misplaced-pointer-arithmetic-in-alloc>` check.

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

+.. title:: clang-tidy - altera-unroll-loops
+
+altera-unroll-loops
+=================
+
+Finds inner loops that have not been unrolled, as well as fully unrolled loops
+with unknown loop bounds or a large number of iterations.
+
+Unrolling inner loops could improve the performance of OpenCL kernels. However,
+if they have unknown loop bounds or a large number of iterations, they 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-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-kernel-name-restriction <altera-kernel-name-restriction.html>`_,
    `altera-struct-pack-align <altera-struct-pack-align.html>`_,
+   `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/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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: kernel performance 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: 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: 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: 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: 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: 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: 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: 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: 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: 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 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 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 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 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 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 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 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 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 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