diff --git a/projects/hip/api/hip/hip_context.cpp b/projects/hip/api/hip/hip_context.cpp index dfeacdc924..8559b641fc 100644 --- a/projects/hip/api/hip/hip_context.cpp +++ b/projects/hip/api/hip/hip_context.cpp @@ -24,8 +24,6 @@ THE SOFTWARE. #include "hip_internal.hpp" #include "platform/runtime.hpp" #include "utils/versions.hpp" -#include -#include std::vector g_devices; @@ -35,7 +33,7 @@ thread_local amd::Context* g_context = nullptr; thread_local std::stack g_ctxtStack; std::once_flag g_ihipInitialized; -std::map g_nullStreams; +std::map g_nullStreams; void init() { if (!amd::Runtime::initialized()) { @@ -66,6 +64,12 @@ void setCurrentContext(unsigned int index) { g_context = g_devices[index]; } +void syncStreams() { + for (const auto& it : streamSet) { + it->finish(); + } +} + amd::HostQueue* getNullStream() { auto stream = g_nullStreams.find(getCurrentContext()); if (stream == g_nullStreams.end()) { @@ -76,6 +80,7 @@ amd::HostQueue* getNullStream() { g_nullStreams[getCurrentContext()] = queue; return queue; } + syncStreams(); return stream->second; } diff --git a/projects/hip/api/hip/hip_internal.hpp b/projects/hip/api/hip/hip_internal.hpp index 099aa6ca25..1f97be451a 100644 --- a/projects/hip/api/hip/hip_internal.hpp +++ b/projects/hip/api/hip/hip_internal.hpp @@ -24,8 +24,9 @@ THE SOFTWARE. #define HIP_SRC_HIP_INTERNAL_H #include "cl_common.hpp" - +#include #include +#include #define HIP_INIT() \ std::call_once(hip::g_ihipInitialized, hip::init); \ @@ -57,9 +58,10 @@ namespace hip { extern void setCurrentContext(unsigned int index); extern amd::HostQueue* getNullStream(); + extern void syncStreams(); }; extern std::vector g_devices; - +extern thread_local std::unordered_set streamSet; extern hipError_t ihipDeviceGetCount(int* count); #endif // HIP_SRC_HIP_INTERNAL_H diff --git a/projects/hip/api/hip/hip_stream.cpp b/projects/hip/api/hip/hip_stream.cpp index 6c2c29fb51..de6083a856 100644 --- a/projects/hip/api/hip/hip_stream.cpp +++ b/projects/hip/api/hip/hip_stream.cpp @@ -21,18 +21,22 @@ THE SOFTWARE. */ #include - #include "hip_internal.hpp" -static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { - assert(flags == 0); // we don't handle flags yet +thread_local std::unordered_set streamSet; +static hipError_t ihipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { amd::Device* device = hip::getCurrentContext()->devices()[0]; amd::HostQueue* queue = new amd::HostQueue(*hip::getCurrentContext(), *device, 0, amd::CommandQueue::RealTimeDisabled, amd::CommandQueue::Priority::Normal); + if (!(flags & hipStreamNonBlocking)) { + hip::syncStreams(); + streamSet.insert(queue); + } + if (queue == nullptr) { return hipErrorOutOfMemory; } @@ -48,23 +52,27 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) { return ihipStreamCreateWithFlags(stream, flags); } - hipError_t hipStreamCreate(hipStream_t *stream) { HIP_INIT_API(stream); return ihipStreamCreateWithFlags(stream, hipStreamDefault); } - hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { HIP_INIT_API(stream, flags); - assert(0 && "Unimplemented"); + amd::HostQueue* hostQueue = reinterpret_cast(stream); + auto it = streamSet.find(hostQueue); - return hipErrorUnknown; + if(flags != nullptr) { + *flags = (it != streamSet.end()) ? hipStreamNonBlocking : hipStreamDefault; + } else { + return hipErrorInvalidValue; + } + + return hipSuccess; } - hipError_t hipStreamSynchronize(hipStream_t stream) { HIP_INIT_API(stream); @@ -85,7 +93,6 @@ hipError_t hipStreamSynchronize(hipStream_t stream) { return hipSuccess; } - hipError_t hipStreamDestroy(hipStream_t stream) { HIP_INIT_API(stream); @@ -93,12 +100,14 @@ hipError_t hipStreamDestroy(hipStream_t stream) { return hipErrorInvalidResourceHandle; } + amd::HostQueue* hostQueue = reinterpret_cast(stream); + streamSet.erase(hostQueue); + as_amd(reinterpret_cast(stream))->release(); return hipSuccess; } - hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_API(stream, event, flags);