Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Aug 29, 2023
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9cfd9cb
Fix the order of parameters for the lsc_slm_atomic_update that caused…
fineg74 Aug 25, 2023
6bc051d
Update function comments to clarify the meaning of the parameters
fineg74 Aug 25, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 6 additions & 6 deletions sycl/include/sycl/ext/intel/experimental/esimd/memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down
2 changes: 0 additions & 2 deletions sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_cmpxchg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
6 changes: 3 additions & 3 deletions sycl/test-e2e/ESIMD/lsc/lsc_slm_atomic_smoke.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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))
;
}
}
Expand Down