| 369 | }; |
| 370 | |
| 371 | void Stream::holdResources(LockResources usedResources) |
| 372 | { |
| 373 | if (!usedResources.empty()) |
| 374 | { |
| 375 | // Looks like a good place to clear the gc bag, as every time we create |
| 376 | // a new closure that eventually gets added to the bag, we empty it. |
| 377 | // The bag shouldn't grow unlimited. |
| 378 | // Calling it before allocating a new closure just avoid having two |
| 379 | // closures not inside a cuda stream that are simultaneously alive, but |
| 380 | // in practice it doesn't seem to matter much. |
| 381 | ClearGCBag(); |
| 382 | |
| 383 | auto closure = std::make_unique<HostFunctionClosure>(); |
| 384 | |
| 385 | closure->stream = this->shared_from_this(); |
| 386 | closure->resources = std::move(usedResources); |
| 387 | |
| 388 | auto fn = [](cudaStream_t stream, cudaError_t error, void *userData) -> void |
| 389 | { |
| 390 | std::unique_ptr<HostFunctionClosure> pclosure(reinterpret_cast<HostFunctionClosure *>(userData)); |
| 391 | NVCV_ASSERT(pclosure != nullptr); |
| 392 | AddToGCBag(std::move(pclosure)); |
| 393 | }; |
| 394 | |
| 395 | // If we naively execute the callback in the main stream (m_handle), the GPU will wait until the callback |
| 396 | // is executed (on host). For correctness, GPU doesn't need to wait - it's the CPU that needs |
| 397 | // to wait for the work already scheduled to complete. |
| 398 | // |
| 399 | // Naive timeline: |
| 400 | // |
| 401 | // stream GPU_kernel1 | Callback | GPU_kernel2 |
| 402 | // GPU activity xxxxxxxxxxx xxxxxxxxxxx |
| 403 | // CPU activity xxxxxxxx |
| 404 | // |
| 405 | // Optimized timeline |
| 406 | // |
| 407 | // |
| 408 | // event -----v |
| 409 | // stream GPU_kernel1 | GPU_kernel2 |
| 410 | // aux_stream waitEvent >| Callback |
| 411 | // |
| 412 | // GPU activity xxxxxxxxxxx xxxxxxxxxxx |
| 413 | // CPU activity xxxxxxxx |
| 414 | |
| 415 | util::CheckThrow(cudaEventRecord(m_event, m_handle)); // add async record the event in the main stream |
| 416 | util::CheckThrow( |
| 417 | cudaStreamWaitEvent(GetAuxStream(), m_event)); // add async wait for the event in the aux stream |
| 418 | |
| 419 | // cudaStreamAddCallback pushes a task to the given stream, which at some point (asynchonously) calls |
| 420 | // the given callback (fn), passing to it the closure we created, among other stream states. |
| 421 | // When fn is executed, the refcnt of all objects that the closure holds will eventually be decremented, which |
| 422 | // will trigger their deletion if refcnt==0. This effectively extends the objects' lifetime until |
| 423 | // all tasks that refer to them are finished. |
| 424 | |
| 425 | // The callback will be executed in the singleton aux stream there may be contention with other callbacks and waitEvents from |
| 426 | // other streams. However the callback is used to release resources from the cache and should not be a performance bottleneck. |
| 427 | // This avoids opening a new aux stream for each stream object. |
| 428 |
no test coverage detected