Skip to content
This repository was archived by the owner on Mar 28, 2023. It is now read-only.

[ESIMD] Fix and improve slm_split_barrier.cpp #542

Merged
merged 3 commits into from
Nov 1, 2021
Merged

[ESIMD] Fix and improve slm_split_barrier.cpp #542

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
0cb8a25
[ESIMD] Fix and improve slm_split_barrier.cpp
kbobrovs Oct 30, 2021
9249c92
Merge branch 'intel' into slm_split_barrier
kbobrovs Oct 30, 2021
6da7bd7
Update SYCL/ESIMD/slm_split_barrier.cpp
kbobrovs Oct 30, 2021
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
14 changes: 8 additions & 6 deletions SYCL/ESIMD/slm_split_barrier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ void load_to_slm(uint grpSize, uint localId, uint slmOffset, char *addr,
vOffsets += (grpSize * 256);
}

esimd::fence(ESIMD_GLOBAL_COHERENT_FENCE);
esimd::fence(fence_mask::global_coherent_fence);
esimd::sbarrier(split_barrier_action::signal);
esimd::sbarrier(split_barrier_action::wait);
}
Expand All @@ -81,10 +81,8 @@ int main(void) {

auto dev = q.get_device();
std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
auto ctxt = q.get_context();
// TODO: release memory in the end of the test
uint *A = static_cast<uint *>(malloc_shared(Size * sizeof(uint), dev, ctxt));
uint *B = static_cast<uint *>(malloc_shared(Size * sizeof(uint), dev, ctxt));
uint *A = malloc_shared<uint>(Size, q);
uint *B = malloc_shared<uint>(Size, q);

// Checking with specific inputs
for (int i = 0; i < NUM_THREADS; i++) {
Expand Down Expand Up @@ -134,7 +132,9 @@ int main(void) {
e.wait();
} catch (cl::sycl::exception const &e) {
std::cout << "SYCL exception caught: " << e.what() << '\n';
return e.get_cl_code();
sycl::free(A, q);
sycl::free(B, q);
return e.code().value();
}

std::cout << "result" << std::endl;
Expand All @@ -160,6 +160,8 @@ int main(void) {
}
std::cout << std::endl;
}
sycl::free(A, q);
sycl::free(B, q);

std::cout << (result < 0 ? "FAILED\n" : "Passed\n");
return result < 0 ? 1 : 0;
Expand Down