Switch util functions for HIP to hipMemcpyWithStream.
This API was introduced in ROCm 3.1 and provides a much faster path for
copies than hipMemcpyAsync followed by stream synchronize for H2D and
D2H copies, especially for unpinned copies.
MODIFY
thrust/system/hip/detail/util.h
thrust/system/hip/detail/util.h
@@ -65,10 +65,7 @@ trivial_copy_from_device(Type* dst, Type const* src, size_t count, hipStream_t s
if(count == 0)
return status;
- status = ::hipMemcpyAsync(dst, src, sizeof(Type) * count, hipMemcpyDeviceToHost, stream);
- if(status != hipSuccess)
- return status;
- status = hipStreamSynchronize(stream);
+ status = ::hipMemcpyWithStream(dst, src, sizeof(Type) * count, hipMemcpyDeviceToHost, stream);
return status;
}
@@ -80,10 +77,7 @@ trivial_copy_to_device(Type* dst, Type const* src, size_t count, hipStream_t str
if(count == 0)
return status;
- status = ::hipMemcpyAsync(dst, src, sizeof(Type) * count, hipMemcpyHostToDevice, stream);
- if(status != hipSuccess)
- return status;
- status = hipStreamSynchronize(stream);
+ status = ::hipMemcpyWithStream(dst, src, sizeof(Type) * count, hipMemcpyHostToDevice, stream);
return status;
}
template <class Policy, class Type>

Switch util functions for HIP to hipMemcpyWithStream. (#94)

This API was introduced in ROCm 3.1 and provides a much faster path for
copies than hipMemcpyAsync followed by stream synchronize for H2D and
D2H copies, especially for unpinned copies.