diff --git a/projects/clr/hipamd/api/hip/hip_context.cpp b/projects/clr/hipamd/api/hip/hip_context.cpp
index d2c37c1b13..2ec1f9db50 100644
--- a/projects/clr/hipamd/api/hip/hip_context.cpp
+++ b/projects/clr/hipamd/api/hip/hip_context.cpp
@@ -75,6 +75,19 @@ void setCurrentContext(unsigned int index) {
   g_context = g_devices[index];
 }
 
+amd::HostQueue* getQueue(hipStream_t stream) {
+ if (stream == nullptr) {
+    syncStreams();
+    return getNullStream();
+  } else {
+    hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+    if ((s->flags & hipStreamNonBlocking) == 0) {
+      getNullStream()->finish();
+    }
+    return s->asHostQueue();
+  }
+}
+
 amd::HostQueue* getNullStream(amd::Context& context) {
   auto stream = g_nullStreams.find(&context);
   if (stream == g_nullStreams.end()) {
@@ -86,7 +99,6 @@ amd::HostQueue* getNullStream(amd::Context& context) {
     g_nullStreams[&context] = queue;
     return queue;
   }
-  syncStreams();
   return stream->second;
 }
 
diff --git a/projects/clr/hipamd/api/hip/hip_event.cpp b/projects/clr/hipamd/api/hip/hip_event.cpp
index 143d2eded9..81ae93e279 100644
--- a/projects/clr/hipamd/api/hip/hip_event.cpp
+++ b/projects/clr/hipamd/api/hip/hip_event.cpp
@@ -231,12 +231,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
 
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
 
-  amd::HostQueue* queue;
-  if (stream == nullptr) {
-    queue = hip::getNullStream();
-  } else {
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   amd::Command* command = queue->getLastQueuedCommand(true);
 
diff --git a/projects/clr/hipamd/api/hip/hip_internal.hpp b/projects/clr/hipamd/api/hip/hip_internal.hpp
index 49b3757dda..7b0ff09ed9 100644
--- a/projects/clr/hipamd/api/hip/hip_internal.hpp
+++ b/projects/clr/hipamd/api/hip/hip_internal.hpp
@@ -72,6 +72,7 @@ namespace hip {
   extern amd::Context* getCurrentContext();
   extern void setCurrentContext(unsigned int index);
 
+  extern amd::HostQueue* getQueue(hipStream_t s);
   extern amd::HostQueue* getNullStream(amd::Context&);
   extern amd::HostQueue* getNullStream();
   extern void syncStreams();
diff --git a/projects/clr/hipamd/api/hip/hip_memory.cpp b/projects/clr/hipamd/api/hip/hip_memory.cpp
index 0d9b1bec06..b1f1f40aec 100644
--- a/projects/clr/hipamd/api/hip/hip_memory.cpp
+++ b/projects/clr/hipamd/api/hip/hip_memory.cpp
@@ -95,6 +95,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
 
   if (((srcMemory == nullptr) && (dstMemory == nullptr)) ||
       (kind == hipMemcpyHostToHost)) {
+    queue.finish();
     memcpy(dst, src, sizeBytes);
     return hipSuccess;
   } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
@@ -169,17 +170,9 @@ hipError_t ihipMemset(void* dst, int value, size_t valueSize, size_t sizeBytes,
   }
 
   size_t offset = 0;
-  amd::HostQueue* queue = nullptr;
+  amd::HostQueue* queue = hip::getQueue(stream);
   amd::Memory* memory = getMemoryObject(dst, offset);
 
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
-
   if (memory != nullptr) {
     // Device memory
     amd::Command::EventWaitList waitList;
@@ -764,15 +757,7 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes,
                           hipMemcpyKind kind, hipStream_t stream) {
   HIP_INIT_API(dst, src, sizeBytes, kind, stream);
 
-  amd::HostQueue* queue = nullptr;
-
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, kind, *queue, true));
 }
@@ -782,15 +767,7 @@ hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes,
                               hipStream_t stream) {
   HIP_INIT_API(dst, src, sizeBytes, stream);
 
-  amd::HostQueue* queue = nullptr;
-
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyHostToDevice,
                     *queue, true));
@@ -800,15 +777,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
                               hipStream_t stream) {
   HIP_INIT_API(dst, src, sizeBytes, stream);
 
-  amd::HostQueue* queue = nullptr;
-
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToDevice,
                    *queue, true));
@@ -818,15 +787,7 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes,
                               hipStream_t stream) {
   HIP_INIT_API(dst, src, sizeBytes, stream);
 
-  amd::HostQueue* queue = nullptr;
-
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemcpy(reinterpret_cast<void*>(dst), (const void*) src, sizeBytes, hipMemcpyDeviceToHost,
                    *queue, true));
@@ -921,15 +882,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
                             size_t height, hipMemcpyKind kind, hipStream_t stream) {
   HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind, stream);
 
-  amd::HostQueue* queue;
-
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind, *queue, true));
 }
@@ -1349,14 +1302,7 @@ hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
                             size_t width, size_t height, hipStream_t stream) {
   HIP_INIT_API(dst, pitch, value, width, height, stream);
 
-  amd::HostQueue* queue = nullptr;
-  if (stream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(stream);
 
   HIP_RETURN(ihipMemset2D(dst, pitch, value, width, height, *queue, true));
 }
diff --git a/projects/clr/hipamd/api/hip/hip_module.cpp b/projects/clr/hipamd/api/hip/hip_module.cpp
index 542cd89551..0cd590886c 100644
--- a/projects/clr/hipamd/api/hip/hip_module.cpp
+++ b/projects/clr/hipamd/api/hip/hip_module.cpp
@@ -218,6 +218,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
                                  hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0,
                                  uint32_t params = 0)
 {
+  HIP_INIT_API(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
+    sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params);
+
   hip::Function* function = hip::Function::asFunction(f);
   amd::Kernel* kernel = function->function_;
   amd::Device* device = hip::getCurrentContext()->devices()[0];
@@ -226,14 +229,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
 
   hip::Event* eStart = reinterpret_cast<hip::Event*>(startEvent);
   hip::Event* eStop = reinterpret_cast<hip::Event*>(stopEvent);
-  amd::HostQueue* queue;
-  if (hStream == nullptr) {
-    hip::syncStreams();
-    queue = hip::getNullStream();
-  } else {
-    hip::getNullStream()->finish();
-    queue = reinterpret_cast<hip::Stream*>(hStream)->asHostQueue();
-  }
+  amd::HostQueue* queue = hip::getQueue(hStream);
+
   if ((params & amd::NDRangeKernelCommand::CooperativeGroups) &&
       !device->info().cooperativeGroups_) {
     return hipErrorLaunchFailure;
diff --git a/projects/clr/hipamd/api/hip/hip_stream.cpp b/projects/clr/hipamd/api/hip/hip_stream.cpp
index bc681d8a7e..ab26b4fc2e 100644
--- a/projects/clr/hipamd/api/hip/hip_stream.cpp
+++ b/projects/clr/hipamd/api/hip/hip_stream.cpp
@@ -173,21 +173,9 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) {
 hipError_t hipStreamSynchronize(hipStream_t stream) {
   HIP_INIT_API(stream);
 
-  amd::HostQueue* hostQueue;
+  amd::HostQueue* hostQueue = hip::getQueue(stream);
+  hostQueue->finish();
 
-  if (stream == nullptr) {
-    hip::syncStreams();
-
-    hostQueue = hip::getNullStream();
-
-    hostQueue->finish();
-  } else {
-    hip::getNullStream()->finish();
-
-    hip::Stream* hStream = reinterpret_cast<hip::Stream*>(stream);
-
-    hStream->finish();
-  }
   HIP_RETURN(hipSuccess);
 }