[SYCL][ESIMD] Fix the order of parameters for the lsc_slm_atomic_update that caused test failures (#10963)

Author: fineg74

sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp

@@ -3077,8 +3077,8 @@ lsc_slm_atomic_update(__ESIMD_NS::simd<uint32_t, N> offsets,
 /// @tparam N is the number of channels (platform dependent).
 /// @tparam DS is the data size.
 /// @param offsets is the zero-based offsets.
-/// @param src0 is the first atomic operand.
-/// @param src1 is the second atomic operand.
+/// @param src0 is the first atomic operand (expected value).
+/// @param src1 is the second atomic operand (new value).
 /// @param pred is predicates.
 ///
 /// @return A vector of the old values at the memory locations before the
@@ -3539,8 +3539,8 @@ lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets,
 /// @tparam AccessorTy is the \ref sycl::accessor type.
 /// @param acc is the SYCL accessor.
 /// @param offsets is the zero-based offsets.
-/// @param src0 is the first atomic operand.
-/// @param src1 is the second atomic operand.
+/// @param src0 is the first atomic operand (expected value).
+/// @param src1 is the second atomic operand (new value).
 /// @param pred is predicates.
 ///
 /// @return A vector of the old values at the memory locations before the
@@ -3598,8 +3598,8 @@ lsc_atomic_update(AccessorTy acc, __ESIMD_NS::simd<Toffset, N> offsets,
 /// @tparam AccessorTy is the \ref sycl::accessor type.
 /// @param acc is the SYCL accessor.
 /// @param offsets is the zero-based offsets.
-/// @param src0 is the first atomic operand.
-/// @param src1 is the second atomic operand.
+/// @param src0 is the first atomic operand (expected value).
+/// @param src1 is the second atomic operand (new value).
 /// @param pred is predicates.
 ///
 /// @return A vector of the old values at the memory locations before the

sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_cmpxchg.cpp

@@ -8,8 +8,6 @@
 // This test checks LSC SLM atomic operations.
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu-intel-pvc
-// XFAIL: gpu-intel-pvc
-// TODO: reenable the test once driver issue with cmpxchg is resolved
 // TODO: esimd_emulator fails due to random timeouts (_XFAIL_: esimd_emulator)
 // TODO: esimd_emulator doesn't support xchg operation
 // UNSUPPORTED: esimd_emulator

sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp

@@ -149,10 +149,10 @@ bool test(queue q) {
                 // arg0 and arg1 must provide values which guarantee the loop
                 // is not endless:
                 for (auto old_val = lsc_slm_atomic_update<op, T>(
-                         offsets, new_val, exp_val, m);
+                         offsets, exp_val, new_val, m);
                      any(old_val < exp_val, !m);
-                     old_val = lsc_slm_atomic_update<op, T>(offsets, new_val,
-                                                            exp_val, m))
+                     old_val = lsc_slm_atomic_update<op, T>(offsets, exp_val,
+                                                            new_val, m))
                   ;
               }
             }