[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/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 "KernelNameRestrictionCheck.h"
+#include "SingleWorkItemBarrierCheck.h"
 #include "StructPackAlignCheck.h"
 
 using namespace clang::ast_matchers;
 
 namespace clang {
 namespace tidy {
 namespace altera {
 
 class AlteraModule : public ClangTidyModule {
 public:
   void addCheckFactories(ClangTidyCheckFactories &CheckFactories) override {
     CheckFactories.registerCheck<KernelNameRestrictionCheck>(
         "altera-kernel-name-restriction");
+    CheckFactories.registerCheck<SingleWorkItemBarrierCheck>(
+        "altera-single-work-item-barrier");
     CheckFactories.registerCheck<StructPackAlignCheck>(
         "altera-struct-pack-align");
   }
 };
 
 } // namespace altera
 
 // Register the AlteraTidyModule using this statically initialized variable.

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

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

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 "../ClangTidyCheck.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 an ID function.
+  // 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 four 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(hasAnyName(
+                                             "barrier", "work_group_barrier"))))
+                                    .bind("barrier")),
+              // But do not call an ID function.
+              unless(hasDescendant(callExpr(callee(functionDecl(
+                  hasAnyName("get_global_id", "get_local_id", "get_group_id",
+                             "get_local_linear_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 (AOCVersion < 1701) {
+    // get_group_id and get_local_linear_id were added at/after v17.01
+    diag(MatchedDecl->getLocation(),
+         "kernel function %0 does not call 'get_global_id' or 'get_local_id' "
+         "and will be treated as a single work-item")
+        << MatchedDecl;
+    diag(MatchedBarrier->getBeginLoc(),
+         "barrier call is in a single work-item and may error out",
+         DiagnosticIDs::Note);
+  } else {
+    // 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;
+    if (MatchedDecl->hasAttr<ReqdWorkGroupSizeAttr>()) {
+      const auto *Attribute = MatchedDecl->getAttr<ReqdWorkGroupSizeAttr>();
+      if (Attribute->getXDim() > 1 || Attribute->getYDim() > 1 ||
+          Attribute->getZDim() > 1)
+        IsNDRange = true;
+    }
+    if (IsNDRange) // No warning if kernel is treated as an NDRange.
+      return;
+    diag(MatchedDecl->getLocation(),
+         "kernel function %0 does not call an ID function and may be a viable "
+         "single work-item, but will be forced to execute as an NDRange")
+        << MatchedDecl;
+    diag(MatchedBarrier->getBeginLoc(),
+         "barrier call will force NDRange execution; if single work-item "
+         "semantics are desired a mem_fence may be more efficient",
+         DiagnosticIDs::Note);
+  }
+}
+
+void SingleWorkItemBarrierCheck::storeOptions(
+    ClangTidyOptions::OptionMap &Opts) {
+  Options.store(Opts, "AOCVersion", AOCVersion);
+}
+
+} // namespace altera
+} // namespace tidy
+} // namespace clang

clang-tools-extra/docs/ReleaseNotes.rst

 ^^^^^^^^^^
 
 - New :doc:`altera-kernel-name-restriction
   <clang-tidy/checks/altera-kernel-name-restriction>` check.
 
   Finds kernel files and include directives whose filename is `kernel.cl`,
   `Verilog.cl`, or `VHDL.cl`.
 
+- 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:`altera-struct-pack-align
   <clang-tidy/checks/altera-struct-pack-align>` check.
 
   Finds structs that are inefficiently packed or aligned, and recommends
   packing and/or aligning of said structs as needed.
 
 - New :doc:`cppcoreguidelines-prefer-member-initializer
   <clang-tidy/checks/cppcoreguidelines-prefer-member-initializer>` check.

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 (``get_local_id``, ``get_local_id``, ``get_group_id``, or
+``get_local_linear_id``).
+
+These kernels may be viable single work-item kernels, but will be forced to
+execute as NDRange kernels if using a newer version of the Altera Offline
+Compiler (>= v17.01).
+
+If using an older version of the Altera Offline Compiler, these kernel
+functions will be treated as single work-item kernels, which could be
+inefficient or lead to errors if NDRange semantics were intended.
+
+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-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-string-find-str-contains <abseil-string-find-str-contains.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-kernel-name-restriction <altera-kernel-name-restriction.html>`_,
+   `altera-single-work-item-barrier <altera-single-work-item-barrier.html>`_, "Yes"
    `altera-struct-pack-align <altera-struct-pack-align.html>`_,
    `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>`_,

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

+// RUN: %check_clang_tidy -check-suffix=OLDCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL1.2 -c --include opencl-c.h -DOLDCLOLDAOC
+// RUN: %check_clang_tidy -check-suffix=NEWCLOLDAOC %s altera-single-work-item-barrier %t -- -header-filter=.* "--" -cl-std=CL2.0 -c --include opencl-c.h -DNEWCLOLDAOC
+// RUN: %check_clang_tidy -check-suffix=OLDCLNEWAOC %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 -DOLDCLNEWAOC
+// RUN: %check_clang_tidy -check-suffix=NEWCLNEWAOC %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 -DNEWCLNEWAOC
+
+#ifdef OLDCLOLDAOC  // OpenCL 1.2 Altera Offline Compiler < 17.1
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+  // CHECK-MESSAGES-OLDCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier]
+  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-OLDCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out
+  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 NEWCLOLDAOC  // OpenCL 2.0 Altera Offline Compiler < 17.1
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+  // CHECK-MESSAGES-NEWCLOLDAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call 'get_global_id' or 'get_local_id' and will be treated as a single work-item [altera-single-work-item-barrier]
+  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-NEWCLOLDAOC: :[[@LINE-1]]:3: note: barrier call is in a single work-item and may error out
+  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 OLDCLNEWAOC  // OpenCL 1.2 Altera Offline Compiler >= 17.1
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+  // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
+  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-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
+  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) {
+  // CHECK-MESSAGES-OLDCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
+  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-OLDCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
+  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 NEWCLNEWAOC  // OpenCL 2.0 Altera Offline Compiler >= 17.1
+void __kernel error_barrier_no_id(__global int * foo, int size) {
+  // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
+  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-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
+  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) {
+  // CHECK-MESSAGES-NEWCLNEWAOC: :[[@LINE-1]]:15: warning: kernel function 'error_barrier_no_id_work_group_size' does not call an ID function and may be a viable single work-item, but will be forced to execute as an NDRange [altera-single-work-item-barrier]
+  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-NEWCLNEWAOC: :[[@LINE-1]]:3: note: barrier call will force NDRange execution; if single work-item semantics are desired a mem_fence may be more efficient
+  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