Further comments on this.

We only use this branch for a device execution space instance which isn't able to access the memory in question anyway. The only problem could arise if we have another host execution space (different from the one used here) accesses the memory from a different thread but we don't promise anything in that case anyway. If we happen to use the same execution space instance, we enqueue correctly.

Originally posted by @masterleinad in #7245 (comment)

If I understand correctly, the following case may cause a problem

using ExecSpace = Kokkos::DefaultExecutionSpace;
auto instances =
      Kokkos::Experimental::partition_space(ExecSpace, 1, 1);

std::thread t1([&]() {
  DeepCopy<HostSpace, HostSpace, ExecSpace>(instances[0], dst, src, size);
});
std::thread t2([&]() {
  DeepCopy<HostSpace, HostSpace, ExecSpace>(instances[1], dst, src2, size);
});
t1.join();
t2.join();

To avoid a race condition, we need to make sure that the internal operations are not executed concurrently.
Do I understand the issue correctly? @crtrott @dalg24 @masterleinad

What behavior would you expect? Without coordinating the two (separate) execution space instances in the code snippet, there is no way to know in when dst is accessed by the two threads and so you would always have a race condition.

The case we are trying to support is for both threads to use the same execution space instance. In this case, we want to guarantee that the deep_copies are enqueued in some order so you wouldn't know if you get src or src2 but at least there is no race condition.

The relevant case here is for the execution space instance not to be able to access host memory (since we enqueue otherwise anyway). The current implementation would fence the execution space instance first and then enqueue the deep copy on the default host execution space. The order of the fences doesn't matter and doesn't create a race condition.

What behavior would you expect? Without coordinating the two (separate) execution space instances in the code snippet, there is no way to know in when dst is accessed by the two threads and so you would always have a race condition.

The case we are trying to support is for both threads to use the same execution space instance. In this case, we want to guarantee that the deep_copies are enqueued in some order so you wouldn't know if you get src or src2 but at least there is no race condition.

Sure, so we consider something like

using ExecSpace = Kokkos::DefaultExecutionSpace;

std::thread t1([&]() {
  DeepCopy<HostSpace, HostSpace, ExecSpace>(ExecSpace(), dst, src, size);
});
std::thread t2([&]() {
  DeepCopy<HostSpace, HostSpace, ExecSpace>(ExecSpace(), dst, src2, size);
});
t1.join();
t2.join();

The relevant case here is for the execution space instance not to be able to access host memory (since we enqueue otherwise anyway). The current implementation would fence the execution space instance first and then enqueue the deep copy on the default host execution space. The order of the fences doesn't matter and doesn't create a race condition.

In order to enqueue, we need to do the following procedure in DeepCopy?

1. std::scoped_lock the mutex behind the execspace.
2. cudaStreamSynchronize(stream) (or equivalent)
3. do hostspace_parallel_deepcopy_async(dst, src, n);

This way we can avoid a race condition between t1 and t2?
I am not sure which mutex can be used for this purpose.

I'm trying to argue that the current implementation is correct and we do what we promise. When I see "enqueue", I mean using the respective mechanism for a backend to make sure that kernels submitted to the same execution space instance don't run concurrently and, if submitted from the same thread, in the order of submission.

I'm trying to argue that the current implementation is correct and we do what we promise. When I see "enqueue", I mean using the respective mechanism for a backend to make sure that kernels submitted to the same execution space instance don't run concurrently and, if submitted from the same thread, in the order of submission.

If we submit tasks from different threads, I guess the internal hostspace_parallel_deepcopy_async can potentially run concurrently. Do you mean that nothing is promised if we submit tasks from different threads?

No, we need to make sure that multiple calls to hostspace_parallel_deepcopy_async using the same execution space instance don't run concurrently but enqueue correctly on that execution space instance.

In the current implementation, hostspace_parallel_deepcopy_async runs on Kokkos::DefaultHostExecutionSpace. Thus, I feel hostspace_parallel_deepcopy_async can potentially run concurrently if this is called from different threads and the Kokkos::DefaultHostExecutionSpace is asynchronous like HPX.

hostspace_parallel_deepcopy_async

Yes, we need to fix the implementation of hostspace_parallel_deepcopy_async but the call site is OK.

Sure. Then, I may propose

1. Keep the call site DeepCopy as is
2. Fix the overload of hostspace_parallel_deepcopy_async
3. Introduce a lock inside to make sure that hostspace_parallel_deepcopy_async does not run concurrently.

Kokkos::DefaultHostExecutionSpace exec;
auto *internal_instance = exec.impl_internal_space_instance();

std::lock_guard<std::mutex> lock(internal_instance->m_instance_mutex);
internal_instance->fence(
    "Kokkos::Impl::hostspace_parallel_deepcopy_async: fence before copy");
hostspace_parallel_deepcopy_async(exec, dst, src, n);
internal_instance->fence(
    "Kokkos::Impl::hostspace_parallel_deepcopy_async: fence after copy");

If it is good for you, I will make a PR for this fix

Sure, you would just need to make sure that all relevant execution spaces (Serial, OpenMP, HPX, Threads) have m_instance_mutex. I don't know about the latter two. Also, you can't have that mutex locked when calling a parallel construct so you shouldn't do that if the implementation just calls a single parallel construct. You would need to refactor if multiple parallel constructs are called or if it's something else touching the memory (memcpy for example) and a parallel construct.

Sure, you would just need to make sure that all relevant execution spaces (Serial, OpenMP, HPX, Threads) have m_instance_mutex.

Sure, I will have a look.

Also, you can't have that mutex locked when calling a parallel construct so you shouldn't do that if the implementation just calls a single parallel construct. You would need to refactor if multiple parallel constructs are called or if it's something else touching the memory (memcpy for example) and a parallel construct.

Internally, the procedure changes depending on the size and alignment. It may perform memcpy, parallel_for or byte based copies for non-aligned part. For asynchronous HPX backend, it always performs parallel_for, so we do not need a lock this case.

The problematic parts are

Thank you for comments. I will start working on this

After some discussions with @dalg24, I am currently trying to get a race condition deliberately to develop a good test for this issue. The problematic lines can only be executed with unaligned data, so I am using short typed view.

However, these lines in while loop are in practice not executed at the same time, since it finishes quite rapidly. It is a thread that executes this part last who will update the values. I can deliberately have a race condition by inserting sleep to line 104 for example, but it does not make sense. Is there an alternative way to force a race condition within two threads?

Have a look at our other thread-safety tests. We are doing something like

You could have one thread execute such a kernel repeatedly and the other thread would copy repeatedly. You will probably need to use more than one element to hit the desired path in the deep_copy, though.

That is a good idea. I was just trying to make a race with deep_copies from two threads, but a race condition has never been observed.