SWDEV-193956

[hipclang-vdi-rocm][perf]~45% to 50% of Performance drop on
rocBLAS_int8 test

- Enable AMD_OPT_FLUSH optimization by default to match HCC
- Disable CPU writes to GPU memory on boards with large bar,
because it requires HDP flush tracking.
- Enable L2 cache on kernel arguments, because L2 will be
invalidated on memory reuse .

Change-Id: I124cf250bdd4d19c523ce542c163813828f8fbdc

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;
 

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;
   }

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,                                         \