| 357 | } |
| 358 | |
| 359 | void ROCmCompNodeImpl::sync() { |
| 360 | activate(); |
| 361 | |
| 362 | // same behavior as cuda |
| 363 | // do not use MGB_ROCM_CHECK(hipStreamSynchronize(m_env->stream)) since |
| 364 | // other threads may be adding operations into the stream, and we only care |
| 365 | // about previous operations in current thread. However docs of |
| 366 | // hipStreamSynchronize did not describe details of such condition, so we |
| 367 | // use manual event implementation |
| 368 | |
| 369 | Event* event; |
| 370 | { |
| 371 | MGB_LOCK_GUARD(m_sync_event_mtx); |
| 372 | if (!m_sync_event) |
| 373 | m_sync_event = create_event(0); |
| 374 | event = m_sync_event.get(); |
| 375 | } |
| 376 | event->record(); |
| 377 | event->host_wait(); |
| 378 | } |
| 379 | |
| 380 | void ROCmCompNodeImpl::enable_peer_access(int dev0, int dev1) { |
| 381 | static bool already_enabled[StaticData::MAX_NR_DEVICE][StaticData::MAX_NR_DEVICE]; |