[SYCL][CUDA] Remove unnecessary memfence (#1935)

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>

Author: bjoernknafla

libclc/ptx-nvidiacl/libspirv/synchronization/barrier.cl

@@ -17,5 +17,4 @@ _CLC_OVERLOAD _CLC_DEF _CLC_CONVERGENT void
 __spirv_ControlBarrier(unsigned int scope, unsigned int memory,
                        unsigned int semantics) {
   __syncthreads();
-  __spirv_MemoryBarrier(memory, semantics);
 }