[clang-tools-extra] 83af66e - new altera ID dependent backward branch check
Aaron Ballman via cfe-commits
cfe-commits at lists.llvm.org
Thu May 6 14:02:37 PDT 2021
Author: Frank Derry Wanye
Date: 2021-05-06T17:01:39-04:00
New Revision: 83af66e18e3d3760d56ea7e3bdbff3428ae7730d
URL: https://github.com/llvm/llvm-project/commit/83af66e18e3d3760d56ea7e3bdbff3428ae7730d
DIFF: https://github.com/llvm/llvm-project/commit/83af66e18e3d3760d56ea7e3bdbff3428ae7730d.diff
LOG: new altera ID dependent backward branch check
This lint check is a part of the FLOCL (FPGA Linters for OpenCL) project
out of the Synergy Lab at Virginia Tech.
FLOCL is a set of lint checks aimed at FPGA developers who write code
in OpenCL.
The altera ID dependent backward branch lint check finds ID dependent
variables and fields used within loops, and warns of their usage. Using
these variables in loops can lead to performance degradation.
Added:
clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp
clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h
clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst
clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp
Modified:
clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
clang-tools-extra/clang-tidy/altera/CMakeLists.txt
clang-tools-extra/docs/ReleaseNotes.rst
clang-tools-extra/docs/clang-tidy/checks/list.rst
Removed:
################################################################################
diff --git a/clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp b/clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
index a6c29e03f7aa8..17458771d6358 100644
--- a/clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
+++ b/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 @@ namespace altera {
class AlteraModule : public ClangTidyModule {
public:
void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
+ CheckFactories.registerCheck<IdDependentBackwardBranchCheck>(
+ "altera-id-dependent-backward-branch");
CheckFactories.registerCheck<KernelNameRestrictionCheck>(
"altera-kernel-name-restriction");
CheckFactories.registerCheck<SingleWorkItemBarrierCheck>(
diff --git a/clang-tools-extra/clang-tidy/altera/CMakeLists.txt b/clang-tools-extra/clang-tidy/altera/CMakeLists.txt
index c8a883d0750c5..cf8104b9d8426 100644
--- a/clang-tools-extra/clang-tidy/altera/CMakeLists.txt
+++ b/clang-tools-extra/clang-tidy/altera/CMakeLists.txt
@@ -5,6 +5,7 @@ set(LLVM_LINK_COMPONENTS
add_clang_library(clangTidyAlteraModule
AlteraTidyModule.cpp
+ IdDependentBackwardBranchCheck.cpp
KernelNameRestrictionCheck.cpp
SingleWorkItemBarrierCheck.cpp
StructPackAlignCheck.cpp
diff --git a/clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp b/clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.cpp
new file mode 100644
index 0000000000000..09cddb220877a
--- /dev/null
+++ b/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<DeclRefExpr>(Expression)) {
+ // It is a DeclRefExpr, so check if it's an ID-dependent variable.
+ const auto *CheckVariable = dyn_cast<VarDecl>(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<Expr>(Child))
+ if (IdDependencyRecord *Result = hasIdDepVar(ChildExpression))
+ return Result;
+ return nullptr;
+}
+
+IdDependentBackwardBranchCheck::IdDependencyRecord *
+IdDependentBackwardBranchCheck::hasIdDepField(const Expr *Expression) {
+ if (const auto *MemberExpression = dyn_cast<MemberExpr>(Expression)) {
+ const auto *CheckField =
+ dyn_cast<FieldDecl>(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<Expr>(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<VarDecl>(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<FieldDecl>(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<VarDecl>(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<FieldDecl>(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<VarDecl>("tid_dep_var");
+ const auto *Field = Result.Nodes.getNodeAs<FieldDecl>("tid_dep_field");
+ const auto *Statement = Result.Nodes.getNodeAs<Stmt>("straight_assignment");
+ const auto *RefExpr = Result.Nodes.getNodeAs<DeclRefExpr>("assign_ref_var");
+ const auto *MemExpr = Result.Nodes.getNodeAs<MemberExpr>("assign_ref_field");
+ const auto *PotentialVar = Result.Nodes.getNodeAs<VarDecl>("pot_tid_var");
+ const auto *PotentialField =
+ Result.Nodes.getNodeAs<FieldDecl>("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<Expr>("cond_expr");
+ const auto *IDCall = Result.Nodes.getNodeAs<CallExpr>("id_call");
+ const auto *Loop = Result.Nodes.getNodeAs<Stmt>("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
diff --git a/clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h b/clang-tools-extra/clang-tidy/altera/IdDependentBackwardBranchCheck.h
new file mode 100644
index 0000000000000..5ebfa8e2f1e07
--- /dev/null
+++ b/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<const VarDecl *, IdDependencyRecord> IdDepVarsMap;
+ // Stores the locations where ID-dependent fields are created.
+ std::map<const FieldDecl *, IdDependencyRecord> 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
diff --git a/clang-tools-extra/docs/ReleaseNotes.rst b/clang-tools-extra/docs/ReleaseNotes.rst
index a9c943795e945..28cb3b2c55529 100644
--- a/clang-tools-extra/docs/ReleaseNotes.rst
+++ b/clang-tools-extra/docs/ReleaseNotes.rst
@@ -95,6 +95,13 @@ New checks
Finds ``pthread_setcanceltype`` function calls where a thread's cancellation
type is set to asynchronous.
+- New :doc:`altera-id-dependent-backward-branch
+ <clang-tidy/checks/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
<clang-tidy/checks/altera-unroll-loops>` check.
diff --git a/clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst b/clang-tools-extra/docs/clang-tidy/checks/altera-id-dependent-backward-branch.rst
new file mode 100644
index 0000000000000..26dd1fb90925b
--- /dev/null
+++ b/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
+<https://www.altera.com/en_US/pdfs/literature/hb/opencl-sdk/aocl_optimization_guide.pdf>`_.
diff --git a/clang-tools-extra/docs/clang-tidy/checks/list.rst b/clang-tools-extra/docs/clang-tidy/checks/list.rst
index 8889be098b6df..913c7fde26a56 100644
--- a/clang-tools-extra/docs/clang-tidy/checks/list.rst
+++ b/clang-tools-extra/docs/clang-tidy/checks/list.rst
@@ -30,6 +30,7 @@ Clang-Tidy Checks
`abseil-time-comparison <abseil-time-comparison.html>`_, "Yes"
`abseil-time-subtraction <abseil-time-subtraction.html>`_, "Yes"
`abseil-upgrade-duration-conversions <abseil-upgrade-duration-conversions.html>`_, "Yes"
+ `altera-id-dependent-backward-branch <altera-id-dependent-backward-branch.html>`_,
`altera-kernel-name-restriction <altera-kernel-name-restriction.html>`_,
`altera-single-work-item-barrier <altera-single-work-item-barrier.html>`_,
`altera-struct-pack-align <altera-struct-pack-align.html>`_, "Yes"
diff --git a/clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp b/clang-tools-extra/test/clang-tidy/checkers/altera-id-dependent-backward-branch.cpp
new file mode 100644
index 0000000000000..a6dbab7b72fce
--- /dev/null
+++ b/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++;
+ }
+ }
+}
More information about the cfe-commits
mailing list