From f1486e33e409ed2dacfc80172447f851954dca2b Mon Sep 17 00:00:00 2001 From: Sourabh Date: Tue, 31 Aug 2021 16:14:13 +0000 Subject: [PATCH] SWDEV-292525 - [hipamd] Use a blit to streamWait Implementation to use a blit kernel to perform a hipStreamWait instead of an AQL packet. Change-Id: I9d9bc1dd793c1454e91e68afb793460749e10d2e [ROCm/clr commit: 4177d20877585ce067fbe1e39f8d789572dc4993] --- projects/clr/hipamd/src/hip_stream_ops.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/src/hip_stream_ops.cpp b/projects/clr/hipamd/src/hip_stream_ops.cpp index 88487358d4..6e5be889af 100644 --- a/projects/clr/hipamd/src/hip_stream_ops.cpp +++ b/projects/clr/hipamd/src/hip_stream_ops.cpp @@ -36,8 +36,8 @@ hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void // 'flags' for now used only for Wait, but in future there will usecases for Write too. if (cmdType == ROCCLR_COMMAND_STREAM_WAIT_VALUE) { - // Wait is only supported on SignalMemory objects - if (!(memory->getMemFlags() & ROCCLR_MEM_HSA_SIGNAL_MEMORY)) { + // Stream Wait on AQL barrier-value type packet is only supported on SignalMemory objects + if (GPU_STREAMOPS_CP_WAIT && (!(memory->getMemFlags() & ROCCLR_MEM_HSA_SIGNAL_MEMORY))) { return hipErrorInvalidValue; } switch (flags) { @@ -88,7 +88,7 @@ hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, uint32_t value, u value, mask, flags, - 0)); // sizeBytes un-used for wait, set it to 0 + sizeof(uint32_t))); } hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags, @@ -101,7 +101,7 @@ hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, uint64_t value, u value, mask, flags, - 0)); // sizeBytes un-used for wait, set it to 0 + sizeof(uint64_t))); } hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags) { @@ -113,7 +113,7 @@ hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, value, 0, // mask un-used set it to 0 0, // flags un-used for now set it to 0 - 4)); + sizeof(uint32_t))); } hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags) { @@ -125,5 +125,5 @@ hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, value, 0, // mask un-used set it to 0 0, // flags un-used for now set it to 0 - 8)); + sizeof(uint64_t))); }