ffrankies updated this revision to Diff 254304.
ffrankies added a comment.
- Updated underlying repo to https://github.com/llvm/llvm-project
- Removed braces from one-line if-statements
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D72241/new/
https://reviews.llvm.org/D72241
Files:
clang-tools-extra/clang-tidy/CMakeLists.txt
clang-tools-extra/clang-tidy/ClangTidyForceLinker.h
clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
clang-tools-extra/clang-tidy/altera/CMakeLists.txt
clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp
clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.h
clang-tools-extra/docs/ReleaseNotes.rst
clang-tools-extra/docs/clang-tidy/checks/altera-single-work-item-barrier.rst
clang-tools-extra/docs/clang-tidy/checks/list.rst
clang-tools-extra/docs/clang-tidy/index.rst
clang-tools-extra/test/clang-tidy/checkers/altera-single-work-item-barrier.cpp
Index: clang-tools-extra/test/clang-tidy/checkers/altera-single-work-item-barrier.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/test/clang-tidy/checkers/altera-single-work-item-barrier.cpp
@@ -0,0 +1,294 @@
+// RUN: %check_clang_tidy -check-suffix=OLD %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLD
+// RUN: %check_clang_tidy -check-suffix=NEW %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DNEW
+// RUN: %check_clang_tidy -check-suffix=AOCOLD %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DAOCOLD
+// RUN: %check_clang_tidy -check-suffix=AOCNEW %s altera-single-work-item-barrier %t -- -config='{CheckOptions: [{key: altera-single-work-item-barrier.AOCVersion, value: 1701}]}' -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DAOCNEW
+
+#ifdef OLD
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-OLD: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id' does not call get_global_id or get_local_id and will be treated as single-work-item.{{[[:space:]]}}Barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 may error out [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void __kernel success_barrier_global_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void __kernel success_barrier_local_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void __kernel success_barrier_both_ids(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+
+void success_nokernel_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void success_nokernel_barrier_global_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void success_nokernel_barrier_local_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void success_nokernel_barrier_both_ids(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+#endif
+
+#ifdef NEW
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-NEW: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id' does not call get_global_id or get_local_id and will be treated as single-work-item.{{[[:space:]]}}Barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 may error out [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void __kernel success_barrier_global_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void __kernel success_barrier_local_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void __kernel success_barrier_both_ids(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+
+void success_nokernel_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void success_nokernel_barrier_global_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void success_nokernel_barrier_local_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void success_nokernel_barrier_both_ids(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+#endif
+
+#ifdef AOCOLD
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-AOCOLD: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id' does not call get_global_id or get_local_id may be a viable single work-item kernel, but barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 will force NDRange execution. If single work-item semantics are desired a mem_fence may be more efficient. [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+__attribute__ ((reqd_work_group_size(1,1,1)))
+void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-AOCOLD: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id_work_group_size' does not call get_global_id or get_local_id may be a viable single work-item kernel, but barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 will force NDRange execution. If single work-item semantics are desired a mem_fence may be more efficient. [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+__attribute__ ((reqd_work_group_size(2,1,1)))
+void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void __kernel success_barrier_global_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void __kernel success_barrier_local_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void __kernel success_barrier_both_ids(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+
+void success_nokernel_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void success_nokernel_barrier_global_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void success_nokernel_barrier_local_id(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void success_nokernel_barrier_both_ids(__global int * foo, int size) {
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+#endif
+
+#ifdef AOCNEW
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-AOCNEW: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id' does not call get_global_id or get_local_id may be a viable single work-item kernel, but barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 will force NDRange execution. If single work-item semantics are desired a mem_fence may be more efficient. [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+__attribute__ ((reqd_work_group_size(1,1,1)))
+void __kernel error_barrier_no_id_work_group_size(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ // CHECK-MESSAGES-AOCNEW: :[[@LINE-7]]:15: warning: Kernel function 'error_barrier_no_id_work_group_size' does not call get_global_id or get_local_id may be a viable single work-item kernel, but barrier call at {{(\/)?([^\/\0]+(\/)?)+}}:[[@LINE-1]]:3 will force NDRange execution. If single work-item semantics are desired a mem_fence may be more efficient. [altera-single-work-item-barrier]
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+__attribute__ ((reqd_work_group_size(2,1,1)))
+void __kernel success_barrier_no_id_work_group_size(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void __kernel success_barrier_global_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void __kernel success_barrier_local_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void __kernel success_barrier_both_ids(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+
+void success_nokernel_barrier_no_id(__global int * foo, int size) {
+ for (int j = 0; j < 256; j++) {
+ for (int i = 256; i < size; i+= 256) {
+ foo[j] += foo[j+i];
+ }
+ }
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ for (int i = 1; i < 256; i++) {
+ foo[0] += foo[i];
+ }
+}
+
+void success_nokernel_barrier_global_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_global_id(0);
+}
+
+void success_nokernel_barrier_local_id(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int tid = get_local_id(0);
+}
+
+void success_nokernel_barrier_both_ids(__global int * foo, int size) {
+ work_group_barrier(CLK_GLOBAL_MEM_FENCE);
+ int gid = get_global_id(0);
+ int lid = get_local_id(0);
+}
+#endif
Index: clang-tools-extra/docs/clang-tidy/index.rst
===================================================================
--- clang-tools-extra/docs/clang-tidy/index.rst
+++ clang-tools-extra/docs/clang-tidy/index.rst
@@ -58,6 +58,7 @@
Name prefix Description
====================== =========================================================
``abseil-`` Checks related to Abseil library.
+``altera-`` Checks related to OpenCL programming for FPGAs.
``android-`` Checks related to Android.
``boost-`` Checks related to Boost library.
``bugprone-`` Checks that target bugprone code constructs.
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
@@ -29,6 +29,7 @@
`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-single-work-item-barrier <altera-single-work-item-barrier.html>`_, "Yes"
`android-cloexec-accept <android-cloexec-accept.html>`_, "Yes"
`android-cloexec-accept4 <android-cloexec-accept4.html>`_,
`android-cloexec-creat <android-cloexec-creat.html>`_, "Yes"
Index: clang-tools-extra/docs/clang-tidy/checks/altera-single-work-item-barrier.rst
===================================================================
--- /dev/null
+++ clang-tools-extra/docs/clang-tidy/checks/altera-single-work-item-barrier.rst
@@ -0,0 +1,52 @@
+.. title:: clang-tidy - altera-single-work-item-barrier
+
+altera-single-work-item-barrier
+=============================
+
+Finds OpenCL kernel functions that call a barrier function but do not call
+an ID function.
+
+These kernel functions will be treated as single work-item kernels, which
+could be inefficient or lead to errors.
+
+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>`_.
+
+Examples:
+
+.. code-block:: c++
+
+ // error: function calls barrier but does not call an ID function.
+ void __kernel barrier_no_id(__global int * foo, int size) {
+ for (int i = 0; i < 100; i++) {
+ foo[i] += 5;
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ }
+
+ // ok: function calls barrier and an ID function.
+ void __kernel barrier_with_id(__global int * foo, int size) {
+ for (int i = 0; i < 100; i++) {
+ int tid = get_global_id(0);
+ foo[tid] += 5;
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ }
+
+ // ok with AOC Version 17.01: the reqd_work_group_size turns this into
+ // an NDRange.
+ __attribute__((reqd_work_group_size(2,2,2)))
+ void __kernel barrier_with_id(__global int * foo, int size) {
+ for (int i = 0; i < 100; i++) {
+ foo[tid] += 5;
+ }
+ barrier(CLK_GLOBAL_MEM_FENCE);
+ }
+
+Options
+-------
+
+.. option:: AOCVersion
+
+ Defines the version of the Altera Offline Compiler. Defaults to `1600`
+ (corresponding to version 16.00).
Index: clang-tools-extra/docs/ReleaseNotes.rst
===================================================================
--- clang-tools-extra/docs/ReleaseNotes.rst
+++ clang-tools-extra/docs/ReleaseNotes.rst
@@ -67,9 +67,24 @@
Improvements to clang-tidy
--------------------------
+New modules
+^^^^^^^^^^^
+
+- New :doc:`altera <clang-tidy/modules/altera>` module.
+
+ Includes checks related to OpenCL for FPGA coding guidelines, 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>`_.
+
New checks
^^^^^^^^^^
+- New :doc:`altera-single-work-item-barrier
+ <clang-tidy/checks/altera-single-work-item-barrier>` check.
+
+ Finds OpenCL kernel functions that call a barrier function but do not call
+ an ID function.
+
- New :doc:`bugprone-misplaced-pointer-arithmetic-in-alloc
<clang-tidy/checks/bugprone-misplaced-pointer-arithmetic-in-alloc>` check.
Index: clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.h
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.h
@@ -0,0 +1,40 @@
+//===--- SingleWorkItemBarrierCheck.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_SINGLE_WORK_ITEM_BARRIER_H
+#define LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_SINGLE_WORK_ITEM_BARRIER_H
+
+#include "../ClangTidy.h"
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+/// Detects OpenCL kernel functions that call a barrier but do not call an
+/// ID-function function. These functions will be treated as single work-item
+/// kernels, which may be inefficient or cause an error.
+///
+/// For the user-facing documentation see:
+/// http://clang.llvm.org/extra/clang-tidy/checks/opencl-single-work-item-barrier.html
+class SingleWorkItemBarrierCheck : public ClangTidyCheck {
+ const unsigned AOCVersion;
+
+public:
+ SingleWorkItemBarrierCheck(StringRef Name, ClangTidyContext *Context)
+ : ClangTidyCheck(Name, Context),
+ AOCVersion(Options.get("AOCVersion", 1600U)) {}
+ void registerMatchers(ast_matchers::MatchFinder *Finder) override;
+ void check(const ast_matchers::MatchFinder::MatchResult &Result) override;
+ void storeOptions(ClangTidyOptions::OptionMap &Opts) override;
+};
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang
+
+#endif // LLVM_CLANG_TOOLS_EXTRA_CLANG_TIDY_ALTERA_SINGLE_WORK_ITEM_BARRIER_H
Index: clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp
@@ -0,0 +1,88 @@
+//===--- SingleWorkItemBarrierCheck.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 "SingleWorkItemBarrierCheck.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+void SingleWorkItemBarrierCheck::registerMatchers(MatchFinder *Finder) {
+ // Find any function that calls barrier but does not call either get_local_id
+ // or get_global_id, and will thus be treated as a single-work-item kernel
+ // hasAttr(attr::Kind::OpenCLKernel) restricts it to only kernel functions.
+ // FIXME: Have it accept all functions but check for a parameter that gets an
+ // ID from one of the two ID functions.
+ Finder->addMatcher(
+ // Find function declarations...
+ functionDecl(
+ allOf(
+ // That are OpenCL kernels...
+ hasAttr(attr::Kind::OpenCLKernel),
+ // And call a barrier function (either 1.x or 2.x version)...
+ forEachDescendant(callExpr(callee(functionDecl(anyOf(
+ hasName("barrier"),
+ hasName("work_group_barrier")))))
+ .bind("barrier")),
+ // But do not call an ID function.
+ unless(hasDescendant(callExpr(callee(functionDecl(anyOf(
+ hasName("get_global_id"), hasName("get_local_id")))))))))
+ .bind("function"),
+ this);
+}
+
+void SingleWorkItemBarrierCheck::check(const MatchFinder::MatchResult &Result) {
+ const auto *MatchedDecl = Result.Nodes.getNodeAs<FunctionDecl>("function");
+ const auto *MatchedBarrier = Result.Nodes.getNodeAs<CallExpr>("barrier");
+ // If reqd_work_group_size is anything other than (1,1,1), it will be
+ // interpreted as an NDRange in AOC version 17.1.
+ bool IsNDRange = false;
+ for (const Attr *Attribute : MatchedDecl->getAttrs()) {
+ if (Attribute->getKind() == attr::Kind::ReqdWorkGroupSize) {
+ const auto *RWGSAttribute =
+ static_cast<const ReqdWorkGroupSizeAttr *>(Attribute);
+ if (RWGSAttribute->getXDim() > 1 || RWGSAttribute->getYDim() > 1 ||
+ RWGSAttribute->getZDim() > 1)
+ IsNDRange = true;
+ break;
+ }
+ }
+ if (AOCVersion < 1701)
+ diag(MatchedDecl->getLocation(),
+ "Kernel function %0 does not call get_global_id or get_local_id and "
+ "will be treated as single-work-item.\nBarrier call at %1 may error "
+ "out")
+ << MatchedDecl
+ << MatchedBarrier->getBeginLoc().printToString(
+ Result.Context->getSourceManager());
+ else {
+ if (IsNDRange) // No warning if kernel is treated as an NDRange.
+ return;
+ diag(MatchedDecl->getLocation(),
+ "Kernel function %0 does not call get_global_id or get_local_id may "
+ "be a viable single work-item kernel, but barrier call at %1 will "
+ "force NDRange execution. If single work-item semantics are desired a "
+ "mem_fence may be more efficient.")
+ << MatchedDecl
+ << MatchedBarrier->getBeginLoc().printToString(
+ Result.Context->getSourceManager());
+ }
+}
+
+void SingleWorkItemBarrierCheck::storeOptions(
+ ClangTidyOptions::OptionMap &Opts) {
+ Options.store(Opts, "AOCVersion", AOCVersion);
+}
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang
Index: clang-tools-extra/clang-tidy/altera/CMakeLists.txt
===================================================================
--- /dev/null
+++ clang-tools-extra/clang-tidy/altera/CMakeLists.txt
@@ -0,0 +1,15 @@
+set(LLVM_LINK_COMPONENTS support)
+
+add_clang_library(clangTidyAlteraModule
+ AlteraTidyModule.cpp
+ SingleWorkItemBarrierCheck.cpp
+
+ LINK_LIBS
+ clangAnalysis
+ clangAST
+ clangASTMatchers
+ clangBasic
+ clangLex
+ clangTidy
+ clangTidyUtils
+ )
Index: clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp
===================================================================
--- /dev/null
+++ clang-tools-extra/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 "SingleWorkItemBarrierCheck.h"
+
+using namespace clang::ast_matchers;
+
+namespace clang {
+namespace tidy {
+namespace altera {
+
+class AlteraModule : public ClangTidyModule {
+public:
+ void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
+ CheckFactories.registerCheck<SingleWorkItemBarrierCheck>(
+ "altera-single-work-item-barrier");
+ }
+};
+
+} // namespace altera
+
+// Register the AlteraTidyModule using this statically initialized variable.
+static ClangTidyModuleRegistry::Add<altera::AlteraModule>
+ 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-tools-extra/clang-tidy/ClangTidyForceLinker.h
===================================================================
--- clang-tools-extra/clang-tidy/ClangTidyForceLinker.h
+++ clang-tools-extra/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-tools-extra/clang-tidy/CMakeLists.txt
===================================================================
--- clang-tools-extra/clang-tidy/CMakeLists.txt
+++ clang-tools-extra/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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
