The SPIR spec states that all OpenCL built-in functions should be
overloadable and mangled, to ensure consistency.
Add the overload attribute to functions which were missing them:
work dimensions, memory barriers and fences, and events.
Reviewed By: tstellar, jenatali
Differential Revision: https://reviews.llvm.org/D82078
This generates a "bar.sync 0” instruction, which not only causes the
threads to wait, but does acts as a memory fence, as required by
OpenCL. The fence does not differentiate between local and global
memory. Unfortunately, there is no similar instruction which does
not include a memory fence. Hence, we cannot optimize the case
where neither CLK_LOCAL_MEM_FENCE nor CLK_GLOBAL_MEM_FENCE is
passed.
llvm-svn: 315228
Daniel Stone