[clang-tidy] new altera single work item barrier check

[ffrankies]

Summary

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 single work item barrier check 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."

Depends on https://reviews.llvm.org/D66564 due to the altera module being introduced there.

[Eugene.Zelenko]

I think will be good idea to separate module code in own review or refer to previous one of previous reviews as dependency.

clang-tidy/altera/SingleWorkItemBarrierCheck.cpp

51 ↗

(On Diff #236289)

const auto *.

69 ↗

(On Diff #236289)

Just IsNDRange. See readability-simplify-boolean-expr.

docs/clang-tidy/checks/altera-single-work-item-barrier.rst

51 ↗

(On Diff #236289)

Please highlight 1600 with single back-quotes.

[ffrankies]

Implemented requested changes by @Eugene.Zelenko

• Changed auto to const auto *
• Changed if(IsNDRange == true) to if(IsNDRange)
• Highlighted 1600 with single back-quotes
• Added link to https://reviews.llvm.org/D66564 as a dependency in Summary (is there someplace else I can refer to that as a dependency?)

[ffrankies]

• Updated underlying repo to https://github.com/llvm/llvm-project
• Removed braces from one-line if-statements

[ffrankies]

• Rebased code and fixed merge conflicts with D66564
• Added SingleWorkItemBarrierCheck.cpp to BUILD.gn

[aaron.ballman]

clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp

35

You can use hasAnyName() and pass both strings rather than two hasName() calls.

39

Same here.

50

Rather than getting all attributes, you can use specific_attrs<> to loop over just the specific attributes you care about on the declaration. That also fixes the non-idiomatic way to convert from an attribute to a specific derived attribute (which should use dyn_cast<> or friends).

62

clang-tidy diagnostics shouldn't start with a capital letter or contain terminating punctuation, the function names should be surrounded in single quotes, and we don't typically use newlines or print out source locations like that in diagnostics. How about:

kernel function %0 does not call 'get_global_id' or 'get_local_id' and is treated as a single work-item

and issue a note diagnostic barrier call may error out at the barrier location?

72

Same here as above.

Will users know what NDRange execution means? (I'm can't really understand what the diagnostic is telling me, but this is outside of my typical domain.)

clang-tools-extra/docs/clang-tidy/checks/altera-single-work-item-barrier.rst

8

Maybe link to documentation describing what an ID function is (or describe it more explicitly here)?

llvm/utils/gn/secondary/clang-tools-extra/clang-tidy/altera/BUILD.gn

16 ↗

(On Diff #295698)

FWIW, you don't need to maintain the gn files and can leave them out of your patch.

[ffrankies]

Implemented changes requested by @aaron.ballman

• using hasAnyName() instead of multiple hasName() calls in the matcher
• switched to a combination of hasAttr<>() and getAttr<>() to remove need for casting and looping over all attributes (did not use specific_attrs<>() because there should only be one ReqdWorkGroupSizeAttr per function)
• re-phrased diagnostic messages
• added all 4 ID functions to the documentation
• removed the update to BUILD.gn

[aaron.ballman]

LGTM!

[ffrankies]

@aaron.ballman Thank you! If there are no further comments, could you please commit this on my behalf? My GitHub username is ffrankies.

clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp

50

I went with hasAttr<> and then getAttr<> to avoid a loop over all attributes - there should only be one ReqdWorkGroupSizeAttr

72

I discussed this with other students in my lab and the consensus there is this should be clear for OpenCL/OpenCL-for-FPGA users. If really needed, we can try to re-phrase it, but then the diagnostics will most likely become more wordy.

[aaron.ballman]

In D72241#2449426, @ffrankies wrote:

@aaron.ballman Thank you! If there are no further comments, could you please commit this on my behalf? My GitHub username is ffrankies.

I'd be happy to except the diff seems to be causing me problems that I'm struggling to figure out how to solve.

F:\llvm-project>git apply "F:\Aaron Ballman\Desktop\D72241.diff"
F:/Aaron Ballman/Desktop/D72241.diff:207: trailing whitespace.
Based on the `Altera SDK for OpenCL: Best Practices Guide
F:/Aaron Ballman/Desktop/D72241.diff:213: trailing whitespace.

F:/Aaron Ballman/Desktop/D72241.diff:230: trailing whitespace.

F:/Aaron Ballman/Desktop/D72241.diff:279: trailing whitespace.
  // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a sing
le work-item and may error out
F:/Aaron Ballman/Desktop/D72241.diff:339: trailing whitespace.
  // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a sing
le work-item and may error out
error: clang-tidy/altera/AlteraTidyModule.cpp: No such file or directory
error: clang-tidy/altera/CMakeLists.txt: No such file or directory
error: docs/ReleaseNotes.rst: No such file or directory
error: docs/clang-tidy/checks/list.rst: No such file or directory

What's so strange is -- those files all exist for me. I tried switching the directory separators to be Windows separators, but then I got errors about invalid file paths. Can you rebase your patch on ToT and maybe that'll apply more cleanly for me?

[ffrankies]

@aaron.ballman hmm, that is strange. I've rebased the patch and updated the diff, let me know if this one doesn't work either or there's something else you'd like me to try. Thanks! For what it's worth it's building just fine and passing the clang-tools tests on my end.

[aaron.ballman]

In D72241#2461921, @ffrankies wrote:

@aaron.ballman hmm, that is strange. I've rebased the patch and updated the diff, let me know if this one doesn't work either or there's something else you'd like me to try. Thanks! For what it's worth it's building just fine and passing the clang-tools tests on my end.

This time everything applied cleanly and worked fine -- not certain what was up before. I've commit on your behalf in e69e551e0e5fddffb6479da6a2998457104ba9e6, thank you for the new check!

[njames93]

In D72241#2461921, @ffrankies wrote:

@aaron.ballman hmm, that is strange. I've rebased the patch and updated the diff, let me know if this one doesn't work either or there's something else you'd like me to try. Thanks! For what it's worth it's building just fine and passing the clang-tools tests on my end.

FWIW This version I was able to pull just fine and run the clang tidy tests without a hitch.

NVM, aaron has already sorted it

[ffrankies]

@aaron.ballman @njames93 thanks for testing this out, the feedback, and the commit!

Could you also take a look at D72235 unroll loops check ? It's been updated and is ready for review.

clang-tools-extra/clang-tidy/CMakeLists.txt

target_link_libraries(clangTidy PRIVATE
     clangStaticAnalyzerFrontend
   )
 endif()
 
 # Checks.
 # 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)
 add_subdirectory(cppcoreguidelines)
 add_subdirectory(darwin)
 add_subdirectory(fuchsia)
 add_subdirectory(google)
 add_subdirectory(hicpp)
 add_subdirectory(linuxkernel)
 add_subdirectory(llvm)
 add_subdirectory(misc)
 add_subdirectory(modernize)
 if(CLANG_ENABLE_STATIC_ANALYZER)
   add_subdirectory(mpi)
 endif()
 add_subdirectory(objc)
 add_subdirectory(openmp)
 add_subdirectory(performance)
 add_subdirectory(portability)
 add_subdirectory(readability)
 add_subdirectory(zircon)
 set(ALL_CLANG_TIDY_CHECKS
   clangTidyAndroidModule
   clangTidyAbseilModule
+  clangTidyAlteraModule
   clangTidyBoostModule
   clangTidyBugproneModule
   clangTidyCERTModule
   clangTidyCppCoreGuidelinesModule
   clangTidyDarwinModule
   clangTidyFuchsiaModule
   clangTidyGoogleModule
   clangTidyHICPPModule

clang-tools-extra/clang-tidy/ClangTidyForceLinker.h

 static int LLVM_ATTRIBUTE_UNUSED CERTModuleAnchorDestination =
     CERTModuleAnchorSource;
 
 // This anchor is used to force the linker to link the AbseilModule.
 extern volatile int AbseilModuleAnchorSource;
 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 =
     BoostModuleAnchorSource;
 
 // This anchor is used to force the linker to link the BugproneModule.
 extern volatile int BugproneModuleAnchorSource;
 static int LLVM_ATTRIBUTE_UNUSED BugproneModuleAnchorDestination =

clang-tools-extra/clang-tidy/altera/AlteraTidyModule.cpp

+//===--- 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

clang-tools-extra/clang-tidy/altera/CMakeLists.txt

+set(LLVM_LINK_COMPONENTS support)
+
+add_clang_library(clangTidyAlteraModule
+  AlteraTidyModule.cpp
+  SingleWorkItemBarrierCheck.cpp
+
+  LINK_LIBS
+  clangAnalysis
+  clangAST
+  clangASTMatchers
+  clangBasic
+  clangLex
+  clangTidy
+  clangTidyUtils
+  )

clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.h

+//===--- 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

clang-tools-extra/clang-tidy/altera/SingleWorkItemBarrierCheck.cpp

+//===--- 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

clang-tools-extra/docs/ReleaseNotes.rst

 Improvements to clang-rename
 ----------------------------
 
 The improvements are...
 
 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.
 
   Finds cases where an integer expression is added to or subtracted from the
   result of a memory allocation function (``malloc()``, ``calloc()``,
   ``realloc()``, ``alloca()``) instead of its argument.
 
 - New :doc:`bugprone-reserved-identifier

clang-tools-extra/docs/clang-tidy/checks/altera-single-work-item-barrier.rst

+.. 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).

clang-tools-extra/docs/clang-tidy/checks/list.rst

.. csv-table::
    `abseil-no-internal-dependencies <abseil-no-internal-dependencies.html>`_,
    `abseil-no-namespace <abseil-no-namespace.html>`_,
    `abseil-redundant-strcat-calls <abseil-redundant-strcat-calls.html>`_, "Yes"
    `abseil-str-cat-append <abseil-str-cat-append.html>`_, "Yes"
    `abseil-string-find-startswith <abseil-string-find-startswith.html>`_, "Yes"
    `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"
    `android-cloexec-dup <android-cloexec-dup.html>`_, "Yes"
    `android-cloexec-epoll-create <android-cloexec-epoll-create.html>`_,
    `android-cloexec-epoll-create1 <android-cloexec-epoll-create1.html>`_,
    `android-cloexec-fopen <android-cloexec-fopen.html>`_,
    `android-cloexec-inotify-init <android-cloexec-inotify-init.html>`_,

clang-tools-extra/docs/clang-tidy/index.rst

 .. _checks-groups-table:
 
 There are currently the following groups of checks:
 
 ====================== =========================================================
 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.
 ``cert-``              Checks related to CERT Secure Coding Guidelines.
 ``cppcoreguidelines-`` Checks related to C++ Core Guidelines.
 ``clang-analyzer-``    Clang Static Analyzer checks.
 ``fuchsia-``           Checks related to Fuchsia coding conventions.
 ``google-``            Checks related to Google coding conventions.

clang-tools-extra/test/clang-tidy/checkers/altera-single-work-item-barrier.cpp

+// 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