From d0b5bee6803f15466c198a38c36d2729c101b2b7 Mon Sep 17 00:00:00 2001 From: "Xie, Pengda" Date: Tue, 13 May 2025 13:45:27 -0700 Subject: [PATCH] SWDEV-527781 - Remove Stream Validation in HIP APIs [ROCm/clr commit: 0457b634f8f9db3e655fd9a6f38f38173dbb0dbf] --- projects/clr/hipamd/src/hip_event.cpp | 3 --- projects/clr/hipamd/src/hip_graph.cpp | 23 +++++---------------- projects/clr/hipamd/src/hip_hmm.cpp | 8 ++------ projects/clr/hipamd/src/hip_memory.cpp | 12 +++-------- projects/clr/hipamd/src/hip_mempool.cpp | 12 +++-------- projects/clr/hipamd/src/hip_peer.cpp | 4 +--- projects/clr/hipamd/src/hip_stream.cpp | 27 +++++++------------------ 7 files changed, 21 insertions(+), 68 deletions(-) diff --git a/projects/clr/hipamd/src/hip_event.cpp b/projects/clr/hipamd/src/hip_event.cpp index 62ed7c1dae..8bce85adbd 100644 --- a/projects/clr/hipamd/src/hip_event.cpp +++ b/projects/clr/hipamd/src/hip_event.cpp @@ -390,9 +390,6 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream, unsigned return hipErrorInvalidHandle; } getStreamPerThread(stream); - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } hip::Event* e = reinterpret_cast(event); hip::Stream* s = reinterpret_cast(stream); hip::Stream* hip_stream = hip::getStream(stream); diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index d31ef055e3..6835f62ec2 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -42,9 +42,7 @@ inline hipError_t ihipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream) if (graphExec == nullptr) { return hipErrorInvalidValue; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); if (!hip::GraphExec::isGraphExecValid(reinterpret_cast(graphExec))) { return hipErrorInvalidValue; } @@ -1026,9 +1024,7 @@ hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatu if (pCaptureStatus == nullptr) { return hipErrorInvalidValue; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); if (hip::Stream::StreamCaptureBlocking() == true && (stream == nullptr || stream == hipStreamLegacy)) { return hipErrorStreamCaptureImplicit; @@ -1069,9 +1065,7 @@ hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) { hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode, hipGraph_t graph = nullptr) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); // capture cannot be initiated on legacy stream if (stream == nullptr || stream == hipStreamLegacy) { return hipErrorStreamCaptureUnsupported; @@ -1591,9 +1585,7 @@ hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) { } hipError_t ihipGraphLaunch(hip::GraphExec* graphExec, hipStream_t stream) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); hip::Stream* launch_stream = hip::getStream(stream); return graphExec->Run(launch_stream); } @@ -1605,9 +1597,6 @@ hipError_t hipGraphLaunch_common(hip::GraphExec* graphExec, hipStream_t stream) if (graphExec->GetNodeCount() == 0) { return hipSuccess; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } return ihipGraphLaunch(graphExec, stream); } @@ -2004,9 +1993,7 @@ hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, if (pCaptureStatus == nullptr) { return hipErrorInvalidValue; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); if (hip::Stream::StreamCaptureBlocking() == true && (stream == nullptr || stream == hipStreamLegacy)) { return hipErrorStreamCaptureImplicit; diff --git a/projects/clr/hipamd/src/hip_hmm.cpp b/projects/clr/hipamd/src/hip_hmm.cpp index 95c4c44a6c..5f636d6b39 100644 --- a/projects/clr/hipamd/src/hip_hmm.cpp +++ b/projects/clr/hipamd/src/hip_hmm.cpp @@ -82,9 +82,7 @@ hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, HIP_RETURN(hipErrorInvalidValue); } - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); - } + getStreamPerThread(stream); size_t offset = 0; amd::Memory* memObj = getMemoryObject(dev_ptr, offset); @@ -238,9 +236,7 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr, HIP_RETURN(hipErrorInvalidValue); } - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); - } + getStreamPerThread(stream); if (flags != hipMemAttachGlobal && flags != hipMemAttachHost && flags != hipMemAttachSingle) { HIP_RETURN(hipErrorInvalidValue); diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4b99b27de0..a5144c3432 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1430,13 +1430,11 @@ hipError_t hipMemcpyAsync_common(void* dst, const void* src, size_t sizeBytes, if (static_cast(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) { return hipErrorInvalidMemcpyDirection; } + getStreamPerThread(stream); hip::Stream* hip_stream = hip::getStream(stream); if (hip_stream == nullptr) { return hipErrorInvalidValue; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } return ihipMemcpy(dst, src, sizeBytes, kind, *hip_stream, true); } @@ -2360,9 +2358,7 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool if (pCopy == nullptr) { return hipErrorInvalidValue; } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); hipMemoryType srcMemoryType; hipMemoryType dstMemoryType; ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType); @@ -2448,9 +2444,7 @@ hipError_t hipMemcpy2DValidateParams(hipMemcpyKind kind, hipStream_t stream = nu return hipErrorInvalidMemcpyDirection; } - if (!hip::isValid(stream)) { - return hipErrorInvalidValue; - } + getStreamPerThread(stream); return hipSuccess; } diff --git a/projects/clr/hipamd/src/hip_mempool.cpp b/projects/clr/hipamd/src/hip_mempool.cpp index 230622f945..ff7d164c70 100644 --- a/projects/clr/hipamd/src/hip_mempool.cpp +++ b/projects/clr/hipamd/src/hip_mempool.cpp @@ -85,9 +85,7 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) { if (dev_ptr == nullptr) { HIP_RETURN(hipErrorInvalidValue); } - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidHandle); - } + getStreamPerThread(stream); if (size == 0) { *dev_ptr = nullptr; HIP_RETURN(hipSuccess); @@ -147,9 +145,7 @@ class FreeAsyncCommand : public amd::Command { hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { HIP_INIT_API(hipFreeAsync, dev_ptr, stream); - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidHandle); - } + getStreamPerThread(stream); hip::Stream* s = reinterpret_cast(stream); auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? @@ -376,9 +372,7 @@ hipError_t hipMallocFromPoolAsync( if ((dev_ptr == nullptr) || (mem_pool == nullptr)) { HIP_RETURN(hipErrorInvalidValue); } - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorInvalidHandle); - } + getStreamPerThread(stream); if (size == 0) { *dev_ptr = nullptr; HIP_RETURN(hipSuccess); diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index b43c450530..6eb2ae0b59 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/projects/clr/hipamd/src/hip_peer.cpp @@ -209,9 +209,7 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src srcDevice < 0 || dstDevice < 0) { HIP_RETURN(hipErrorInvalidDevice); } - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); hip::Stream* hip_stream = hip::getStream(stream); if (hip_stream == nullptr) { return hipErrorInvalidValue; diff --git a/projects/clr/hipamd/src/hip_stream.cpp b/projects/clr/hipamd/src/hip_stream.cpp index b0ea00ec74..6fbd15a86c 100644 --- a/projects/clr/hipamd/src/hip_stream.cpp +++ b/projects/clr/hipamd/src/hip_stream.cpp @@ -323,9 +323,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio // ================================================================================================ hipError_t hipStreamGetFlags_common(hipStream_t stream, unsigned int* flags) { if ((flags != nullptr) && (stream != nullptr)) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); *flags = reinterpret_cast(stream)->Flags(); } else { return hipErrorInvalidValue; @@ -349,9 +347,7 @@ hipError_t hipStreamGetFlags_spt(hipStream_t stream, unsigned int* flags) { // ================================================================================================ hipError_t hipStreamSynchronize_common(hipStream_t stream) { - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); - } + getStreamPerThread(stream); if (stream != nullptr && stream != hipStreamLegacy) { // If still capturing return error if (hip::Stream::StreamCaptureOngoing(stream) == true) { @@ -398,9 +394,6 @@ hipError_t hipStreamDestroy(hipStream_t stream) { if (stream == hipStreamPerThread || stream == hipStreamLegacy) { HIP_RETURN(hipErrorInvalidResourceHandle); } - if (!hip::isValid(stream)) { - HIP_RETURN(hipErrorContextIsDestroyed); - } hip::Stream* s = reinterpret_cast(stream); if (s->GetCaptureStatus() != hipStreamCaptureStatusNone) { if (s->GetParentStream() != nullptr) { @@ -448,9 +441,10 @@ void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_d // ================================================================================================ hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsigned int flags) { hipError_t status = hipSuccess; - if (event == nullptr || !hip::isValid(stream)) { + if (event == nullptr) { return hipErrorInvalidHandle; } + getStreamPerThread(stream); hip::Stream* waitStream = hip::getStream(stream); hip::Event* e = reinterpret_cast(event); auto eventStreamHandle = reinterpret_cast(e->GetCaptureStream()); @@ -511,9 +505,7 @@ hipError_t hipStreamWaitEvent_spt(hipStream_t stream, hipEvent_t event, unsigned // ================================================================================================ hipError_t hipStreamQuery_common(hipStream_t stream) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); if (stream != nullptr) { // If still capturing return error if (hip::Stream::StreamCaptureOngoing(stream) == true) { @@ -566,10 +558,7 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) { } hipError_t streamCallback_common(hipStream_t stream, StreamCallback* cbo, void* userData) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } - + getStreamPerThread(stream); hip::Stream* hip_stream = hip::getStream(stream); amd::Command* last_command = hip_stream->getLastQueuedCommand(true); amd::Command::EventWaitList eventWaitList; @@ -688,9 +677,7 @@ hipError_t hipStreamGetPriority_common(hipStream_t stream, int* priority) { } if ((priority != nullptr) && (stream != nullptr)) { - if (!hip::isValid(stream)) { - return hipErrorContextIsDestroyed; - } + getStreamPerThread(stream); *priority = static_cast(reinterpret_cast(stream)->GetPriority()); } else { return hipErrorInvalidValue;