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

SWDEV-381402 - Derive hip::Stream from amd::HostQueue

Browse Source 
Change-Id: I6c1aca5eb350c32d974ae4ffcc725705355956d8


[ROCm/clr commit: e3633dc8f4]
This commit is contained in:
Ioannis Assiouras
2023-02-08 20:18:11 +00:00
parent b91c27aef1
commit 2633e62801
19 changed files with 403 additions and 732 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_memory.cpp
+145 -145
View File
@@ -78,9 +78,9 @@ hipError_t ihipFree(void *ptr) {
auto dev = g_devices[device_id];
// Skip stream allocation, since if it wasn't allocated until free, then the device wasn't used
constexpr bool SkipStreamAlloc = true;
amd::HostQueue* queue = dev->NullStream(SkipStreamAlloc);
if (queue != nullptr) {
queue->finish();
hip::Stream* stream = dev->NullStream(SkipStreamAlloc);
if (stream != nullptr) {
stream->finish();
}
hip::Stream::syncNonBlockingStreams(device_id);
// Find out if memory belongs to any memory pool
@@ -195,15 +195,15 @@ hipError_t hipSignalExternalSemaphoresAsync(
if (extSemArray == nullptr || paramsArray == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
for (unsigned int i = 0; i < numExtSems; i++) {
if (extSemArray[i] != nullptr) {
amd::ExternalSemaphoreCmd* command =
new amd::ExternalSemaphoreCmd(*queue, extSemArray[i], paramsArray[i].params.fence.value,
new amd::ExternalSemaphoreCmd(*hip_stream, extSemArray[i], paramsArray[i].params.fence.value,
amd::ExternalSemaphoreCmd::COMMAND_SIGNAL_EXTSEMAPHORE);
if (command == nullptr) {
return hipErrorOutOfMemory;
@@ -227,15 +227,15 @@ hipError_t hipWaitExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemAr
if (extSemArray == nullptr || paramsArray == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
for (unsigned int i = 0; i < numExtSems; i++) {
if (extSemArray[i] != nullptr) {
amd::ExternalSemaphoreCmd* command =
new amd::ExternalSemaphoreCmd(*queue, extSemArray[i], paramsArray[i].params.fence.value,
new amd::ExternalSemaphoreCmd(*hip_stream, extSemArray[i], paramsArray[i].params.fence.value,
amd::ExternalSemaphoreCmd::COMMAND_WAIT_EXTSEMAPHORE);
if (command == nullptr) {
return hipErrorOutOfMemory;
@@ -343,35 +343,35 @@ hipError_t ihipMemcpy_validate(void* dst, const void* src, size_t sizeBytes,
}
hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, amd::HostQueue& queue, bool isAsync) {
hipMemcpyKind kind, hip::Stream& stream, bool isAsync) {
amd::Command::EventWaitList waitList;
size_t sOffset = 0;
amd::Memory* srcMemory = getMemoryObject(src, sOffset);
size_t dOffset = 0;
amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
amd::Device* queueDevice = &queue.device();
amd::Device* queueDevice = &stream.device();
amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
amd::HostQueue* pQueue = &queue;
hip::Stream* pStream = &stream;
if (queueDevice != dstMemory->getContext().devices()[0]) {
pQueue = hip::getNullStream(dstMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
pStream = hip::getNullStream(dstMemory->getContext());
amd::Command* cmd = stream.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
command = new amd::WriteMemoryCommand(*pQueue, CL_COMMAND_WRITE_BUFFER, waitList,
command = new amd::WriteMemoryCommand(*pStream, CL_COMMAND_WRITE_BUFFER, waitList,
*dstMemory->asBuffer(), dOffset, sizeBytes, src, 0, 0, copyMetadata);
} else if ((srcMemory != nullptr) && (dstMemory == nullptr)) {
amd::HostQueue* pQueue = &queue;
hip::Stream* pStream = &stream;
if (queueDevice != srcMemory->getContext().devices()[0]) {
pQueue = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
pStream = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = stream.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
command = new amd::ReadMemoryCommand(*pQueue, CL_COMMAND_READ_BUFFER, waitList,
command = new amd::ReadMemoryCommand(*pStream, CL_COMMAND_READ_BUFFER, waitList,
*srcMemory->asBuffer(), sOffset, sizeBytes, dst, 0, 0, copyMetadata);
} else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
// Check if the queue device doesn't match the device on any memory object.
@@ -380,7 +380,7 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
if ((srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) &&
((srcMemory->getContext().devices().size() == 1) &&
(dstMemory->getContext().devices().size() == 1))) {
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
command = new amd::CopyMemoryP2PCommand(stream, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
if (command == nullptr) {
return hipErrorOutOfMemory;
@@ -392,12 +392,12 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
return hipErrorInvalidValue;
}
} else {
amd::HostQueue* pQueue = &queue;
hip::Stream* pStream = &stream;
if ((srcMemory->getContext().devices()[0] == dstMemory->getContext().devices()[0]) &&
(queueDevice != srcMemory->getContext().devices()[0])) {
copyMetadata.copyEnginePreference_ = amd::CopyMetadata::CopyEnginePreference::NONE;
pQueue = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
pStream = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = stream.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
@@ -405,22 +405,22 @@ hipError_t ihipMemcpyCommand(amd::Command*& command, void* dst, const void* src,
// Scenarios such as DtoH where dst is pinned memory
if ((queueDevice != srcMemory->getContext().devices()[0]) &&
(dstMemory->getContext().devices().size() != 1)) {
pQueue = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
pStream = hip::getNullStream(srcMemory->getContext());
amd::Command* cmd = stream.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
// Scenarios such as HtoD where src is pinned memory
} else if ((queueDevice != dstMemory->getContext().devices()[0]) &&
(srcMemory->getContext().devices().size() != 1)) {
pQueue = hip::getNullStream(dstMemory->getContext());
amd::Command* cmd = queue.getLastQueuedCommand(true);
pStream = hip::getNullStream(dstMemory->getContext());
amd::Command* cmd = stream.getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
}
}
}
command = new amd::CopyMemoryCommand(*pQueue, CL_COMMAND_COPY_BUFFER, waitList,
command = new amd::CopyMemoryCommand(*pStream, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes,
copyMetadata);
}
@@ -445,13 +445,13 @@ bool IsHtoHMemcpy(void* dst, const void* src, hipMemcpyKind kind) {
}
return false;
}
void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, amd::HostQueue& queue) {
queue.finish();
void ihipHtoHMemcpy(void* dst, const void* src, size_t sizeBytes, hip::Stream& stream) {
stream.finish();
memcpy(dst, src, sizeBytes);
}
// ================================================================================================
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
amd::HostQueue& queue, bool isAsync = false) {
hip::Stream& stream, bool isAsync = false) {
hipError_t status;
if (sizeBytes == 0) {
// Skip if nothing needs writing.
@@ -470,7 +470,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
size_t dOffset = 0;
amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
if (srcMemory == nullptr && dstMemory == nullptr) {
ihipHtoHMemcpy(dst, src, sizeBytes, queue);
ihipHtoHMemcpy(dst, src, sizeBytes, stream);
return hipSuccess;
} else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
isAsync = false;
@@ -483,7 +483,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
isP2P = true;
}
amd::Command* command = nullptr;
status = ihipMemcpyCommand(command, dst, src, sizeBytes, kind, queue, isAsync);
status = ihipMemcpyCommand(command, dst, src, sizeBytes, kind, stream, isAsync);
if (status != hipSuccess) {
return status;
}
@@ -491,22 +491,22 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
if (!isAsync) {
command->awaitCompletion();
} else if (isP2P) {
amd::HostQueue* pQueue = hip::getNullStream(dstMemory->getContext());
hip::Stream* pStream = hip::getNullStream(dstMemory->getContext());
amd::Command::EventWaitList waitList;
waitList.push_back(command);
amd::Command* depdentMarker = new amd::Marker(*pQueue, false, waitList);
amd::Command* depdentMarker = new amd::Marker(*pStream, false, waitList);
if (depdentMarker != nullptr) {
depdentMarker->enqueue();
depdentMarker->release();
}
} else {
amd::HostQueue* newQueue = command->queue();
if (newQueue != &queue) {
if (newQueue != &stream) {
amd::Command::EventWaitList waitList;
amd::Command* cmd = newQueue->getLastQueuedCommand(true);
if (cmd != nullptr) {
waitList.push_back(cmd);
amd::Command* depdentMarker = new amd::Marker(queue, true, waitList);
amd::Command* depdentMarker = new amd::Marker(stream, true, waitList);
if (depdentMarker != nullptr) {
depdentMarker->enqueue();
depdentMarker->release();
@@ -611,18 +611,18 @@ hipError_t hipFree(void* ptr) {
hipError_t hipMemcpy_common(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream = nullptr) {
CHECK_STREAM_CAPTURING();
amd::HostQueue* queue = nullptr;
hip::Stream* hip_stream = nullptr;
if (stream != nullptr) {
queue = hip::getQueue(stream);
hip_stream = hip::getStream(stream);
} else {
queue = hip::getNullStream();
hip_stream = hip::getNullStream();
}
if (queue == nullptr) {
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
return ihipMemcpy(dst, src, sizeBytes, kind, *queue);
return ihipMemcpy(dst, src, sizeBytes, kind, *hip_stream);
}
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
@@ -643,12 +643,12 @@ hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
HIP_RETURN(hipErrorContextIsDestroyed);
}
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(ihipMemcpy(dst, src, sizeBytes, kind, *queue, false));
HIP_RETURN_DURATION(ihipMemcpy(dst, src, sizeBytes, kind, *hip_stream, false));
}
hipError_t hipMemPtrGetInfo(void *ptr, size_t *size) {
@@ -697,9 +697,9 @@ hipError_t ihipArrayDestroy(hipArray* array) {
}
for (auto& dev : g_devices) {
amd::HostQueue* queue = dev->NullStream(true);
if (queue != nullptr) {
queue->finish();
hip::Stream* stream = dev->NullStream(true);
if (stream != nullptr) {
stream->finish();
}
}
@@ -1205,9 +1205,9 @@ hipError_t ihipHostUnregister(void* hostPtr) {
// Wait on the device, associated with the current memory object during allocation
auto device_id = mem->getUserData().deviceId;
amd::HostQueue* queue = g_devices[device_id]->NullStream(true);
if (queue != nullptr) {
queue->finish();
hip::Stream* stream = g_devices[device_id]->NullStream(true);
if (stream != nullptr) {
stream->finish();
}
amd::MemObjMap::RemoveMemObj(hostPtr);
@@ -1392,11 +1392,11 @@ hipError_t hipMemcpyHtoD(hipDeviceptr_t dstDevice,
size_t ByteCount) {
HIP_INIT_API(hipMemcpyHtoD, dstDevice, srcHost, ByteCount);
CHECK_STREAM_CAPTURING();
amd::HostQueue* queue = hip::getQueue(nullptr);
if (queue == nullptr) {
hip::Stream* stream = hip::getStream(nullptr);
if (stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(ihipMemcpy(dstDevice, srcHost, ByteCount, hipMemcpyHostToDevice, *queue));
HIP_RETURN_DURATION(ihipMemcpy(dstDevice, srcHost, ByteCount, hipMemcpyHostToDevice, *stream));
}
hipError_t hipMemcpyDtoH(void* dstHost,
@@ -1404,11 +1404,11 @@ hipError_t hipMemcpyDtoH(void* dstHost,
size_t ByteCount) {
HIP_INIT_API(hipMemcpyDtoH, dstHost, srcDevice, ByteCount);
CHECK_STREAM_CAPTURING();
amd::HostQueue* queue = hip::getQueue(nullptr);
if (queue == nullptr) {
hip::Stream* stream = hip::getStream(nullptr);
if (stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *queue));
HIP_RETURN_DURATION(ihipMemcpy(dstHost, srcDevice, ByteCount, hipMemcpyDeviceToHost, *stream));
}
hipError_t hipMemcpyDtoD(hipDeviceptr_t dstDevice,
@@ -1416,22 +1416,22 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dstDevice,
size_t ByteCount) {
HIP_INIT_API(hipMemcpyDtoD, dstDevice, srcDevice, ByteCount);
CHECK_STREAM_CAPTURING();
amd::HostQueue* queue = hip::getQueue(nullptr);
if (queue == nullptr) {
hip::Stream* stream = hip::getStream(nullptr);
if (stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(ihipMemcpy(dstDevice, srcDevice, ByteCount, hipMemcpyDeviceToDevice, *queue));
HIP_RETURN_DURATION(ihipMemcpy(dstDevice, srcDevice, ByteCount, hipMemcpyDeviceToDevice, *stream));
}
hipError_t hipMemcpyAsync_common(void* dst, const void* src, size_t sizeBytes,
hipMemcpyKind kind, hipStream_t stream) {
STREAM_CAPTURE(hipMemcpyAsync, stream, dst, src, sizeBytes, kind);
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
return ihipMemcpy(dst, src, sizeBytes, kind, *queue, true);
return ihipMemcpy(dst, src, sizeBytes, kind, *hip_stream, true);
}
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
@@ -1452,12 +1452,12 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dstDevice, void* srcHost, size_t By
HIP_INIT_API(hipMemcpyHtoDAsync, dstDevice, srcHost, ByteCount, stream);
hipMemcpyKind kind = hipMemcpyHostToDevice;
STREAM_CAPTURE(hipMemcpyHtoDAsync, stream, dstDevice, srcHost, ByteCount, kind);
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(
ihipMemcpy(dstDevice, srcHost, ByteCount, kind, *queue, true));
ihipMemcpy(dstDevice, srcHost, ByteCount, kind, *hip_stream, true));
}
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount,
@@ -1465,12 +1465,12 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice
HIP_INIT_API(hipMemcpyDtoDAsync, dstDevice, srcDevice, ByteCount, stream);
hipMemcpyKind kind = hipMemcpyDeviceToDevice;
STREAM_CAPTURE(hipMemcpyDtoDAsync, stream, dstDevice, srcDevice, ByteCount, kind);
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(
ihipMemcpy(dstDevice, srcDevice, ByteCount, kind, *queue, true));
ihipMemcpy(dstDevice, srcDevice, ByteCount, kind, *hip_stream, true));
}
hipError_t hipMemcpyDtoHAsync(void* dstHost, hipDeviceptr_t srcDevice, size_t ByteCount,
@@ -1478,12 +1478,12 @@ hipError_t hipMemcpyDtoHAsync(void* dstHost, hipDeviceptr_t srcDevice, size_t By
HIP_INIT_API(hipMemcpyDtoHAsync, dstHost, srcDevice, ByteCount, stream);
hipMemcpyKind kind = hipMemcpyDeviceToHost;
STREAM_CAPTURE(hipMemcpyDtoHAsync, stream, dstHost, srcDevice, ByteCount, kind);
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN_DURATION(
ihipMemcpy(dstHost, srcDevice, ByteCount, kind, *queue, true));
ihipMemcpy(dstHost, srcDevice, ByteCount, kind, *hip_stream, true));
}
hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coord3D& srcOrigin,
@@ -1532,7 +1532,7 @@ hipError_t ihipMemcpyAtoDValidate(hipArray* srcArray, void* dstDevice, amd::Coor
hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, void* dstDevice,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
amd::HostQueue* queue) {
hip::Stream* stream) {
amd::BufferRect srcRect;
amd::BufferRect dstRect;
amd::Memory* dstMemory;
@@ -1544,7 +1544,7 @@ hipError_t ihipMemcpyAtoDCommand(amd::Command*& command, hipArray* srcArray, voi
return status;
}
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*stream, CL_COMMAND_COPY_IMAGE_TO_BUFFER,
amd::Command::EventWaitList{}, *srcImage, *dstMemory,
srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
@@ -1606,7 +1606,7 @@ hipError_t ihipMemcpyDtoAValidate(void* srcDevice, hipArray* dstArray, amd::Coor
hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArray* dstArray,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
amd::HostQueue* queue) {
hip::Stream* stream) {
amd::Image* dstImage;
amd::Memory* srcMemory;
amd::BufferRect dstRect;
@@ -1617,7 +1617,7 @@ hipError_t ihipMemcpyDtoACommand(amd::Command*& command, void* srcDevice, hipArr
if (status != hipSuccess) {
return status;
}
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_BUFFER_TO_IMAGE,
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*stream, CL_COMMAND_COPY_BUFFER_TO_IMAGE,
amd::Command::EventWaitList{}, *srcMemory, *dstImage,
srcOrigin, dstOrigin, copyRegion, srcRect, dstRect);
@@ -1679,7 +1679,7 @@ hipError_t ihipMemcpyDtoDValidate(void* srcDevice, void* dstDevice, amd::Coord3D
hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void* dstDevice,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue) {
size_t dstRowPitch, size_t dstSlicePitch, hip::Stream* stream) {
amd::Memory* srcMemory;
amd::Memory* dstMemory;
amd::BufferRect srcRect;
@@ -1694,7 +1694,7 @@ hipError_t ihipMemcpyDtoDCommand(amd::Command*& command, void* srcDevice, void*
amd::Coord3D srcStart(srcRect.start_, 0, 0);
amd::Coord3D dstStart(dstRect.start_, 0, 0);
amd::CopyMemoryCommand* copyCommand = new amd::CopyMemoryCommand(
*queue, CL_COMMAND_COPY_BUFFER_RECT, amd::Command::EventWaitList{}, *srcMemory, *dstMemory,
*stream, CL_COMMAND_COPY_BUFFER_RECT, amd::Command::EventWaitList{}, *srcMemory, *dstMemory,
srcStart, dstStart, copyRegion, srcRect, dstRect);
if (copyCommand == nullptr) {
@@ -1744,7 +1744,7 @@ hipError_t ihipMemcpyDtoHValidate(void* srcDevice, void* dstHost, amd::Coord3D&
hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void* dstHost,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue,
size_t dstRowPitch, size_t dstSlicePitch, hip::Stream* stream,
bool isAsync = false) {
amd::Memory* srcMemory;
amd::BufferRect srcRect;
@@ -1758,7 +1758,7 @@ hipError_t ihipMemcpyDtoHCommand(amd::Command*& command, void* srcDevice, void*
amd::Coord3D srcStart(srcRect.start_, 0, 0);
amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
amd::ReadMemoryCommand* readCommand =
new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER_RECT, amd::Command::EventWaitList{},
new amd::ReadMemoryCommand(*stream, CL_COMMAND_READ_BUFFER_RECT, amd::Command::EventWaitList{},
*srcMemory, srcStart, copyRegion, dstHost, srcRect, dstRect,
copyMetadata);
@@ -1809,7 +1809,7 @@ hipError_t ihipMemcpyHtoDValidate(const void* srcHost, void* dstDevice, amd::Coo
hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, void* dstDevice,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
size_t dstRowPitch, size_t dstSlicePitch, amd::HostQueue* queue,
size_t dstRowPitch, size_t dstSlicePitch, hip::Stream* stream,
bool isAsync = false) {
amd::Memory* dstMemory;
amd::BufferRect srcRect;
@@ -1824,7 +1824,7 @@ hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, vo
amd::Coord3D dstStart(dstRect.start_, 0, 0);
amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
amd::WriteMemoryCommand* writeCommand = new amd::WriteMemoryCommand(
*queue, CL_COMMAND_WRITE_BUFFER_RECT, amd::Command::EventWaitList{}, *dstMemory, dstStart,
*stream, CL_COMMAND_WRITE_BUFFER_RECT, amd::Command::EventWaitList{}, *dstMemory, dstStart,
copyRegion, srcHost, dstRect, srcRect, copyMetadata);
if (writeCommand == nullptr) {
@@ -1842,7 +1842,7 @@ hipError_t ihipMemcpyHtoDCommand(amd::Command*& command, const void* srcHost, vo
hipError_t ihipMemcpyHtoH(const void* srcHost, void* dstHost, amd::Coord3D srcOrigin,
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch,
size_t srcSlicePitch, size_t dstRowPitch, size_t dstSlicePitch,
amd::HostQueue* queue) {
hip::Stream* stream) {
if ((srcHost == nullptr) || (dstHost == nullptr)) {
return hipErrorInvalidValue;
}
@@ -1859,8 +1859,8 @@ hipError_t ihipMemcpyHtoH(const void* srcHost, void* dstHost, amd::Coord3D srcOr
return hipErrorInvalidValue;
}
if (queue) {
queue->finish();
if (stream) {
stream->finish();
}
for (size_t slice = 0; slice < copyRegion[2]; slice++) {
@@ -1909,7 +1909,7 @@ hipError_t ihipMemcpyAtoAValidate(hipArray* srcArray, hipArray* dstArray, amd::C
hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hipArray* dstArray,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, amd::HostQueue* queue) {
amd::Coord3D copyRegion, hip::Stream* stream) {
amd::Image* srcImage;
amd::Image* dstImage;
@@ -1919,7 +1919,7 @@ hipError_t ihipMemcpyAtoACommand(amd::Command*& command, hipArray* srcArray, hip
return status;
}
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*queue, CL_COMMAND_COPY_IMAGE,
amd::CopyMemoryCommand* cpyMemCmd = new amd::CopyMemoryCommand(*stream, CL_COMMAND_COPY_IMAGE,
amd::Command::EventWaitList{}, *srcImage, *dstImage,
srcOrigin, dstOrigin, copyRegion);
@@ -1968,7 +1968,7 @@ hipError_t ihipMemcpyHtoAValidate(const void* srcHost, hipArray* dstArray,
hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hipArray* dstArray,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t srcRowPitch, size_t srcSlicePitch,
amd::HostQueue* queue, bool isAsync = false) {
hip::Stream* stream, bool isAsync = false) {
amd::Image* dstImage;
amd::BufferRect srcRect;
@@ -1980,7 +1980,7 @@ hipError_t ihipMemcpyHtoACommand(amd::Command*& command, const void* srcHost, hi
amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
amd::WriteMemoryCommand* writeMemCmd = new amd::WriteMemoryCommand(
*queue, CL_COMMAND_WRITE_IMAGE, amd::Command::EventWaitList{}, *dstImage, dstOrigin,
*stream, CL_COMMAND_WRITE_IMAGE, amd::Command::EventWaitList{}, *dstImage, dstOrigin,
copyRegion, static_cast<const char*>(srcHost) + srcRect.start_, srcRowPitch, srcSlicePitch,
copyMetadata);
@@ -2029,7 +2029,7 @@ hipError_t ihipMemcpyAtoHValidate(hipArray* srcArray, void* dstHost, amd::Coord3
hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, void* dstHost,
amd::Coord3D srcOrigin, amd::Coord3D dstOrigin,
amd::Coord3D copyRegion, size_t dstRowPitch, size_t dstSlicePitch,
amd::HostQueue* queue, bool isAsync = false) {
hip::Stream* stream, bool isAsync = false) {
amd::Image* srcImage;
amd::BufferRect dstRect;
amd::CopyMetadata copyMetadata(isAsync, amd::CopyMetadata::CopyEnginePreference::SDMA);
@@ -2041,7 +2041,7 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, voi
}
amd::ReadMemoryCommand* readMemCmd = new amd::ReadMemoryCommand(
*queue, CL_COMMAND_READ_IMAGE, amd::Command::EventWaitList{}, *srcImage, srcOrigin,
*stream, CL_COMMAND_READ_IMAGE, amd::Command::EventWaitList{}, *srcImage, srcOrigin,
copyRegion, static_cast<char*>(dstHost) + dstRect.start_, dstRowPitch, dstSlicePitch,
copyMetadata);
@@ -2058,7 +2058,7 @@ hipError_t ihipMemcpyAtoHCommand(amd::Command*& command, hipArray* srcArray, voi
}
hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3D* pCopy,
amd::HostQueue* queue) {
hip::Stream* stream) {
// If {src/dst}MemoryType is hipMemoryTypeUnified, {src/dst}Device and {src/dst}Pitch specify the
// (unified virtual address space) base address of the source data and the bytes per row to apply.
// {src/dst}Array is ignored.
@@ -2106,41 +2106,41 @@ hipError_t ihipGetMemcpyParam3DCommand(amd::Command*& command, const HIP_MEMCPY3
// Host to Device.
return ihipMemcpyHtoDCommand(command, pCopy->srcHost, pCopy->dstDevice, srcOrigin, dstOrigin,
copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, queue);
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream);
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeHost)) {
// Device to Host.
return ihipMemcpyDtoHCommand(command, pCopy->srcDevice, pCopy->dstHost, srcOrigin, dstOrigin,
copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, queue);
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream);
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeDevice)) {
// Device to Device.
return ihipMemcpyDtoDCommand(command, pCopy->srcDevice, pCopy->dstDevice, srcOrigin, dstOrigin,
copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, queue);
pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight, stream);
} else if ((srcMemoryType == hipMemoryTypeHost) && (dstMemoryType == hipMemoryTypeArray)) {
// Host to Image.
return ihipMemcpyHtoACommand(command, pCopy->srcHost, pCopy->dstArray, srcOrigin, dstOrigin,
copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
queue);
stream);
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeHost)) {
// Image to Host.
return ihipMemcpyAtoHCommand(command, pCopy->srcArray, pCopy->dstHost, srcOrigin, dstOrigin,
copyRegion, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight,
queue);
stream);
} else if ((srcMemoryType == hipMemoryTypeDevice) && (dstMemoryType == hipMemoryTypeArray)) {
// Device to Image.
return ihipMemcpyDtoACommand(command, pCopy->srcDevice, pCopy->dstArray, srcOrigin, dstOrigin,
copyRegion, pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight,
queue);
stream);
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeDevice)) {
// Image to Device.
return ihipMemcpyAtoDCommand(command, pCopy->srcArray, pCopy->dstDevice, srcOrigin, dstOrigin,
copyRegion, pCopy->dstPitch, pCopy->dstPitch * pCopy->dstHeight,
queue);
stream);
} else if ((srcMemoryType == hipMemoryTypeArray) && (dstMemoryType == hipMemoryTypeArray)) {
// Image to Image.
return ihipMemcpyAtoACommand(command, pCopy->srcArray, pCopy->dstArray, srcOrigin, dstOrigin,
copyRegion, queue);
copyRegion, stream);
} else {
ShouldNotReachHere();
}
@@ -2212,14 +2212,14 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
// Host to Host.
return ihipMemcpyHtoH(pCopy->srcHost, pCopy->dstHost, srcOrigin, dstOrigin, copyRegion,
pCopy->srcPitch, pCopy->srcPitch * pCopy->srcHeight, pCopy->dstPitch,
pCopy->dstPitch * pCopy->dstHeight, hip::getQueue(stream));
pCopy->dstPitch * pCopy->dstHeight, hip::getStream(stream));
} else {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
status = ihipGetMemcpyParam3DCommand(command, pCopy, queue);
status = ihipGetMemcpyParam3DCommand(command, pCopy, hip_stream);
if (status != hipSuccess) return status;
// Transfers from device memory to pageable host memory and transfers from any host memory to any host memory
@@ -2507,13 +2507,13 @@ hipError_t ihipMemcpyAtoD(hipArray* srcArray, void* dstDevice, amd::Coord3D srcO
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch,
size_t dstSlicePitch, hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status =
ihipMemcpyAtoDCommand(command, srcArray, dstDevice, srcOrigin, dstOrigin, copyRegion,
dstRowPitch, dstSlicePitch, queue);
dstRowPitch, dstSlicePitch, hip_stream);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2521,13 +2521,13 @@ hipError_t ihipMemcpyDtoA(void* srcDevice, hipArray* dstArray, amd::Coord3D srcO
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch,
size_t srcSlicePitch, hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status =
ihipMemcpyDtoACommand(command, srcDevice, dstArray, srcOrigin, dstOrigin, copyRegion,
srcRowPitch, srcSlicePitch, queue);
srcRowPitch, srcSlicePitch, hip_stream);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2536,13 +2536,13 @@ hipError_t ihipMemcpyDtoD(void* srcDevice, void* dstDevice, amd::Coord3D srcOrig
size_t srcSlicePitch, size_t dstRowPitch, size_t dstSlicePitch,
hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status = ihipMemcpyDtoDCommand(command, srcDevice, dstDevice, srcOrigin, dstOrigin,
copyRegion, srcRowPitch, srcSlicePitch, dstRowPitch,
dstSlicePitch, queue);
dstSlicePitch, hip_stream);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2551,13 +2551,13 @@ hipError_t ihipMemcpyDtoH(void* srcDevice, void* dstHost, amd::Coord3D srcOrigin
size_t srcSlicePitch, size_t dstRowPitch, size_t dstSlicePitch,
hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status = ihipMemcpyDtoHCommand(command, srcDevice, dstHost, srcOrigin, dstOrigin,
copyRegion, srcRowPitch, srcSlicePitch, dstRowPitch,
dstSlicePitch, queue, isAsync);
dstSlicePitch, hip_stream, isAsync);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2566,13 +2566,13 @@ hipError_t ihipMemcpyHtoD(const void* srcHost, void* dstDevice, amd::Coord3D src
size_t srcSlicePitch, size_t dstRowPitch, size_t dstSlicePitch,
hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status = ihipMemcpyHtoDCommand(command, srcHost, dstDevice, srcOrigin, dstOrigin,
copyRegion, srcRowPitch, srcSlicePitch, dstRowPitch,
dstSlicePitch, queue, isAsync);
dstSlicePitch, hip_stream, isAsync);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2580,12 +2580,12 @@ hipError_t ihipMemcpyAtoA(hipArray* srcArray, hipArray* dstArray, amd::Coord3D s
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, hipStream_t stream,
bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status = ihipMemcpyAtoACommand(command, srcArray, dstArray, srcOrigin, dstOrigin,
copyRegion, queue);
copyRegion, hip_stream);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2593,13 +2593,13 @@ hipError_t ihipMemcpyHtoA(const void* srcHost, hipArray* dstArray, amd::Coord3D
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t srcRowPitch,
size_t srcSlicePitch, hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status =
ihipMemcpyHtoACommand(command, srcHost, dstArray, srcOrigin, dstOrigin, copyRegion,
srcRowPitch, srcSlicePitch, queue, isAsync);
srcRowPitch, srcSlicePitch, hip_stream, isAsync);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2607,13 +2607,13 @@ hipError_t ihipMemcpyAtoH(hipArray* srcArray, void* dstHost, amd::Coord3D srcOri
amd::Coord3D dstOrigin, amd::Coord3D copyRegion, size_t dstRowPitch,
size_t dstSlicePitch, hipStream_t stream, bool isAsync = false) {
amd::Command* command;
amd::HostQueue* queue = hip::getQueue(stream);
if (queue == nullptr) {
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipError_t status =
ihipMemcpyAtoHCommand(command, srcArray, dstHost, srcOrigin, dstOrigin, copyRegion,
dstRowPitch, dstSlicePitch, queue, isAsync);
dstRowPitch, dstSlicePitch, hip_stream, isAsync);
if (status != hipSuccess) return status;
return ihipMemcpyCmdEnqueue(command, isAsync);
}
@@ -2673,9 +2673,9 @@ hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p) {
}
hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p,
amd::HostQueue* queue) {
hip::Stream* stream) {
const HIP_MEMCPY3D desc = hip::getDrvMemcpy3DDesc(*p);
return ihipGetMemcpyParam3DCommand(command, &desc, queue);
return ihipGetMemcpyParam3DCommand(command, &desc, stream);
}
hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) {
@@ -2733,8 +2733,8 @@ hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream) {
hipError_t packFillMemoryCommand(amd::Command*& command, amd::Memory* memory, size_t offset,
int64_t value, size_t valueSize, size_t sizeBytes,
amd::HostQueue* queue) {
if ((memory == nullptr) || (queue == nullptr)) {
hip::Stream* stream) {
if ((memory == nullptr) || (stream == nullptr)) {
return hipErrorInvalidValue;
}
@@ -2744,7 +2744,7 @@ hipError_t packFillMemoryCommand(amd::Command*& command, amd::Memory* memory, si
// surface=[pitch, width, height]
amd::Coord3D surface(sizeBytes, sizeBytes, 1);
amd::FillMemoryCommand* fillMemCommand =
new amd::FillMemoryCommand(*queue, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(),
new amd::FillMemoryCommand(*stream, CL_COMMAND_FILL_BUFFER, waitList, *memory->asBuffer(),
&value, valueSize, fillOffset, fillSize, surface);
if (fillMemCommand == nullptr) {
return hipErrorOutOfMemory;
@@ -2810,7 +2810,7 @@ hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams) {
}
hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, int64_t value,
size_t valueSize, size_t sizeBytes, amd::HostQueue* queue) {
size_t valueSize, size_t sizeBytes, hip::Stream* stream) {
hipError_t hip_error = hipSuccess;
auto aligned_dst = amd::alignUp(reinterpret_cast<address>(dst), sizeof(uint64_t));
size_t offset = 0;
@@ -2820,7 +2820,7 @@ hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, in
amd::Command* command;
hip_error = packFillMemoryCommand(command, memory, offset, value, valueSize, sizeBytes,
queue);
stream);
commands.push_back(command);
return hip_error;
@@ -2854,8 +2854,8 @@ hipError_t ihipMemset(void* dst, int64_t value, size_t valueSize, size_t sizeByt
}
}
std::vector<amd::Command*> commands;
amd::HostQueue* queue = hip::getQueue(stream);
hip_error = ihipMemsetCommand(commands, dst, value, valueSize, sizeBytes, queue);
hip::Stream* hip_stream = hip::getStream(stream);
hip_error = ihipMemsetCommand(commands, dst, value, valueSize, sizeBytes, hip_stream);
if (hip_error != hipSuccess) {
break;
}
@@ -2972,13 +2972,13 @@ hipError_t ihipMemset3D_validate(hipPitchedPtr pitchedDevPtr, int value, hipExte
}
hipError_t ihipMemset3DCommand(std::vector<amd::Command*> &commands, hipPitchedPtr pitchedDevPtr,
int value, hipExtent extent, amd::HostQueue* queue, size_t elementSize = 1) {
int value, hipExtent extent, hip::Stream* stream, size_t elementSize = 1) {
size_t offset = 0;
auto sizeBytes = extent.width * extent.height * extent.depth;
amd::Memory* memory = getMemoryObject(pitchedDevPtr.ptr, offset);
if (pitchedDevPtr.pitch == extent.width) {
return ihipMemsetCommand(commands, pitchedDevPtr.ptr, value, elementSize,
static_cast<size_t>(sizeBytes), queue);
static_cast<size_t>(sizeBytes), stream);
}
// Workaround for cases when pitch > row until fill kernel will be updated to support pitch.
// Fall back to filling one row at a time.
@@ -2994,7 +2994,7 @@ hipError_t ihipMemset3DCommand(std::vector<amd::Command*> &commands, hipPitchedP
}
amd::FillMemoryCommand* command;
command = new amd::FillMemoryCommand(
*queue, CL_COMMAND_FILL_BUFFER, amd::Command::EventWaitList{}, *memory->asBuffer(),
*stream, CL_COMMAND_FILL_BUFFER, amd::Command::EventWaitList{}, *memory->asBuffer(),
&value, elementSize, origin, region, surface);
commands.push_back(command);
return hipSuccess;
@@ -3025,9 +3025,9 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent
isAsync = true;
}
}
amd::HostQueue* queue = hip::getQueue(stream);
hip::Stream* hip_stream = hip::getStream(stream);
std::vector<amd::Command*> commands;
status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, queue);
status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, hip_stream);
if (status != hipSuccess) {
return status;
}
@@ -3946,9 +3946,9 @@ hipError_t ihipMipmappedArrayDestroy(hipMipmappedArray_t mipmapped_array_ptr) {
}
for (auto& dev : g_devices) {
amd::HostQueue* queue = dev->NullStream(true);
if (queue != nullptr) {
queue->finish();
hip::Stream* stream = dev->NullStream(true);
if (stream != nullptr) {
stream->finish();
}
}