Index: clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp =================================================================== --- clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp +++ clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp @@ -11,6 +11,7 @@ #include "../ClangTidyModuleRegistry.h" #include "KernelNameRestrictionCheck.h" #include "StructPackAlignCheck.h" +#include "UnrollLoopsCheck.h" using namespace clang::ast_matchers; @@ -25,6 +26,8 @@ "altera-kernel-name-restriction"); CheckFactories.registerCheck( "altera-struct-pack-align"); + CheckFactories.registerCheck( + "altera-unroll-loops"); } }; Index: clang-tools-extra/clang-tidy/altera/CMakeLists.txt =================================================================== --- clang-tools-extra/clang-tidy/altera/CMakeLists.txt +++ clang-tools-extra/clang-tidy/altera/CMakeLists.txt @@ -7,6 +7,7 @@ AlteraTidyModule.cpp KernelNameRestrictionCheck.cpp StructPackAlignCheck.cpp + UnrollLoopsCheck.cpp LINK_LIBS clangTidy Index: clang-tools-extra/clang-tidy/altera/UnrollLoopsCheck.h =================================================================== --- /dev/null +++ clang-tools-extra/clang-tidy/altera/UnrollLoopsCheck.h @@ -0,0 +1,71 @@ +//===--- 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 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 Index: clang-tools-extra/clang-tidy/altera/UnrollLoopsCheck.cpp =================================================================== --- /dev/null +++ clang-tools-extra/clang-tidy/altera/UnrollLoopsCheck.cpp @@ -0,0 +1,165 @@ +//===--- 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("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 directive"); + return; + } + return; + } + diag(MatchedLoop->getBeginLoc(), + "full unrolling requested, but loop bounds are not known; to " + "partially unroll this loop, use the #pragma unroll directive"); + } +} + +enum UnrollLoopsCheck::UnrollType +UnrollLoopsCheck::unrollType(const Stmt *Statement, ASTContext *Context) { + const clang::DynTypedNodeList Parents = Context->getParents(*Statement); + for (const ast_type_traits::DynTypedNode &Parent : Parents) { + const auto *ParentStmt = Parent.get(); + 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(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(Conditional)) { + const auto *BinaryOp = static_cast(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(Statement)) { + const auto *ForLoop = static_cast(Statement); + Conditional = ForLoop->getCond(); + } + if (isa(Statement)) { + const auto *WhileLoop = static_cast(Statement); + Conditional = WhileLoop->getCond(); + } + if (isa(Statement)) { + const auto *DoWhileLoop = static_cast(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(Conditional)) { + const auto *BinaryOp = static_cast(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 Index: clang-tools-extra/docs/ReleaseNotes.rst =================================================================== --- clang-tools-extra/docs/ReleaseNotes.rst +++ clang-tools-extra/docs/ReleaseNotes.rst @@ -102,6 +102,12 @@ Finds structs that are inefficiently packed or aligned, and recommends packing and/or aligning of said structs as needed. +- New :doc:`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 ` check. Index: clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst =================================================================== --- /dev/null +++ clang-tools-extra/docs/clang-tidy/checks/altera-unroll-loops.rst @@ -0,0 +1,72 @@ +.. 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 +`_. + +.. 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 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. Index: clang-tools-extra/docs/clang-tidy/checks/list.rst =================================================================== --- clang-tools-extra/docs/clang-tidy/checks/list.rst +++ clang-tools-extra/docs/clang-tidy/checks/list.rst @@ -32,6 +32,7 @@ `abseil-upgrade-duration-conversions `_, "Yes" `altera-kernel-name-restriction `_, `altera-struct-pack-align `_, + `altera-unroll-loops `_, `android-cloexec-accept `_, "Yes" `android-cloexec-accept4 `_, `android-cloexec-creat `_, "Yes" Index: clang-tools-extra/test/clang-tidy/checkers/altera-unroll-loops.cpp =================================================================== --- /dev/null +++ clang-tools-extra/test/clang-tidy/checkers/altera-unroll-loops.cpp @@ -0,0 +1,359 @@ +// 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 directive [altera-unroll-loops] + printf("%d", someVector); + } while (j < vectorSize); +} +// There are no fix-its for this check