Skip to content

Commit

Permalink
Fix atomic_ref_stress / atomicity_with_host_code
Browse files Browse the repository at this point in the history 
atomicity_with_host_code requires a memory_scope argument (always
`system`). atomic_ref<float> has no operator++ and must use fetch_add()
instead.
  • Loading branch information
@fknorr
fknorr committed Jan 8, 2025
1 parent 9dc609d commit e28ab11
Showing 1 changed file with 33 additions and 28 deletions.
61 changes: 33 additions & 28 deletions tests/atomic_ref_stress/atomic_ref_stress_common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -42,10 +42,10 @@ class atomicity_device_scope {
AddressSpaceT::value;

public:
void operator()(const std::string &type_name,
const std::string &memory_order_name,
const std::string &memory_scope_name,
const std::string &address_space_name) {
void operator()(const std::string& type_name,
const std::string& memory_order_name,
const std::string& memory_scope_name,
const std::string& address_space_name) {
INFO(atomic_ref::tests::common::get_section_name(
type_name, memory_order_name, memory_scope_name, address_space_name,
"atomicity_device_scope"));
Expand All @@ -59,7 +59,7 @@ class atomicity_device_scope {
max_size);
{
sycl::buffer buf{&val, {1}};
queue.submit([&](sycl::handler &cgh) {
queue.submit([&](sycl::handler& cgh) {
sycl::accessor acc{buf, cgh};
cgh.parallel_for({size}, [=](auto i) {
sycl::atomic_ref<T, MemoryOrder, MemoryScope, AddressSpace> a_r{
Expand All @@ -86,10 +86,10 @@ class atomicity_work_group_scope {
AddressSpaceT::value;

public:
void operator()(const std::string &type_name,
const std::string &memory_order_name,
const std::string &memory_scope_name,
const std::string &address_space_name) {
void operator()(const std::string& type_name,
const std::string& memory_order_name,
const std::string& memory_scope_name,
const std::string& address_space_name) {
INFO(atomic_ref::tests::common::get_section_name(
type_name, memory_order_name, memory_scope_name, address_space_name,
"atomicity_work_group_scope"));
Expand All @@ -106,7 +106,7 @@ class atomicity_work_group_scope {
{
sycl::buffer buf{vals.data(), {group_range}};
queue
.submit([&](sycl::handler &cgh) {
.submit([&](sycl::handler& cgh) {
sycl::accessor acc{buf, cgh};
sycl::local_accessor<T> lacc{{1}, cgh};
cgh.parallel_for(
Expand Down Expand Up @@ -142,10 +142,10 @@ class aquire_release {
AddressSpaceT::value;

public:
void operator()(const std::string &type_name,
const std::string &memory_order_name,
const std::string &memory_scope_name,
const std::string &address_space_name) {
void operator()(const std::string& type_name,
const std::string& memory_order_name,
const std::string& memory_scope_name,
const std::string& address_space_name) {
INFO(atomic_ref::tests::common::get_section_name(
type_name, memory_order_name, memory_scope_name, address_space_name,
"aquire_release"));
Expand All @@ -160,7 +160,7 @@ class aquire_release {
res.fill(false);
{
sycl::buffer buf{res.data(), {global_range / local_range}};
queue.submit([&](sycl::handler &cgh) {
queue.submit([&](sycl::handler& cgh) {
sycl::accessor res_acc{buf, cgh};
sycl::local_accessor<T, 0> x{cgh};
sycl::local_accessor<T, 0> y{cgh};
Expand Down Expand Up @@ -202,10 +202,10 @@ class ordering {
AddressSpaceT::value;

public:
void operator()(const std::string &type_name,
const std::string &memory_order_name,
const std::string &memory_scope_name,
const std::string &address_space_name) {
void operator()(const std::string& type_name,
const std::string& memory_order_name,
const std::string& memory_scope_name,
const std::string& address_space_name) {
INFO(atomic_ref::tests::common::get_section_name(
type_name, memory_order_name, memory_scope_name, address_space_name,
"ordering"));
Expand All @@ -225,7 +225,7 @@ class ordering {
res.fill(false);
{
sycl::buffer buf{res.data(), {global_range}};
queue.submit([&](sycl::handler &cgh) {
queue.submit([&](sycl::handler& cgh) {
sycl::accessor res_acc{buf, cgh};
sycl::local_accessor<T, 0> local_acc{cgh};
sycl::local_accessor<bool> arr_acc{{local_range}, cgh};
Expand All @@ -246,9 +246,11 @@ class ordering {
}
};
#ifdef __cpp_lib_atomic_ref
template <typename T, typename MemoryOrderT, typename AddressSpaceT>
template <typename T, typename MemoryOrderT, typename MemoryScopeT,
typename AddressSpaceT>
class atomicity_with_host_code {
static constexpr sycl::memory_order MemoryOrder = MemoryOrderT::value;
static constexpr sycl::memory_scope MemoryScope = MemoryScopeT::value;
static constexpr sycl::access::address_space AddressSpace =
AddressSpaceT::value;

Expand Down Expand Up @@ -278,18 +280,18 @@ class atomicity_with_host_code {
sycl::atomic_ref<T, MemoryOrder, sycl::memory_scope::system,
AddressSpace>
a_dev{*pval};
a_dev++;
a_dev.fetch_add(1);
});
});
for (int i = 0; i < count; i++) a_host++;
for (int i = 0; i < count; i++) a_host.fetch_add(1);
event.wait();
CHECK(*pval == size + count);
}
};
#endif
template <typename T>
struct run_atomicity_device_scope {
void operator()(const std::string &type_name) {
void operator()(const std::string& type_name) {
const auto memory_orders =
value_pack<sycl::memory_order, sycl::memory_order::relaxed,
sycl::memory_order::acq_rel,
Expand All @@ -308,7 +310,7 @@ struct run_atomicity_device_scope {

template <typename T>
struct run_atomicity_work_group_scope {
void operator()(const std::string &type_name) {
void operator()(const std::string& type_name) {
const auto memory_orders =
value_pack<sycl::memory_order, sycl::memory_order::relaxed,
sycl::memory_order::acq_rel,
Expand All @@ -328,7 +330,7 @@ struct run_atomicity_work_group_scope {

template <typename T>
struct run_aquire_release {
void operator()(const std::string &type_name) {
void operator()(const std::string& type_name) {
const auto memory_orders =
value_pack<sycl::memory_order, sycl::memory_order::acq_rel,
sycl::memory_order::seq_cst>::generate_named();
Expand All @@ -347,7 +349,7 @@ struct run_aquire_release {

template <typename T>
struct run_ordering {
void operator()(const std::string &type_name) {
void operator()(const std::string& type_name) {
const auto memory_orders =
value_pack<sycl::memory_order, sycl::memory_order::release,
sycl::memory_order::seq_cst>::generate_named();
Expand All @@ -371,12 +373,15 @@ struct run_atomicity_with_host_code {
value_pack<sycl::memory_order, sycl::memory_order::relaxed,
sycl::memory_order::acq_rel,
sycl::memory_order::seq_cst>::generate_named();
const auto memory_scopes =
value_pack<sycl::memory_scope,
sycl::memory_scope::system>::generate_named();
const auto address_spaces = value_pack<
sycl::access::address_space, sycl::access::address_space::global_space,
sycl::access::address_space::generic_space>::generate_named();

for_all_combinations<atomicity_with_host_code, T>(
memory_orders, address_spaces, type_name);
memory_orders, memory_scopes, address_spaces, type_name);
}
};
#endif
Expand Down

0 comments on commit e28ab11

Please sign in to comment.