[SYCL][CUDA] Remove unnecessary memfence

[bjoernknafla]

Remove unnecessary memory fence after a CUDA memory barrier
(__syncthreads).

The emitted bar.sync 0 PTX instruction ensures that all memory
accesses of threads involved in the barrier 0 have been performed and
that no new memory accesses happen before the barrier completes.

The removed memory fence reduced performance without adding any
functionality to the barrier memory behavior.

Signed-off-by: Bjoern Knafla bjoern@codeplay.com
Co-authored-be: Victor Lomuller victor@codeplay.com

[bader]

Remove unnecessary memory fence after a CUDA memory barrier
(__syncthreads).

I have doubts that it's unnecessary.
Please, take a look at #1258.

Tagging @againull, @Naghasan.

[againull]

@bader, thanks for double checking.
I agree that it is unnecessary. And keeping or removing it doesn't affect/resolve #1258

#1258 is about the following:
sync_threads() is lowered to llvm.nvvm.barrier0 and __spirv_MemoryBarrier() is lowered to llvm.nvvm.membar. LLVM middle-end doesn't know that llvm.nvvm.barrier0 and llvm.nvvm.membar are cuda barriers, it treats them as a regular llvm intrinsic and can move memory accesses across any of these intrinsics -> barrier logic will be broken. See workaround: https://github.com/intel/llvm/pull/1334/files

[Naghasan]

Remove unnecessary memory fence after a CUDA memory barrier
(__syncthreads).

I have doubts that it's unnecessary.

it is, from the ptx doc prior memory accesses requested by this thread are performed relative to all threads participating in the barrier

Please, take a look at #1258.

as @againull pointed out, the issue is about llvm not understanding the semantic of the these instructions.

[bader]

Sorry, I forgot about the workaround implemented by #1334.

LIT failure seems to be related to #1919. I'll re-run it to verify.