diff --git a/rocclr/device/rocm/rocmemory.cpp b/rocclr/device/rocm/rocmemory.cpp
index ab91f2046a..5dab89b17d 100644
--- a/rocclr/device/rocm/rocmemory.cpp
+++ b/rocclr/device/rocm/rocmemory.cpp
@@ -707,7 +707,9 @@ bool Buffer::create() {
     if (dev().forceFineGrain(owner()) ||
         dev().isFineGrainedSystem(true)) {
       memFlags |= CL_MEM_SVM_FINE_GRAIN_BUFFER;
-      flags_ |= HostMemoryDirectAccess;
+      // Don't enable direct access to GPU memory with large bar, because
+      // there is no tracking of HDP flush after CPU writes
+      // flags_ |= HostMemoryDirectAccess;
     }
     const bool isFineGrain = memFlags & CL_MEM_SVM_FINE_GRAIN_BUFFER;
 
diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp
index 530451cdeb..2a18f97d8e 100644
--- a/rocclr/device/rocm/rocvirtual.cpp
+++ b/rocclr/device/rocm/rocvirtual.cpp
@@ -452,7 +452,9 @@ bool VirtualGPU::dispatchGenericAqlPacket(
   }
 
   // Make sure the slot is free for usage
-  while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask);
+  while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) {
+    amd::Os::yield();
+  }
 
   // Add blocking command if the original value of read index was behind of the queue size
   if (blocking || (index - read) >= queueMask) {
@@ -485,9 +487,6 @@ bool VirtualGPU::dispatchGenericAqlPacket(
       LogPrintfError("Failed signal [0x%lx] wait", signal.handle);
       return false;
     }
-
-    // Release the pool, since runtime just drained the entire queue
-    resetKernArgPool();
   }
 
   return true;
@@ -750,7 +749,7 @@ bool VirtualGPU::create(bool profilingEna) {
 
 bool VirtualGPU::initPool(size_t kernarg_pool_size, uint signal_pool_count) {
   kernarg_pool_size_ = kernarg_pool_size;
-  kernarg_pool_base_ = reinterpret_cast<char*>(roc_device_.hostAlloc(kernarg_pool_size_, 1));
+  kernarg_pool_base_ = reinterpret_cast<char*>(roc_device_.hostAlloc(kernarg_pool_size_, false));
   if (kernarg_pool_base_ == nullptr) {
     return false;
   }
diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp
index a37c0a27c9..ee2dc77c26 100644
--- a/rocclr/utils/flags.hpp
+++ b/rocclr/utils/flags.hpp
@@ -226,7 +226,7 @@ release(bool, PAL_ALWAYS_RESIDENT, false,                                     \
 release(uint, HIP_HOST_COHERENT, 0,                                           \
         "Coherent memory in hipHostMalloc, 0x1 = memory is coherent with host"\
         "0x0 = memory is not coherent between host and GPU")                  \
-release(uint, AMD_OPT_FLUSH, 0,                                               \
+release(uint, AMD_OPT_FLUSH, 1,                                               \
         "Kernel flush option , 0x0 = Use system-scope fence operations."      \
         "0x1 = Use device-scope fence operations when possible.")             \
 release(uint, HIP_HIDDEN_FREE_MEM, 0,                                         \