diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.cpp b/SYCL/AtomicRef/atomic_fence_capabilities.cpp new file mode 100644 index 0000000000..74e8617e40 --- /dev/null +++ b/SYCL/AtomicRef/atomic_fence_capabilities.cpp @@ -0,0 +1,116 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +// NOTE: General tests for atomic fence capabilities. + +#include "atomic_fence_capabilities.h" +#include +#include +using namespace sycl; + +void checkFenceBehaviour(memory_order order, memory_scope scope) { + auto q = queue(); + // Both read and write being release or acquire is wrong. In case order is + // release or acquire we need read to be acquire and write to be release. + // If we flip both acquire and release, we will be checking the same case + // (read == acquire, write == release) twice, so we just skip one case and + // flip for the other. + if (order == memory_order::acquire) + return; + memory_order order_read = order; + memory_order order_write = order; + if (order == memory_order::release) { + order_read = memory_order::acquire; + } + + // Count of retries in the check cycle + constexpr size_t RETRY_COUNT = 256; + constexpr int expected_val = 42; + + bool res = true; + int sync = 0; + int data = 0; + int value = expected_val; + + // These global_range and local_range values provide a check in one group + // when test_type = single_group, and between four groups when + // test_type = between_groups + range<1> global_range(2); + range<1> local_range(2); + + { + buffer res_buf(&res, range<1>(1)); + buffer sync_buffer(&sync, range<1>(1)); + buffer data_buffer(&data, range<1>(1)); + q.submit([&](handler &cgh) { + auto res_acc = res_buf.template get_access(cgh); + auto sync_flag_acc = + sync_buffer.template get_access(cgh); + auto data_acc = + data_buffer.template get_access(cgh); + cgh.parallel_for( + nd_range<1>(global_range, local_range), [=](nd_item<1> nditem) { + atomic_ref + sync_flag(sync_flag_acc[0]); + int *data = &data_acc[0]; + // Only one nditem should perform non-atomic write. + // All other nditems should perform non-atomic + // reads + if (nditem.get_global_linear_id() == 0) { + // Non-atomic write to data + *data = value; + // Used atomic_fence to guarantee the order + // instructions execution + atomic_fence(order_write, scope); + // Used atomic sync flag to avoid data racing + sync_flag = 1; + } else { + bool write_happened = false; + for (size_t i = 0; i < RETRY_COUNT; i++) { + if (sync_flag == 1) { + write_happened = true; + break; + } + } + atomic_fence(order_read, scope); + // After the fence safe non-atomic reading + if (write_happened) { + // Non-atomic read of data + if (*data != value) + res_acc[0] = false; + } + } + }); + }); + } + assert(res); +} + +int main() { + queue q; + + std::vector supported_memory_orders = + q.get_device().get_info(); + + // Relaxed, acquire, release and acq_rel memory order must be supported. + assert(is_supported_order(supported_memory_orders, memory_order::relaxed)); + assert(is_supported_order(supported_memory_orders, memory_order::acquire)); + assert(is_supported_order(supported_memory_orders, memory_order::release)); + assert(is_supported_order(supported_memory_orders, memory_order::acq_rel)); + + std::vector supported_memory_scopes = + q.get_device().get_info(); + + // Work_group, sub_group and work_item memory order must be supported. + assert(is_supported_scope(supported_memory_scopes, memory_scope::work_item)); + assert(is_supported_scope(supported_memory_scopes, memory_scope::sub_group)); + assert(is_supported_scope(supported_memory_scopes, memory_scope::work_group)); + + for (auto order : supported_memory_orders) + for (auto scope : supported_memory_scopes) + checkFenceBehaviour(order, scope); + + std::cout << "Test passed." << std::endl; +} diff --git a/SYCL/AtomicRef/atomic_fence_capabilities.h b/SYCL/AtomicRef/atomic_fence_capabilities.h new file mode 100644 index 0000000000..6688ecc320 --- /dev/null +++ b/SYCL/AtomicRef/atomic_fence_capabilities.h @@ -0,0 +1,18 @@ +#pragma once + +#include +#include + +using namespace sycl; + +bool is_supported_order(const std::vector &capabilities, + memory_order mem_order) { + return std::find(capabilities.begin(), capabilities.end(), mem_order) != + capabilities.end(); +} + +bool is_supported_scope(const std::vector &capabilities, + memory_scope mem_scope) { + return std::find(capabilities.begin(), capabilities.end(), mem_scope) != + capabilities.end(); +}