diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index 6289a9891c..107711d8ca 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -91,7 +91,7 @@ void setCurrentDevice(unsigned int index) { } hip::Stream* getStream(hipStream_t stream, bool wait) { - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { return getNullStream(wait); } else { hip::Stream* hip_stream = reinterpret_cast(stream); diff --git a/hipamd/src/hip_graph.cpp b/hipamd/src/hip_graph.cpp index 84cfc23ae7..c1d4d71043 100644 --- a/hipamd/src/hip_graph.cpp +++ b/hipamd/src/hip_graph.cpp @@ -855,10 +855,11 @@ hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatu if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } - if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) { + if (hip::Stream::StreamCaptureBlocking() == true && + (stream == nullptr || stream == hipStreamLegacy)) { return hipErrorStreamCaptureImplicit; } - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { *pCaptureStatus = hipStreamCaptureStatusNone; } else { *pCaptureStatus = reinterpret_cast(stream)->GetCaptureStatus(); @@ -898,7 +899,7 @@ hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode return hipErrorContextIsDestroyed; } // capture cannot be initiated on legacy stream - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { return hipErrorStreamCaptureUnsupported; } if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) { @@ -977,7 +978,7 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) { if (pGraph == nullptr) { return hipErrorInvalidValue; } - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { return hipErrorIllegalState; } if (!hip::isValid(stream)) { @@ -1771,10 +1772,11 @@ hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, if (!hip::isValid(stream)) { return hipErrorContextIsDestroyed; } - if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) { + if (hip::Stream::StreamCaptureBlocking() == true && + (stream == nullptr || stream == hipStreamLegacy)) { return hipErrorStreamCaptureImplicit; } - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { *pCaptureStatus = hipStreamCaptureStatusNone; return hipSuccess; } @@ -1807,10 +1809,11 @@ hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream, if (captureStatus_out == nullptr) { return hipErrorInvalidValue; } - if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) { + if (hip::Stream::StreamCaptureBlocking() == true && + (stream == nullptr || stream == hipStreamLegacy)) { return hipErrorStreamCaptureImplicit; } - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { *captureStatus_out = hipStreamCaptureStatusNone; return hipSuccess; } diff --git a/hipamd/src/hip_hmm.cpp b/hipamd/src/hip_hmm.cpp index c6e5733b2a..95c4c44a6c 100644 --- a/hipamd/src/hip_hmm.cpp +++ b/hipamd/src/hip_hmm.cpp @@ -108,10 +108,12 @@ hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, // Pick the specified stream or Null one from the provided device if (device == hipCpuDeviceId) { cpu_access = true; - hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() : hip::getStream(stream); + hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + hip::getCurrentDevice()->NullStream() : hip::getStream(stream); } else { dev = g_devices[device]->devices()[0]; - hip_stream = (stream == nullptr) ? g_devices[device]->NullStream() : hip::getStream(stream); + hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + g_devices[device]->NullStream() : hip::getStream(stream); } if (hip_stream == nullptr) { @@ -250,8 +252,8 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr, // host-accessible region of system-allocated pageable memory. // This type of memory may only be specified if the device associated with the // stream reports a non-zero value for the device attribute hipDevAttrPageableMemoryAccess. - hip::Stream* hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() - : hip::getStream(stream); + hip::Stream* hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + hip::getCurrentDevice()->NullStream() : hip::getStream(stream); size_t offset = 0; amd::Memory* memObj = getMemoryObject(dev_ptr, offset); if (memObj == nullptr) { diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp index 0ae7c393e4..fd9a61385f 100644 --- a/hipamd/src/hip_internal.hpp +++ b/hipamd/src/hip_internal.hpp @@ -230,7 +230,7 @@ const char* ihipGetErrorName(hipError_t hip_error); #define STREAM_CAPTURE(name, stream, ...) \ hip::getStreamPerThread(stream); \ - if (stream != nullptr && \ + if (stream != nullptr && stream != hipStreamLegacy && \ reinterpret_cast(stream)->GetCaptureStatus() == \ hipStreamCaptureStatusActive) { \ hipError_t status = hip::capture##name(stream, ##__VA_ARGS__); \ @@ -242,7 +242,7 @@ const char* ihipGetErrorName(hipError_t hip_error); } #define PER_THREAD_DEFAULT_STREAM(stream) \ - if (stream == nullptr) { \ + if (stream == nullptr || stream == hipStreamLegacy) { \ stream = getPerThreadDefaultStream(); \ } diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index edb7178e76..7bcc1ba48a 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -674,7 +674,7 @@ hipError_t hipMemcpy_common(void* dst, const void* src, size_t sizeBytes, CHECK_STREAM_CAPTURING(); hip::Stream* hip_stream = nullptr; - if (stream != nullptr) { + if (stream != nullptr && stream != hipStreamLegacy) { hip_stream = hip::getStream(stream); } else { hip_stream = hip::getNullStream(); diff --git a/hipamd/src/hip_mempool.cpp b/hipamd/src/hip_mempool.cpp index 87bc3e41ab..c7f302cbf7 100644 --- a/hipamd/src/hip_mempool.cpp +++ b/hipamd/src/hip_mempool.cpp @@ -92,8 +92,8 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) { *dev_ptr = nullptr; HIP_RETURN(hipSuccess); } - auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() : - reinterpret_cast(stream); + auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + hip::getCurrentDevice()->NullStream() : reinterpret_cast(stream); auto device = hip_stream->GetDevice(); auto mem_pool = device->GetCurrentMemoryPool(); @@ -147,8 +147,8 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { STREAM_CAPTURE(hipFreeAsync, stream, dev_ptr); - auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() - : reinterpret_cast(stream); + auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + hip::getCurrentDevice()->NullStream(): reinterpret_cast(stream); hip::Event* event = nullptr; bool graph_in_use = false; @@ -192,7 +192,7 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) { } } } - + auto cmd = new FreeAsyncCommand(*hip_stream, dev_ptr, event); if (cmd == nullptr) { HIP_RETURN(hipErrorUnknown); @@ -367,8 +367,8 @@ hipError_t hipMallocFromPoolAsync( STREAM_CAPTURE(hipMallocAsync, stream, mem_pool, size, dev_ptr); auto mpool = reinterpret_cast(mem_pool); - auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() : - reinterpret_cast(stream); + auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ? + hip::getCurrentDevice()->NullStream() : reinterpret_cast(stream); *dev_ptr = mpool->AllocateMemory(size, hip_stream); if (*dev_ptr == nullptr) { HIP_RETURN(hipErrorOutOfMemory); diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index f73a9275cd..7152ed1ba2 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -80,7 +80,7 @@ void Stream::Destroy(hip::Stream* stream) { // ================================================================================================ bool isValid(hipStream_t& stream) { // NULL stream is always valid - if (stream == nullptr) { + if (stream == nullptr || stream == hipStreamLegacy) { return true; }