您好，我们已经收到了您的问题，会安排技术人员尽快解答您的问题，请耐心等待。请您再次检查是否提供了清晰的问题描述、复现代码、环境&版本、报错信息等。同时，您也可以通过查看官网API文档、常见问题、历史Issue、AI社区来寻求解答。祝您生活愉快～

Hi! We've received your issue and please be patient to get responded. We will arrange technicians to answer your questions as soon as possible. Please make sure that you have posted enough message to demo your request. You may also check out the API，FAQ，Github Issue and AI community to get the answer.Have a nice day!

add PIC for Paddle with DALI integration - @zhiqiu

Hi, @jeng1220@JanuszL I tried to compile your PR3305 but failed these days. Could you give me a built wheel package for debug?

And after reading your PR, I suggest that you can try to use cudaEventRecord on compute stream +cudaStreamWaitEvent on dali stream instead of fluid.core._cuda_synchronize (which sync all the paddle compute streams) between mutable_data and feed_ndarray.

pesudo code:

ptr = lod_tensor._mutable_data(category_place[cat],
                                               category_pd_type[cat])
event = cudaCreateEvent()
cudaEventRecord(event, paddle.device.cuda.current_stream(device_id))
cudaStreamWaitEvent(copy_stream, event)
feed_ndarray(tensor, ptr, copy_stream)

Just like paddle's dataloder doing:

Paddle/paddle/fluid/operators/reader/buffered_reader.cc

Lines 158 to 198 in 97a73e1

| | // NOTE(zjl): cudaStreamWaitEvent() must be called after all |
| | // cuda[i].mutable_data() is called, since some ops release |
| | // cuda memory immediately without waiting cuda kernel ends |
| | platform::SetDeviceId( |
| | BOOST_GET_CONST(platform::CUDAPlace, place_).device); |
| | #ifdef PADDLE_WITH_HIP |
| | PADDLE_ENFORCE_CUDA_SUCCESS( |
| | hipEventRecord(events_[i].get(), compute_stream_)); |
| | PADDLE_ENFORCE_CUDA_SUCCESS( |
| | hipStreamWaitEvent(stream_.get(), events_[i].get(), 0)); |
| | #else |
| | PADDLE_ENFORCE_CUDA_SUCCESS( |
| | cudaEventRecord(events_[i].get(), compute_stream_)); |
| | PADDLE_ENFORCE_CUDA_SUCCESS( |
| | cudaStreamWaitEvent(stream_.get(), events_[i].get(), 0)); |
| | #endif |
| | |
| | platform::RecordEvent record_event("BufferedReader:MemoryCopy"); |
| | for (size_t i = 0; i < cpu.size(); ++i) { |
| | auto cpu_place = cpu[i].place(); |
| | auto cpu_ptr = cpu[i].data(); |
| | auto gpu_ptr = gpu_ptrs[i]; |
| | auto size = |
| | cpu[i].numel() * paddle::framework::SizeOfType(cpu[i].type()); |
| | if (platform::is_cuda_pinned_place(cpu_place)) { |
| | memory::Copy(BOOST_GET_CONST(platform::CUDAPlace, place_), gpu_ptr, |
| | BOOST_GET_CONST(platform::CUDAPinnedPlace, cpu_place), |
| | cpu_ptr, size, stream_.get()); |
| | } elseif ((platform::is_gpu_place(cpu_place))) { |
| | memory::Copy(BOOST_GET_CONST(platform::CUDAPlace, place_), gpu_ptr, |
| | BOOST_GET_CONST(platform::CUDAPlace, cpu_place), |
| | cpu_ptr, size, stream_.get()); |
| | } else { |
| | platform::CUDAPinnedPlace cuda_pinned_place; |
| | framework::LoDTensor cuda_pinned_tensor; |
| | cuda_pinned_tensor.Resize(cpu[i].dims()); |
| | auto cuda_pinned_ptr = cuda_pinned_tensor.mutable_data( |
| | cuda_pinned_place, cpu[i].type()); |
| | memory::Copy(cuda_pinned_place, cuda_pinned_ptr, |
| | BOOST_GET_CONST(platform::CPUPlace, cpu_place), |
| | cpu_ptr, size); |

Hi @zhiqiu,

Thank you for looking into that issue. To repro it, you don't have to build DALI, the changes from the mentioned PR relate only to the python part, so you can install the latest DALI and just update the relevant python files in the location where your wheel is installed (in my case /usr/local/lib/python3.6/dist-packages/nvidia/dali/plugin/).

Just to be sure, do I need to call cudaEventRecord->cudaStreamWaitEvent or can I just feed_ndarray(tensor, ptr, paddle.device.cuda.current_stream(device_id)). This way DALI would do the copy on the same stream as PaddlePaddle ops operate on so it should be safe.

Another question - when the python API that exposes CUDA streams is going to be available (I see that is available only on the development branch).

Thanks. I will try your method.

I think the better way is to use two stream and use event to sync them, in that case, the data copy stream and compute stream can overlap to improve performance.

The paddle.device.cuda.current_stream(device_id)) will be released in Paddle-2.2rc (next month). If you need to use it, you can install the nightly-build version. Here is the command for cuda-11.2 version:
python -m pip install paddlepaddle-gpu==0.0.0.post112 -f https://www.paddlepaddle.org.cn/whl/linux/gpu/develop.html

@zhiqiu,

I'll give it a try. Thanks for the help.

Hi @zhiqiu,

I tested the method by creating and recording the event. It seems to work fine.

In the case of DALI device2device copy is performed and it cannot overlap with other kernels. However I cannot do feed_ndarray(tensor, ptr, paddle.device.cuda.current_stream(device_id)) as current_stream call returns PaddlePaddle object and it is not possible to extract a raw cuda stream that DALI expects. I checked the API and I was not able to find way, even raw_stream method from cuda_stream.h is not exposed in the binary so I cannot call it indirectly (using ctypes.CDLL).
Do you plan to expose streams and events raw pointer so other libraries can use them?

OK, I got it. I will find how to expose the raw pointer of current_stream.

Hi, @JanuszL Please refer to: #35813

Hi @zhiqiu,

That should do. I will check the nightly build as soon as it is merged. Thanks a lot.