Index: clang-tidy/CMakeLists.txt =================================================================== --- clang-tidy/CMakeLists.txt +++ clang-tidy/CMakeLists.txt @@ -41,6 +41,7 @@ # 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) @@ -65,6 +66,7 @@ set(ALL_CLANG_TIDY_CHECKS clangTidyAndroidModule clangTidyAbseilModule + clangTidyAlteraModule clangTidyBoostModule clangTidyBugproneModule clangTidyCERTModule Index: clang-tidy/ClangTidyForceLinker.h =================================================================== --- clang-tidy/ClangTidyForceLinker.h +++ clang-tidy/ClangTidyForceLinker.h @@ -25,6 +25,11 @@ 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 = Index: clang-tidy/altera/AlteraTidyModule.cpp =================================================================== --- /dev/null +++ clang-tidy/altera/AlteraTidyModule.cpp @@ -0,0 +1,39 @@ +//===--- 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( + "altera-unroll-loops"); + } +}; + +} // namespace altera + +// Register the AlteraTidyModule using this statically initialized variable. +static ClangTidyModuleRegistry::Add + 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 Index: clang-tidy/altera/CMakeLists.txt =================================================================== --- /dev/null +++ clang-tidy/altera/CMakeLists.txt @@ -0,0 +1,15 @@ +set(LLVM_LINK_COMPONENTS support) + +add_clang_library(clangTidyAlteraModule + AlteraTidyModule.cpp + UnrollLoopsCheck.cpp + + LINK_LIBS + clangAnalysis + clangAST + clangASTMatchers + clangBasic + clangLex + clangTidy + clangTidyUtils + ) Index: clang-tidy/altera/UnrollLoopsCheck.h =================================================================== --- /dev/null +++ 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 "../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 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-tidy/altera/UnrollLoopsCheck.cpp =================================================================== --- /dev/null +++ clang-tidy/altera/UnrollLoopsCheck.cpp @@ -0,0 +1,176 @@ +//===--- 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/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 Stmt *MatchedLoop = Result.Nodes.getNodeAs("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 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 directive"); + } +} + +enum UnrollLoopsCheck::UnrollType +UnrollLoopsCheck::unrollType(const Stmt *Statement, ASTContext *Context) { + const ASTContext::DynTypedNodeList Parents = Context->getParents(*Statement); + for (const ast_type_traits::DynTypedNode &Parent: Parents) { + const AttributedStmt *ParentStmt = Parent.get(); + if (!ParentStmt) { + continue; + } + for (const Attr* Attribute: ParentStmt->getAttrs()) { + const LoopHintAttr *LoopHint; + if ((LoopHint = static_cast(Attribute))) { + 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 (std::string(Conditional->getStmtClassName()).compare("BinaryOperator") == + 0) { + const BinaryOperator *BinaryOp = + static_cast(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(Statement)) { + const ForStmt *ForLoop = static_cast(Statement); + Conditional = ForLoop->getCond(); + } + if (isa(Statement)) { + const WhileStmt *WhileLoop = static_cast(Statement); + Conditional = WhileLoop->getCond(); + } + if (isa(Statement)) { + const DoStmt *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)) { + // if (std::string(Conditional->getStmtClassName()).compare("BinaryOperator") == + // 0) { + const BinaryOperator *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: docs/ReleaseNotes.rst =================================================================== --- docs/ReleaseNotes.rst +++ docs/ReleaseNotes.rst @@ -67,6 +67,18 @@ Improvements to clang-tidy -------------------------- +- New :doc:`altera ` module. + + Includes checks related to OpenCL for FPGA coding guidelines, based on the + `Altera SDK for OpenCL: Best Practices Guide + `_. + +- 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:`bugprone-dynamic-static-initializers ` check. Index: docs/clang-tidy/checks/altera-unroll-loops.rst =================================================================== --- /dev/null +++ docs/clang-tidy/checks/altera-unroll-loops.rst @@ -0,0 +1,72 @@ +.. 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 +`_. + +.. 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. + + In practice, this refers to the integer value of the upper bound + within the loop statement's condition expression. Index: docs/clang-tidy/checks/list.rst =================================================================== --- docs/clang-tidy/checks/list.rst +++ docs/clang-tidy/checks/list.rst @@ -22,6 +22,7 @@ abseil-time-comparison abseil-time-subtraction abseil-upgrade-duration-conversions + altera-unroll-loops android-cloexec-accept android-cloexec-accept4 android-cloexec-creat Index: docs/clang-tidy/index.rst =================================================================== --- docs/clang-tidy/index.rst +++ docs/clang-tidy/index.rst @@ -59,6 +59,7 @@ ====================== ========================================================= ``abseil-`` Checks related to Abseil library. ``android-`` Checks related to Android. +``altera-`` Checks related to OpenCL programming for FPGAs. ``boost-`` Checks related to Boost library. ``bugprone-`` Checks that target bugprone code constructs. ``cert-`` Checks related to CERT Secure Coding Guidelines. Index: test/clang-tidy/checkers/altera-unroll-loops.cpp =================================================================== --- /dev/null +++ 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: 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 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 directive [altera-unroll-loops] + printf("%d", someVector); + } while (j < vectorSize); +} +// There are no fix-its for this check