Deprecate and remove `nd_item::barrier`.

[bader]

SYCL 2020 removed this function completely, in favor of sycl::group_barrier.

Originally posted by @Pennycook in #12236 (comment)

[jinz2014]

Could you please explain why nd_item::barrier needs to be replaced by sycl::group_barrier ? I am not a customer mentioned in the comments, but I think many users' SYCL programs still call item.barrier().

[bader]

There is no nd_item::barrier in SYCL 2020 specification.

[kris-rowe]

Removing the existing nd_item::barrier will have a non-trivial performance impact on applications that use this with the access::fence_space::local_space flag. In this case, __spirv_ControlBarrier is called using the WorkgroupMemory mask returned from detail::getSPIRVMemorySemanticsMask.

In the SYCL 2020 spec group_barrier doesn't accept an equivalent to the local_space flag. The current implementation uses SPIRV CrossWorkgroupMemory memory semantics for all calls to group_barrier.

[bader]

@Pennycook, is there any public document addressing performance concerns of moving from nd_item::barrier to group_barrier?

[Pennycook]

@Pennycook, is there any public document addressing performance concerns of moving from nd_item::barrier to group_barrier?

Not that I'm aware of. Any performance impact of the change is both implementation- and hardware-specific.

[jinz2014]

@bader I may confuse item.barrier() with nd_item::barrier()

[MikeDvorskiy]

@bader
Alexey, which version macro (see below) will be changed and what is a new value for the changed macro after the PR will be merged?
__LIBSYCL_MAJOR_VERSION
__LIBSYCL_MINOR_VERSION
__LIBSYCL_PATCH_VERSION

[bader]

@MikeDvorskiy, as this is API breaking change, all of them will change. I can't say what will be new values at this moment as there is PR with this change yet (as far as I know).
Tagging folks who can follow-up on that issue: @sergey-semenov, @AlexeySachkov, @steffenlarsen.

[al42and]

@Pennycook, is there any public document addressing performance concerns of moving from nd_item::barrier to group_barrier?

A quick test on GROMACS shows that replacing a bunch of itemIdx.barrier(fence_space::local_space) with sycl::group_barrier(itemIdx.get_group()) slows down one of the kernels (PmeGatherKernel<4, true, true, 1, false, (ThreadsPerAtom)0, 16>) nearly 1.8x (from 270µs to 470µs for 768k Grappa system) on PVC 1100.

This kernel is not super-optimized (it is responsible for ~4% of the runtime, at least when using nd_item::barrier) and perhaps this slowdown can be worked around. But still, the performance concerns are real.

Some earlier microbenchmarks: #7069 (comment)

[Pennycook]

@bader, @al42and: The performance issue was recently root-caused to a bug in IGC. With the fix from intel/intel-graphics-compiler@ed639f6, you should still expect some small overhead from the difference in fence semantics, but performance should be significantly better than it was before.