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 "IdDependentBackwardBranchCheck.h" + +using namespace clang::ast_matchers; + +namespace clang { +namespace tidy { +namespace altera { + +class AlteraModule : public ClangTidyModule { +public: + void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override { + CheckFactories.registerCheck( + "altera-id-dependent-backward-branch"); + } +}; + +} // 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 + IdDependentBackwardBranchCheck.cpp + + LINK_LIBS + clangAnalysis + clangAST + clangASTMatchers + clangBasic + clangLex + clangTidy + clangTidyUtils + ) Index: clang-tidy/altera/IdDependentBackwardBranchCheck.h =================================================================== --- /dev/null +++ clang-tidy/altera/IdDependentBackwardBranchCheck.h @@ -0,0 +1,82 @@ +//===--- IdDependentBackwardBranchCheck.h - clang-tidy-----------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_ID_DEPENDENT_BACKWARD_BRANCH_H +#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_ID_DEPENDENT_BACKWARD_BRANCH_H + +#include "../ClangTidy.h" + +namespace clang { +namespace tidy { +namespace altera { + +/// Finds ID-dependent variables and fields used within loops, and warns of +/// their usage. Using these variables in loops can lead to performance +/// degradation. +/// +/// For the user-facing documentation see: +/// http://clang.llvm.org/extra/clang-tidy/checks/altera-id-dependent-backward-branch.html +class IdDependentBackwardBranchCheck : public ClangTidyCheck { +private: + enum LoopType { UNK_LOOP = -1, DO_LOOP = 0, WHILE_LOOP = 1, FOR_LOOP = 2 }; + // Stores information necessary for printing out source of error. + struct IdDependencyRecord { + IdDependencyRecord(const VarDecl *Declaration, SourceLocation Location, + std::string Message) + : VariableDeclaration(Declaration), Location(Location), + Message(Message) {} + IdDependencyRecord(const FieldDecl *Declaration, SourceLocation Location, + std::string Message) + : FieldDeclaration(Declaration), Location(Location), Message(Message) {} + IdDependencyRecord() {} + const VarDecl *VariableDeclaration; + const FieldDecl *FieldDeclaration; + SourceLocation Location; + std::string Message; + }; + // Stores the locations where ID-dependent variables are created. + std::map IdDepVarsMap; + // Stores the locations where ID-dependent fields are created. + std::map IdDepFieldsMap; + /// Returns an IdDependencyRecord if the Expression contains an ID-dependent + /// variable, returns a nullptr otherwise. + IdDependencyRecord *hasIdDepVar(const Expr *Expression); + /// Returns an IdDependencyRecord if the Expression contains an ID-dependent + /// field, returns a nullptr otherwise. + IdDependencyRecord *hasIdDepField(const Expr *Expression); + /// Stores the location an ID-dependent variable is created from a call to + /// an ID function in IdDepVarsMap. + void saveIdDepVar(const Stmt *Statement, const VarDecl *Variable); + /// Stores the location an ID-dependent field is created from a call to an ID + /// function in IdDepFieldsMap. + void saveIdDepField(const Stmt *Statement, const FieldDecl *Field); + /// Stores the location an ID-dependent variable is created from a reference + /// to another ID-dependent variable or field in IdDepVarsMap. + void saveIdDepVarFromReference(const DeclRefExpr *RefExpr, + const MemberExpr *MemExpr, + const VarDecl *PotentialVar); + /// Stores the location an ID-dependent field is created from a reference to + /// another ID-dependent variable or field in IdDepFieldsMap. + void saveIdDepFieldFromReference(const DeclRefExpr *RefExpr, + const MemberExpr *MemExpr, + const FieldDecl *PotentialField); + /// Returns the loop type. + LoopType getLoopType(const Stmt *Loop); + +public: + IdDependentBackwardBranchCheck(StringRef Name, ClangTidyContext *Context) + : ClangTidyCheck(Name, Context) {} + void registerMatchers(ast_matchers::MatchFinder *Finder) override; + void check(const ast_matchers::MatchFinder::MatchResult &Result) override; +}; + +} // namespace altera +} // namespace tidy +} // namespace clang + +#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_ID_DEPENDENT_BACKWARD_BRANCH_H Index: clang-tidy/altera/IdDependentBackwardBranchCheck.cpp =================================================================== --- /dev/null +++ clang-tidy/altera/IdDependentBackwardBranchCheck.cpp @@ -0,0 +1,295 @@ +//===--- IdDependentBackwardBranchCheck.cpp - clang-tidy ------------------===// +// +// The LLVM Compiler Infrastructure +// +// This file is distributed under the University of Illinois Open Source +// License. See LICENSE.TXT for details. +// +//===----------------------------------------------------------------------===// + +#include "IdDependentBackwardBranchCheck.h" +#include "clang/AST/ASTContext.h" +#include "clang/ASTMatchers/ASTMatchFinder.h" +#include +#include + +using namespace clang::ast_matchers; + +namespace clang { +namespace tidy { +namespace altera { + +void IdDependentBackwardBranchCheck::registerMatchers(MatchFinder *Finder) { + // Prototype to identify all variables which hold a thread-variant ID + // First Matcher just finds all the direct assignments of either ID call + const auto TID_RHS = expr(hasDescendant(callExpr(callee(functionDecl( + anyOf(hasName("get_global_id"), hasName("get_local_id"))))))); + + const auto ANY_ASSIGN = anyOf( + hasOperatorName("="), hasOperatorName("*="), hasOperatorName("/="), + hasOperatorName("%="), hasOperatorName("+="), hasOperatorName("-="), + hasOperatorName("<<="), hasOperatorName(">>="), hasOperatorName("&="), + hasOperatorName("^="), hasOperatorName("|=")); + + Finder->addMatcher( + compoundStmt( + // Bind on actual get_local/global_id calls + forEachDescendant( + stmt(anyOf(declStmt(hasDescendant(varDecl(hasInitializer(TID_RHS)) + .bind("tid_dep_var"))), + binaryOperator(allOf( + ANY_ASSIGN, hasRHS(TID_RHS), + hasLHS(anyOf( + declRefExpr(to(varDecl().bind("tid_dep_var"))), + memberExpr(member( + fieldDecl().bind("tid_dep_field"))))))))) + .bind("straight_assignment"))), + this); + + // Bind all VarDecls that include an initializer with a variable DeclRefExpr + // (incase it is ID-dependent) + Finder->addMatcher( + stmt(forEachDescendant( + varDecl( + hasInitializer(forEachDescendant(stmt(anyOf( + declRefExpr(to(varDecl())).bind("assign_ref_var"), + memberExpr(member(fieldDecl())).bind("assign_ref_field")))))) + .bind("pot_tid_var"))), + this); + + // Bind all VarDecls that are assigned a value with a variable DeclRefExpr (in + // case it is ID-dependent) + Finder->addMatcher( + stmt(forEachDescendant(binaryOperator(allOf( + ANY_ASSIGN, + hasRHS(forEachDescendant(stmt(anyOf( + declRefExpr(to(varDecl())).bind("assign_ref_var"), + memberExpr(member(fieldDecl())).bind("assign_ref_field"))))), + hasLHS( + anyOf(declRefExpr(to(varDecl().bind("pot_tid_var"))), + memberExpr(member(fieldDecl().bind("pot_tid_field"))))))))), + this); + + // Second Matcher looks for branch statements inside of loops and bind on the + // condition expression IF it either calls an ID function or has a variable + // DeclRefExpr DeclRefExprs are checked later to confirm whether the variable + // is ID-dependent + const auto COND_EXPR = + expr(anyOf(hasDescendant(callExpr(callee(functionDecl( + anyOf(hasName("get_global_id"), + hasName("get_local_id"))))) + .bind("id_call")), + hasDescendant(stmt(anyOf(declRefExpr(to(varDecl())), + memberExpr(member(fieldDecl()))))))) + .bind("cond_expr"); + Finder->addMatcher(stmt(anyOf(forStmt(hasCondition(COND_EXPR)), + doStmt(hasCondition(COND_EXPR)), + whileStmt(hasCondition(COND_EXPR)))) + .bind("backward_branch"), + this); +} + +IdDependentBackwardBranchCheck::IdDependencyRecord * +IdDependentBackwardBranchCheck::hasIdDepVar(const Expr *Expression) { + if (const DeclRefExpr *expr = dyn_cast(Expression)) { + // It is a DeclRefExpr, so check if it's an ID-dependent variable + const VarDecl *CheckVariable = dyn_cast(expr->getDecl()); + auto FoundVariable = IdDepVarsMap.find(CheckVariable); + if (FoundVariable == IdDepVarsMap.end()) { + return nullptr; + } + return &(FoundVariable->second); + } + for (auto i = Expression->child_begin(), e = Expression->child_end(); i != e; + ++i) { + if (auto *ChildExpression = dyn_cast(*i)) { + auto Result = hasIdDepVar(ChildExpression); + if (Result) { + return Result; + } + } + } + return nullptr; +} + +IdDependentBackwardBranchCheck::IdDependencyRecord * +IdDependentBackwardBranchCheck::hasIdDepField(const Expr *Expression) { + if (const MemberExpr *MemberExpression = dyn_cast(Expression)) { + const FieldDecl *CheckField = + dyn_cast(MemberExpression->getMemberDecl()); + auto FoundField = IdDepFieldsMap.find(CheckField); + if (FoundField == IdDepFieldsMap.end()) { + return nullptr; + } + return &(FoundField->second); + } + for (auto I = Expression->child_begin(), E = Expression->child_end(); I != E; + ++I) { + if (auto *ChildExpression = dyn_cast(*I)) { + auto Result = hasIdDepField(ChildExpression); + if (Result) { + return Result; + } + } + } + return nullptr; +} + +void IdDependentBackwardBranchCheck::saveIdDepVar(const Stmt *Statement, + const VarDecl *Variable) { + // Record that this variable is thread-dependent + std::ostringstream StringStream; + StringStream << "assignment of ID-dependent variable " + << Variable->getNameAsString(); + IdDepVarsMap[Variable] = + IdDependencyRecord(Variable, Variable->getBeginLoc(), StringStream.str()); +} + +void IdDependentBackwardBranchCheck::saveIdDepField(const Stmt *Statement, + const FieldDecl *Field) { + std::ostringstream StringStream; + StringStream << "assignment of ID-dependent field " + << Field->getNameAsString(); + IdDepFieldsMap[Field] = + IdDependencyRecord(Field, Statement->getBeginLoc(), StringStream.str()); +} + +void IdDependentBackwardBranchCheck::saveIdDepVarFromReference( + const DeclRefExpr *RefExpr, const MemberExpr *MemExpr, + const VarDecl *PotentialVar) { + // If the variable is already in IdDepVarsMap, ignore it + if (IdDepVarsMap.find(PotentialVar) != IdDepVarsMap.end()) { + return; + } + std::ostringstream StringStream; + StringStream << "inferred assignment of ID-dependent value from " + "ID-dependent "; + if (RefExpr) { + const auto RefVar = dyn_cast(RefExpr->getDecl()); + // If variable isn't ID-dependent, but refVar is + if (IdDepVarsMap.find(RefVar) != IdDepVarsMap.end()) { + StringStream << "variable " << RefVar->getNameAsString(); + } + } + if (MemExpr) { + const auto RefField = dyn_cast(MemExpr->getMemberDecl()); + // If variable isn't ID-dependent, but refField is + if (IdDepFieldsMap.find(RefField) != IdDepFieldsMap.end()) { + StringStream << "member " << RefField->getNameAsString(); + } + } + IdDepVarsMap[PotentialVar] = IdDependencyRecord( + PotentialVar, PotentialVar->getBeginLoc(), StringStream.str()); +} + +void IdDependentBackwardBranchCheck::saveIdDepFieldFromReference( + const DeclRefExpr *RefExpr, const MemberExpr *MemExpr, + const FieldDecl *PotentialField) { + // If the field is already in IdDepFieldsMap, ignore it + if (IdDepFieldsMap.find(PotentialField) != IdDepFieldsMap.end()) { + return; + } + std::ostringstream StringStream; + StringStream << "inferred assignment of ID-dependent member from " + "ID-dependent "; + if (RefExpr) { + const auto RefVar = dyn_cast(RefExpr->getDecl()); + // If field isn't ID-dependent, but RefVar is + if (IdDepVarsMap.find(RefVar) != IdDepVarsMap.end()) { + StringStream << "variable " << RefVar->getNameAsString(); + } + } + if (MemExpr) { + const auto RefField = dyn_cast(MemExpr->getMemberDecl()); + if (IdDepFieldsMap.find(RefField) != IdDepFieldsMap.end()) { + StringStream << "member " << RefField->getNameAsString(); + } + } + IdDepFieldsMap[PotentialField] = IdDependencyRecord( + PotentialField, PotentialField->getBeginLoc(), StringStream.str()); +} + +IdDependentBackwardBranchCheck::LoopType +IdDependentBackwardBranchCheck::getLoopType(const Stmt *Loop) { + if (const auto DoLoop = dyn_cast(Loop)) { + return DO_LOOP; // loop_type = 0; + } else if (const auto WhileLoop = dyn_cast(Loop)) { + return WHILE_LOOP; // loop_type = 1; + } else if (const auto ForLoop = dyn_cast(Loop)) { + return FOR_LOOP; // loop_type = 2; + } + return UNK_LOOP; +} + +void IdDependentBackwardBranchCheck::check( + const MatchFinder::MatchResult &Result) { + // The first half of the callback only deals with identifying and propagating + // ID-dependency information into the IdDepVars vector + const auto *Variable = Result.Nodes.getNodeAs("tid_dep_var"); + const auto *Field = Result.Nodes.getNodeAs("tid_dep_field"); + const auto *Statement = Result.Nodes.getNodeAs("straight_assignment"); + const auto *RefExpr = Result.Nodes.getNodeAs("assign_ref_var"); + const auto *MemExpr = Result.Nodes.getNodeAs("assign_ref_field"); + const auto *PotentialVar = Result.Nodes.getNodeAs("pot_tid_var"); + const auto *PotentialField = + Result.Nodes.getNodeAs("pot_tid_field"); + + // Save variables and fields assigned directly through ID function calls + if (Statement && (Variable || Field)) { + if (Variable) { + saveIdDepVar(Statement, Variable); + } else if (Field) { + saveIdDepField(Statement, Field); + } + } + + // Save variables assigned to values of Id-dependent variables and fields + if ((RefExpr || MemExpr) && PotentialVar) { + saveIdDepVarFromReference(RefExpr, MemExpr, PotentialVar); + } + + // Save fields assigned to values of ID-dependent variables and fields + if ((RefExpr || MemExpr) && PotentialField) { + saveIdDepFieldFromReference(RefExpr, MemExpr, PotentialField); + } + + // The second part of the callback deals with checking if a branch inside a + // loop is thread dependent + const auto *CondExpr = Result.Nodes.getNodeAs("cond_expr"); + const auto *IDCall = Result.Nodes.getNodeAs("id_call"); + const auto *Loop = Result.Nodes.getNodeAs("backward_branch"); + if (!Loop) { + return; + } + LoopType Type = getLoopType(Loop); + if (CondExpr) { + if (IDCall) { // If conditional expression calls an ID function directly + diag(CondExpr->getBeginLoc(), + "backward branch (%select{do|while|for}0 loop) is ID-dependent due " + "to ID function call and may cause performance degradation") + << Type; + } else { // Conditional expression has DeclRefExpr(s), check ID-dependency + IdDependencyRecord *IdDepVar = hasIdDepVar(CondExpr); + IdDependencyRecord *IdDepField = hasIdDepField(CondExpr); + if (IdDepVar) { + diag(IdDepVar->Location, IdDepVar->Message); + diag( + CondExpr->getBeginLoc(), + "backward branch (%select{do|while|for}0 loop) is ID-dependent due " + "to variable reference to %1 and may cause performance degradation") + << Type << IdDepVar->VariableDeclaration; + } else if (IdDepField) { + diag(IdDepField->Location, IdDepField->Message); + diag( + CondExpr->getBeginLoc(), + "backward branch (%select{do|while|for}0 loop) is ID-dependent due " + "to member reference to %1 and may cause performance degradation") + << Type << IdDepField->FieldDeclaration; + } + } + } +} + +} // namespace altera +} // namespace tidy +} // namespace clang Index: docs/ReleaseNotes.rst =================================================================== --- docs/ReleaseNotes.rst +++ docs/ReleaseNotes.rst @@ -67,6 +67,17 @@ 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-id-dependent-backward-branch + ` check. + + Finds ID-dependent variables and fields that are used within loops. + - New :doc:`bugprone-dynamic-static-initializers ` check. Index: docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst =================================================================== --- /dev/null +++ docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst @@ -0,0 +1,29 @@ +.. title:: clang-tidy - fpga-id-dependent-backward-branch + +fpga-id-dependent-backward-branch +================================= + +Finds ID-dependent variables and fields that are used within loops. This causes +branches to occur inside the loops, and thus leads to performance degradation. + +Based on the `Altera SDK for OpenCL: Best Practices Guide +`_. + +.. code-block:: c++ + + // The following code will produce a warning because this ID-dependent + // variable is used in a loop condition statement. + int ThreadID = get_local_id(0); + + // The following loop will produce a warning because the loop condition + // statement depends on an ID-dependent variable. + for (int i = 0; i < ThreadID; ++i) { + std::cout << i << std::endl; + } + + // The following loop will not produce a warning, because the ID-dependent + // variable is not used in the loop condition statement. + for (int i = 0; i < 100; ++i) { + std::cout << ThreadID << std::endl; + } + 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-id-dependent-backward-branch.cpp =================================================================== --- /dev/null +++ test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp @@ -0,0 +1,83 @@ +// RUN: %check_clang_tidy %s altera-id-dependent-backward-branch %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h + +typedef struct ExampleStruct { + int IDDepField; +} ExampleStruct; + +void error() { + // ==== Assignments ==== + int ThreadID = get_local_id(0); +// CHECK-NOTES: :[[@LINE-1]]:3: warning: assignment of ID-dependent variable ThreadID [altera-id-dependent-backward-branch] + + ExampleStruct Example; + Example.IDDepField = get_local_id(0); +// CHECK-NOTES: :[[@LINE-1]]:3: warning: assignment of ID-dependent field IDDepField [altera-id-dependent-backward-branch] + + // ==== Inferred Assignments ==== + int ThreadID2 = ThreadID * get_local_size(0); +// CHECK-NOTES: :[[@LINE-1]]:3: warning: inferred assignment of ID-dependent value from ID-dependent variable ThreadID [altera-id-dependent-backward-branch] + + int ThreadID3 = Example.IDDepField; // OK: not used in any loops + + ExampleStruct UnusedStruct = { + ThreadID * 2 // OK: not used in any loops + }; + + // ==== Conditional Expressions ==== + int accumulator = 0; + for (int i = 0; i < get_local_id(0); i++) { + // CHECK-NOTES: :[[@LINE-1]]:19: warning: backward branch (for loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + int j = 0; + while (j < get_local_id(0)) { + // CHECK-NOTES: :[[@LINE-1]]:10: warning: backward branch (while loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + do { + accumulator++; + } while (j < get_local_id(0)); + // CHECK-NOTES: :[[@LINE-1]]:12: warning: backward branch (do loop) is ID-dependent due to ID function call and may cause performance degradation [altera-id-dependent-backward-branch] + + for (int i = 0; i < ThreadID2; i++) { + // CHECK-NOTES: :[[@LINE-1]]:19: warning: backward branch (for loop) is ID-dependent due to variable reference to 'ThreadID2' and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + while (j < ThreadID) { + // CHECK-NOTES: :[[@LINE-1]]:10: warning: backward branch (while loop) is ID-dependent due to variable reference to 'ThreadID' and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + do { + accumulator++; + } while (j < ThreadID); + // CHECK-NOTES: :[[@LINE-1]]:12: warning: backward branch (do loop) is ID-dependent due to variable reference to 'ThreadID' and may cause performance degradation [altera-id-dependent-backward-branch] + + for (int i = 0; i < Example.IDDepField; i++) { + // CHECK-NOTES: :[[@LINE-1]]:19: warning: backward branch (for loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + while (j < Example.IDDepField) { + // CHECK-NOTES: :[[@LINE-1]]:10: warning: backward branch (while loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch] + accumulator++; + } + + do { + accumulator++; + } while (j < Example.IDDepField); + // CHECK-NOTES: :[[@LINE-1]]:12: warning: backward branch (do loop) is ID-dependent due to member reference to 'IDDepField' and may cause performance degradation [altera-id-dependent-backward-branch] +} + +void success() { + int accumulator = 0; + + for (int i = 0; i < 1000; i++) { + if (i < get_local_id(0)) { + accumulator++; + } + } +}