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