From 260a83c5461e7854d5ec2a07ba25543e2f17bd38 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Wed, 24 Jun 2020 03:24:27 -0400 Subject: [PATCH] Replacing deprecated HSA API calls with newer ones Change-Id: Iebe2c00e717ab0e47c61611752b717966c719994 [ROCm/clr commit: cd1e364911ce66ddf858ed93159629981bab7971] --- projects/clr/rocclr/device/rocm/rocblit.cpp | 14 +++++++------- projects/clr/rocclr/device/rocm/rocdevice.cpp | 2 +- projects/clr/rocclr/device/rocm/rochostcall.cpp | 4 ++-- projects/clr/rocclr/device/rocm/rocvirtual.cpp | 12 ++++++------ 4 files changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index fbd4ee6366..de60713165 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -448,7 +448,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d } - hsa_signal_value_t val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); if (val != 0) { LogError("Async copy failed"); @@ -476,7 +476,7 @@ bool DmaBlitManager::copyBufferRect(device::Memory& srcMemory, device::Memory& d } } - hsa_signal_value_t val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); if (val != 0) { LogError("Async copy failed"); @@ -651,10 +651,10 @@ bool DmaBlitManager::hsaCopy(const Memory& srcMemory, const Memory& dstMemory, constexpr size_t small_transfer_size = 4 * Mi; if (size[0] < small_transfer_size) { - val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, std::numeric_limits::max(), HSA_WAIT_STATE_ACTIVE); } else { - val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, + val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, std::numeric_limits::max(), HSA_WAIT_STATE_BLOCKED); } if (val != (kInitVal - 1)) { @@ -706,7 +706,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ status = hsa_amd_memory_async_copy(hostDst + offset, dev().getBackendDevice(), hsaBuffer, srcAgent, size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { - hsa_signal_value_t val = hsa_signal_wait_acquire( + hsa_signal_value_t val = hsa_signal_wait_scacquire( completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); if (val != (kInitVal - 1)) { @@ -733,7 +733,7 @@ bool DmaBlitManager::hsaCopyStaged(const_address hostSrc, address hostDst, size_ hsa_amd_memory_async_copy(hsaBuffer, dstAgent, hostSrc + offset, dev().getBackendDevice(), size, 0, nullptr, completion_signal_); if (status == HSA_STATUS_SUCCESS) { - hsa_signal_value_t val = hsa_signal_wait_acquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, + hsa_signal_value_t val = hsa_signal_wait_scacquire(completion_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED); if (val != (kInitVal - 1)) { @@ -2340,7 +2340,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, } releaseArguments(parameters); - if (hsa_signal_wait_acquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1), + if (hsa_signal_wait_scacquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogWarning("Failed schedulerSignal wait"); return false; diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index 3a48de13d8..9493cda2cf 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -2157,7 +2157,7 @@ bool Device::SvmAllocInit(void* memory, size_t size) const { } // Wait for the prefetch - if (hsa_signal_wait_acquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + if (hsa_signal_wait_scacquire(prefetch_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogError("Barrier packet submission failed"); return false; diff --git a/projects/clr/rocclr/device/rocm/rochostcall.cpp b/projects/clr/rocclr/device/rocm/rochostcall.cpp index 743f24f139..cdebdba6c4 100644 --- a/projects/clr/rocclr/device/rocm/rochostcall.cpp +++ b/projects/clr/rocclr/device/rocm/rochostcall.cpp @@ -316,7 +316,7 @@ void HostcallListener::consumePackets() { while (true) { while (true) { - uint64_t new_value = hsa_signal_wait_acquire(doorbell_, HSA_SIGNAL_CONDITION_NE, signal_value, timeout, + uint64_t new_value = hsa_signal_wait_scacquire(doorbell_, HSA_SIGNAL_CONDITION_NE, signal_value, timeout, HSA_WAIT_STATE_BLOCKED); if (new_value != signal_value) { signal_value = new_value; @@ -344,7 +344,7 @@ void HostcallListener::terminate() { return; } - hsa_signal_store_release(doorbell_, SIGNAL_DONE); + hsa_signal_store_screlease(doorbell_, SIGNAL_DONE); // FIXME_lmoriche: fix termination handshake while (thread_.state() < Thread::FINISHED) { diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index 57cc07de43..7a25670b5d 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -512,11 +512,11 @@ bool VirtualGPU::dispatchGenericAqlPacket( } //hsa_queue_store_write_index_release(gpu_queue_, index); - hsa_signal_store_release(gpu_queue_->doorbell_signal, index - 1); + hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index - 1); // Wait on signal ? if (blocking) { - if (hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), + if (hsa_signal_wait_scacquire(signal, HSA_SIGNAL_CONDITION_LT, 1, uint64_t(-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogPrintfError("Failed signal [0x%lx] wait", signal.handle); return false; @@ -577,7 +577,7 @@ void VirtualGPU::dispatchBarrierPacket(const hsa_barrier_and_packet_t* packet) { *aql_loc = *packet; __atomic_store_n(reinterpret_cast(aql_loc), kBarrierPacketHeader, __ATOMIC_RELEASE); - hsa_signal_store_release(gpu_queue_->doorbell_signal, index); + hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); ClPrint(amd::LOG_DEBUG, amd::LOG_AQL, "[%zx] HWq=0x%zx, BarrierAND Header = 0x%x (type=%d, barrier=%d, acquire=%d, release=%d), " "dep_signal=[0x%zx, 0x%zx, 0x%zx, 0x%zx, 0x%zx], completion_signal=0x%zx", @@ -612,7 +612,7 @@ bool VirtualGPU::releaseGpuMemoryFence() { // Dispatch barrier packet into the queue and wait till it finishes. dispatchBarrierPacket(&barrier_packet_); - if (hsa_signal_wait_acquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogError("Barrier packet submission failed"); return false; @@ -849,7 +849,7 @@ void* VirtualGPU::allocKernArg(size_t size, size_t alignment) { // Dispatch barrier packet into the queue and wait till it finishes. dispatchBarrierPacket(&barrier_packet_); - if (hsa_signal_wait_acquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogError("Kernel arguments reset failed"); } @@ -1205,7 +1205,7 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { hsa_amd_svm_prefetch_async(cmd.dev_prt(), cmd.count(), agent, 0, nullptr, barrier_signal_); // Wait for the prefetch - if (hsa_signal_wait_acquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), + if (hsa_signal_wait_scacquire(barrier_signal_, HSA_SIGNAL_CONDITION_EQ, 0, uint64_t(-1), HSA_WAIT_STATE_BLOCKED) != 0) { LogError("Barrier packet submission failed"); return false;