DPCT1078

Message

Consider replacing memory_order::acq_rel with memory_order::seq_cst for correctness if strong memory order restrictions are needed.

Detailed Help

memory_order::acq_rel is a light-weight fence that is sufficient for memory synchronization in most programs. If a program needs total sequentially consistent memory order to ensure correctness, replace memory_order::acq_rel with memory_order::seq_cst.

Suggestions to Fix

Replace memory_order::acq_rel with memory_order::seq_cst if stricter memory order is needed.

For example, this original CUDA* code:

__device__ void foo_dev() {
  ...
  __threadfence();
  ...
}

results in the following migrated SYCL code:

void foo_dev() {
  ...
  /*
  DPCT1078:0: Consider replacing memory_order::acq_rel with
  memory_order::seq_cst for correctness if strong memory order restrictions are
  needed.
  */
  sycl::atomic_fence(sycl::memory_order::acq_rel, sycl::memory_scope::device);
  ...
}

which is rewritten to:

void foo_dev() {
  ...
  // Assuming strong memory order restrictions are need here
  sycl::atomic_fence(sycl::memory_order::seq_cst, sycl::memory_scope::device);
  ...
}