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 @@ -9,6 +9,7 @@ #include "../ClangTidy.h" #include "../ClangTidyModule.h" #include "../ClangTidyModuleRegistry.h" +#include "IdDependentBackwardBranchCheck.h" #include "KernelNameRestrictionCheck.h" #include "SingleWorkItemBarrierCheck.h" #include "StructPackAlignCheck.h" @@ -23,6 +24,8 @@ class AlteraModule : public ClangTidyModule { public: void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override { + CheckFactories.registerCheck( + "altera-id-dependent-backward-branch"); CheckFactories.registerCheck( "altera-kernel-name-restriction"); CheckFactories.registerCheck( 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 @@ -5,6 +5,7 @@ add_clang_library(clangTidyAlteraModule AlteraTidyModule.cpp + IdDependentBackwardBranchCheck.cpp KernelNameRestrictionCheck.cpp SingleWorkItemBarrierCheck.cpp StructPackAlignCheck.cpp Index: clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h =================================================================== --- /dev/null +++ clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h @@ -0,0 +1,83 @@ +//===--- 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_IDDEPENDENTBACKWARDBRANCHCHECK_H +#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_IDDEPENDENTBACKWARDBRANCHCHECK_H + +#include "../ClangTidyCheck.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 { UnknownLoop = -1, DoLoop = 0, WhileLoop = 1, ForLoop = 2 }; + // Stores information necessary for printing out source of error. + struct IdDependencyRecord { + IdDependencyRecord(const VarDecl *Declaration, SourceLocation Location, + const llvm::Twine &Message) + : VariableDeclaration(Declaration), Location(Location), + Message(Message.str()) {} + IdDependencyRecord(const FieldDecl *Declaration, SourceLocation Location, + const llvm::Twine &Message) + : FieldDeclaration(Declaration), Location(Location), + Message(Message.str()) {} + IdDependencyRecord() = default; + const VarDecl *VariableDeclaration = nullptr; + const FieldDecl *FieldDeclaration = nullptr; + 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_IDDEPENDENTBACKWARDBRANCHCHECK_H Index: clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp =================================================================== --- /dev/null +++ clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp @@ -0,0 +1,264 @@ +//===--- IdDependentBackwardBranchCheck.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 "IdDependentBackwardBranchCheck.h" +#include "clang/AST/ASTContext.h" +#include "clang/ASTMatchers/ASTMatchFinder.h" + +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 ThreadID = expr(hasDescendant(callExpr(callee(functionDecl( + anyOf(hasName("get_global_id"), hasName("get_local_id"))))))); + + const auto RefVarOrField = forEachDescendant( + stmt(anyOf(declRefExpr(to(varDecl())).bind("assign_ref_var"), + memberExpr(member(fieldDecl())).bind("assign_ref_field")))); + + Finder->addMatcher( + compoundStmt( + // Bind on actual get_local/global_id calls. + forEachDescendant( + stmt( + anyOf(declStmt(hasDescendant(varDecl(hasInitializer(ThreadID)) + .bind("tid_dep_var"))), + binaryOperator(allOf( + isAssignmentOperator(), hasRHS(ThreadID), + 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 + // (in case it is ID-dependent). + Finder->addMatcher( + stmt(forEachDescendant( + varDecl(hasInitializer(RefVarOrField)).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(isAssignmentOperator(), hasRHS(RefVarOrField), + 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 CondExpr = + 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(CondExpr)), + doStmt(hasCondition(CondExpr)), + whileStmt(hasCondition(CondExpr)))) + .bind("backward_branch"), + this); +} + +IdDependentBackwardBranchCheck::IdDependencyRecord * +IdDependentBackwardBranchCheck::hasIdDepVar(const Expr *Expression) { + if (const auto *Declaration = dyn_cast(Expression)) { + // It is a DeclRefExpr, so check if it's an ID-dependent variable. + const auto *CheckVariable = dyn_cast(Declaration->getDecl()); + auto FoundVariable = IdDepVarsMap.find(CheckVariable); + if (FoundVariable == IdDepVarsMap.end()) + return nullptr; + return &(FoundVariable->second); + } + for (const auto *Child : Expression->children()) + if (const auto *ChildExpression = dyn_cast(Child)) + if (IdDependencyRecord *Result = hasIdDepVar(ChildExpression)) + return Result; + return nullptr; +} + +IdDependentBackwardBranchCheck::IdDependencyRecord * +IdDependentBackwardBranchCheck::hasIdDepField(const Expr *Expression) { + if (const auto *MemberExpression = dyn_cast(Expression)) { + const auto *CheckField = + dyn_cast(MemberExpression->getMemberDecl()); + auto FoundField = IdDepFieldsMap.find(CheckField); + if (FoundField == IdDepFieldsMap.end()) + return nullptr; + return &(FoundField->second); + } + for (const auto *Child : Expression->children()) + if (const auto *ChildExpression = dyn_cast(Child)) + if (IdDependencyRecord *Result = hasIdDepField(ChildExpression)) + return Result; + return nullptr; +} + +void IdDependentBackwardBranchCheck::saveIdDepVar(const Stmt *Statement, + const VarDecl *Variable) { + // Record that this variable is thread-dependent. + IdDepVarsMap[Variable] = + IdDependencyRecord(Variable, Variable->getBeginLoc(), + Twine("assignment of ID-dependent variable ") + + Variable->getNameAsString()); +} + +void IdDependentBackwardBranchCheck::saveIdDepField(const Stmt *Statement, + const FieldDecl *Field) { + // Record that this field is thread-dependent. + IdDepFieldsMap[Field] = IdDependencyRecord( + Field, Statement->getBeginLoc(), + Twine("assignment of ID-dependent field ") + Field->getNameAsString()); +} + +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::string Message; + llvm::raw_string_ostream StringStream(Message); + 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(), Message); +} + +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::string Message; + llvm::raw_string_ostream StringStream(Message); + 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(), Message); +} + +IdDependentBackwardBranchCheck::LoopType +IdDependentBackwardBranchCheck::getLoopType(const Stmt *Loop) { + switch (Loop->getStmtClass()) { + case Stmt::DoStmtClass: + return DoLoop; + case Stmt::WhileStmtClass: + return WhileLoop; + case Stmt::ForStmtClass: + return ForLoop; + default: + return UnknownLoop; + } +} + +void IdDependentBackwardBranchCheck::check( + const MatchFinder::MatchResult &Result) { + // The first half of the callback only deals with identifying and storing + // ID-dependency information into the IdDepVars and IdDepFields maps. + 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) { // 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; + return; + } + // Conditional expression has DeclRefExpr(s), check ID-dependency. + IdDependencyRecord *IdDepVar = hasIdDepVar(CondExpr); + IdDependencyRecord *IdDepField = hasIdDepField(CondExpr); + if (IdDepVar) { + // Change one of these to a Note + diag(IdDepVar->Location, IdDepVar->Message, DiagnosticIDs::Note); + 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, DiagnosticIDs::Note); + 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: clang-tools-extra/docs/ReleaseNotes.rst =================================================================== --- clang-tools-extra/docs/ReleaseNotes.rst +++ clang-tools-extra/docs/ReleaseNotes.rst @@ -83,6 +83,13 @@ Finds ``pthread_setcanceltype`` function calls where a thread's cancellation type is set to asynchronous. +- New :doc:`altera-id-dependent-backward-branch + ` check. + + 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. + - New :doc:`altera-unroll-loops ` check. Index: clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst =================================================================== --- /dev/null +++ clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst @@ -0,0 +1,28 @@ +.. title:: clang-tidy - altera-id-dependent-backward-branch + +altera-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. + +.. 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; + } + +Based on the `Altera SDK for OpenCL: Best Practices Guide +`_. 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 @@ -30,6 +30,7 @@ `abseil-time-comparison `_, "Yes" `abseil-time-subtraction `_, "Yes" `abseil-upgrade-duration-conversions `_, "Yes" + `altera-id-dependent-backward-branch `_, `altera-kernel-name-restriction `_, `altera-single-work-item-barrier `_, `altera-struct-pack-align `_, "Yes" Index: clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp =================================================================== --- /dev/null +++ clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp @@ -0,0 +1,86 @@ +// 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() { + // ==== 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] + + // ==== Assignments ==== + int ThreadID = get_local_id(0); + + while (j < ThreadID) { + // CHECK-NOTES: :[[@LINE-3]]:3: note: assignment of ID-dependent variable ThreadID + // CHECK-NOTES: :[[@LINE-2]]: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++; + } + + ExampleStruct Example; + Example.IDDepField = get_local_id(0); + + // ==== Inferred Assignments ==== + int ThreadID2 = ThreadID * get_local_size(0); + + int ThreadID3 = Example.IDDepField; // OK: not used in any loops + + ExampleStruct UnusedStruct = { + ThreadID * 2 // OK: not used in any loops + }; + + for (int i = 0; i < ThreadID2; i++) { + // CHECK-NOTES: :[[@LINE-9]]:3: note: inferred assignment of ID-dependent value from ID-dependent variable ThreadID + // CHECK-NOTES: :[[@LINE-2]]: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++; + } + + do { + accumulator++; + } while (j < ThreadID); + // CHECK-NOTES: :[[@LINE-29]]:3: note: assignment of ID-dependent variable ThreadID + // CHECK-NOTES: :[[@LINE-2]]: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-24]]:3: note: assignment of ID-dependent field IDDepField + // CHECK-NOTES: :[[@LINE-2]]: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-30]]:3: note: assignment of ID-dependent field IDDepField + // CHECK-NOTES: :[[@LINE-2]]: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-38]]:3: note: assignment of ID-dependent field IDDepField + // CHECK-NOTES: :[[@LINE-2]]: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++; + } + } +}