| 525 | } |
| 526 | |
| 527 | void ROCmCompNode::EventImpl::do_device_wait_by(Impl* cn_impl) { |
| 528 | if (cn_impl->dyn_typeinfo() == ROCmCompNodeImpl::typeinfo()) { |
| 529 | auto imp = static_cast<ROCmCompNodeImpl*>(cn_impl); |
| 530 | auto stream = imp->m_env.rocm_env().stream; |
| 531 | imp->activate(); |
| 532 | MGB_ROCM_CHECK(hipStreamWaitEvent(stream, m_hip_event, 0)); |
| 533 | return; |
| 534 | } |
| 535 | if (cn_impl->env().property().type == DeviceType::CPU) { |
| 536 | auto waiter = [this]() { MGB_ROCM_CHECK(hipEventSynchronize(m_hip_event)); }; |
| 537 | cn_impl->add_callback(std::move(waiter)); |
| 538 | return; |
| 539 | } |
| 540 | mgb_throw(MegBrainError, "unimplemented event device_wait_by config"); |
| 541 | } |
| 542 | |
| 543 | /* ===================== ROCmCompNode static methods ===================== */ |
| 544 |
nothing calls this directly
no test coverage detected