| 608 | } |
| 609 | |
| 610 | void CudaCompNodeImpl::sync() { |
| 611 | activate(); |
| 612 | |
| 613 | // do not use MGB_CUDA_CHECK(cudaStreamSynchronize(m_env->stream)) since |
| 614 | // other threads may be adding operations into the stream, and we only care |
| 615 | // about previous operations in current thread. However docs of |
| 616 | // cudaStreamSynchronize did not describe details of such condition, so we |
| 617 | // use manual event implementation |
| 618 | |
| 619 | Event* event; |
| 620 | { |
| 621 | MGB_LOCK_GUARD(m_sync_event_mtx); |
| 622 | if (!m_sync_event) |
| 623 | m_sync_event = create_event(0); |
| 624 | event = m_sync_event.get(); |
| 625 | } |
| 626 | event->record(); |
| 627 | event->host_wait(); |
| 628 | } |
| 629 | |
| 630 | void CudaCompNodeImpl::enable_peer_access(int dev0, int dev1) { |
| 631 | static bool already_enabled[MAX_NR_DEVICE][MAX_NR_DEVICE]; |