AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Replacing deprecated HSA API calls with newer ones

Browse Source 
Change-Id: Iebe2c00e717ab0e47c61611752b717966c719994


[ROCm/clr commit: cd1e364911]
This commit is contained in:
Jatin Chaudhary
2020-06-24 03:24:27 -04:00
committed by Jatin Chaudhary
parent 8d083a3c61
commit 260a83c546
4 changed files with 16 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/rocclr/device/rocm/rocblit.cpp
+7 -7
View File
@@ -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<uint64_t>::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<uint64_t>::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;
projects/clr/rocclr/device/rocm/rocdevice.cpp
+1 -1
View File
@@ -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;
projects/clr/rocclr/device/rocm/rochostcall.cpp
+2 -2
View File
@@ -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) {
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+6 -6
View File
@@ -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<uint32_t*>(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;