diff --git a/SYCL/AtomicRef/assignment.cpp b/SYCL/AtomicRef/assignment.cpp new file mode 100644 index 0000000000..ecb8d742f7 --- /dev/null +++ b/SYCL/AtomicRef/assignment.cpp @@ -0,0 +1,56 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %HOST_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include +#include +#include +#include +#include +using namespace sycl; +using namespace sycl::ONEAPI; + +template class assignment_kernel; + +template void assignment_test(queue q, size_t N) { + T initial = T(N); + T assignment = initial; + { + buffer assignment_buf(&assignment, 1); + q.submit([&](handler &cgh) { + auto st = + assignment_buf.template get_access(cgh); + cgh.parallel_for>(range<1>(N), [=](item<1> it) { + size_t gid = it.get_id(0); + auto atm = atomic_ref(st[0]); + atm = T(gid); + }); + }); + } + + // The initial value should have been overwritten by a work-item ID + // Atomicity isn't tested here, but support for assignment() is + assert(assignment != initial); + assert(assignment >= T(0) && assignment <= T(N - 1)); +} + +int main() { + queue q; + std::string version = q.get_device().get_info(); + + constexpr int N = 32; + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + assignment_test(q, N); + + std::cout << "Test passed." << std::endl; +}