From ff23c5b71626cb31aed0e4fa87264b5b67ee9f13 Mon Sep 17 00:00:00 2001
From: German Andryeyev <German.Andryeyev@amd.com>
Date: Thu, 7 Jul 2022 13:05:35 -0400
Subject: [PATCH] SWDEV-311271 - Avoid a crash with the default queue

Move device access after the queue look-up

Change-Id: I799f23ff0299b20437c9a8edae42c38daa3e9ceb
---
 hipamd/src/hip_mempool.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/hipamd/src/hip_mempool.cpp b/hipamd/src/hip_mempool.cpp
index cc7ba10865..bb11708264 100644
--- a/hipamd/src/hip_mempool.cpp
+++ b/hipamd/src/hip_mempool.cpp
@@ -61,10 +61,10 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
   if ((dev_ptr == nullptr) || (size == 0) || (!hip::isValid(stream))) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  auto device = reinterpret_cast<hip::Stream*>(stream)->GetDevice();
-  auto mem_pool = device->GetCurrentMemoryPool();
   auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->GetNullStream() :
     reinterpret_cast<hip::Stream*>(stream);
+  auto device = hip_stream->GetDevice();
+  auto mem_pool = device->GetCurrentMemoryPool();
   *dev_ptr = reinterpret_cast<hip::MemoryPool*>(mem_pool)->AllocateMemory(size, hip_stream);
   HIP_RETURN(hipSuccess);
 }