1.. title:: clang-tidy - altera-single-work-item-barrier
2
3altera-single-work-item-barrier
4===============================
5
6Finds OpenCL kernel functions that call a barrier function but do not call
7an ID function (``get_local_id``, ``get_local_id``, ``get_group_id``, or
8``get_local_linear_id``).
9
10These kernels may be viable single work-item kernels, but will be forced to
11execute as NDRange kernels if using a newer version of the Altera Offline
12Compiler (>= v17.01).
13
14If using an older version of the Altera Offline Compiler, these kernel
15functions will be treated as single work-item kernels, which could be
16inefficient or lead to errors if NDRange semantics were intended.
17
18Based on the `Altera SDK for OpenCL: Best Practices Guide
19<https://www.altera.com/en_US/pdfs/literature/hb/opencl-sdk/aocl_optimization_guide.pdf>`_.
20
21Examples:
22
23.. code-block:: c++
24
25  // error: function calls barrier but does not call an ID function.
26  void __kernel barrier_no_id(__global int * foo, int size) {
27    for (int i = 0; i < 100; i++) {
28      foo[i] += 5;
29    }
30    barrier(CLK_GLOBAL_MEM_FENCE);
31  }
32
33  // ok: function calls barrier and an ID function.
34  void __kernel barrier_with_id(__global int * foo, int size) {
35    for (int i = 0; i < 100; i++) {
36      int tid = get_global_id(0);
37      foo[tid] += 5;
38    }
39    barrier(CLK_GLOBAL_MEM_FENCE);
40  }
41
42  // ok with AOC Version 17.01: the reqd_work_group_size turns this into
43  // an NDRange.
44  __attribute__((reqd_work_group_size(2,2,2)))
45  void __kernel barrier_with_id(__global int * foo, int size) {
46    for (int i = 0; i < 100; i++) {
47      foo[tid] += 5;
48    }
49    barrier(CLK_GLOBAL_MEM_FENCE);
50  }
51
52Options
53-------
54
55.. option:: AOCVersion
56
57   Defines the version of the Altera Offline Compiler. Defaults to ``1600``
58   (corresponding to version 16.00).
59