diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index a7a3efd26e..1bd6a0ff00 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -2691,7 +2691,10 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, sp->vqueue_header = vqVM; sp->parentAQL = sp->kernarg_address + sizeof(SchedulerParam); - sp->eng_clk = (1000 * 1024) / dev().info().maxEngineClockFrequency_; + + if (dev().info().maxEngineClockFrequency_ > 0) { + sp->eng_clk = (1000 * 1024) / dev().info().maxEngineClockFrequency_; + } // Use a device side global atomics to workaround the reliance of PCIe 3 atomics sp->write_index = hsa_queue_load_write_index_relaxed(schedulerQueue); diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index e63ab77eb0..e352577d02 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -1147,7 +1147,9 @@ bool Device::populateOCLDeviceConstants() { //TODO: add the assert statement for Raven if (!(isa().versionMajor() == 9 && isa().versionMinor() == 0 && isa().versionStepping() == 2)) { - assert(info_.maxEngineClockFrequency_ > 0); + if (info_.maxEngineClockFrequency_ <= 0) { + LogError("maxEngineClockFrequency_ is NOT positive!"); + } } if (HSA_STATUS_SUCCESS !=