[BUG]: Ensure cudaMemcpy is called by thrust::copy

[gonzalobg]

We should ensure that this:

void test(double* in, double* out, size_t n) {
    thrust::copy(thrust::device, in, in + n, out);
}

calls cudaMemcpy or cudaMemcpyAsync with cudaMemcpyDefault.

Right now it does not seem to be happening. @jrhemstad

[miscco]

Thanks a lot for raising this potential performance issue. We have moved to our new mono repo, I will duplicate the issue there.

[jrhemstad]

Thanks a lot for raising this potential performance issue. We have moved to our new mono repo, I will duplicate the issue there.

I'll just transfer it.

[jrhemstad]

So it seems that what we're missing here is that we should detect when the input/output iterator satisfy is_contiguous_iterator we should just use a memcpy?

[miscco]

We are currently in the process of merging the first part of <ranges> so detection of contiguous_iterator should be easy enough soon^TM

[jrhemstad]

We may want to land this change sooner, and Thrust already has an is_contiguous_iterator trait, so we can use that in the mean time.

[gevtushenko]

So it seems that what we're missing here is that we should detect when the input/output iterator satisfy is_contiguous_iterator we should just use a memcpy?

@jrhemstad we already identify when it's safe to use memcpy: 1, 2. To my understanding, the issue is that we are using explicit cudaMemcpyHostToDevice and cudaMemcpyDeviceToHost instead of cudaMemcpyDefault that's requested in this issue. @gonzalobg to verify.

[gevtushenko]

The issue was accidentally close, reopening it.

[jrhemstad]

@gonzalobg could you provide a reproducer for this so we can be sure it is addressed correctly?

[gonzalobg]

The only way I could imagine testing this is with benchmarks: benchmarking memory BW and making sure it matches cudaMemcpy. Is that what you are looking for?

[miscco]

I believe we were talking about potential bugs, where managed memory is involved and we cal cudaMemcpyAsync with cudaMemcpyDeviceToDevice and it crashes

[gevtushenko]

@jrhemstad we already identify when it's safe to use memcpy: 1, 2. To my understanding, the issue is that we are using explicit cudaMemcpyHostToDevice and cudaMemcpyDeviceToHost instead of cudaMemcpyDefault that's requested in this issue. @gonzalobg to verify.

@gonzalobg I thought the issue is about using cudaMemcpyDefault. When you specify device execution policy and provide pointers, we'll currently segfault because cudaMemcpyAsync is used with explicit direction:

#include <thrust/device_vector.h>
#include <thrust/execution_policy.h>

int main() {
  constexpr int n = 10;
  thrust::device_vector<int> src(n);
  int dst[n];

  int *src_ptr = thrust::raw_pointer_cast(src.data());
  int *dst_ptr = dst;

  thrust::copy_n(thrust::device, src_ptr, n, dst_ptr);
}

// terminate called after throwing an instance of 'thrust::system::system_error'
// what():  __copy:: D->D: failed: cudaErrorInvalidValue: invalid argument

[gonzalobg]

The issue I am running into is H2D copies being slow for iterators of raw pointer type (T*) because instead of using cudaMemcpy with cudaMemcpyDefault we are doing something else in those cases.

These pointers are allocated with cudaMallocManaged or malloc.