From 0f819eed6f758d8d2d61c87721a6fc4c192446bd Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 10 Jan 2018 10:17:06 -0500
Subject: [PATCH] P4 to Git Change 1502108 by wchau@wchau_OCL_boltzmann on
 2018/01/10 10:10:39

	SWDEV-142081 - Performance drop while running GE HealthCare - Penality test test due to CL#1498909
	- Port the logic to consider cache line size for the local workgroup detection (CL#1496286)

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#49 edit
---
 rocclr/runtime/device/rocm/rocvirtual.cpp | 35 +++++++++++++++++------
 1 file changed, 26 insertions(+), 9 deletions(-)

diff --git a/rocclr/runtime/device/rocm/rocvirtual.cpp b/rocclr/runtime/device/rocm/rocvirtual.cpp
index 58466e091a..947eaa4c27 100644
--- a/rocclr/runtime/device/rocm/rocvirtual.cpp
+++ b/rocclr/runtime/device/rocm/rocvirtual.cpp
@@ -1435,10 +1435,15 @@ void setRuntimeCompilerLocalSize(hsa_kernel_dispatch_packet_t& dispatchPacket,
           tmp /= div;
       }
 
+      // Assuming DWORD access
+      const uint cacheLineMatch = dev.info().globalMemCacheLineSize_ >> 2;
+
       // Check if partial dispatch is enabled and
       if (dev.settings().partialDispatch_ &&
           // we couldn't find optimal workload
-          (sizes.local().product() % devKernel->workGroupInfo()->wavefrontSize_) != 0) {
+          ((sizes.local().product() % devKernel->workGroupInfo()->wavefrontSize_) != 0 ||
+                // or size is too small for the cache line
+           (sizes.local()[0] < cacheLineMatch))) {
         size_t maxSize = 0;
         size_t maxDim = 0;
         for (uint d = 0; d < sizes.dimensions(); ++d) {
@@ -1447,18 +1452,30 @@ void setRuntimeCompilerLocalSize(hsa_kernel_dispatch_packet_t& dispatchPacket,
             maxDim = d;
           }
         }
-        // Check if a local workgroup has the most optimal size
-        if (thrPerGrp > maxSize) {
-          thrPerGrp = maxSize;
+
+        if ((maxDim != 0) && (sizes.global()[0] >= (cacheLineMatch / 2))) {
+          sizes.local()[0] = cacheLineMatch;
+          thrPerGrp /= cacheLineMatch;
+          sizes.local()[maxDim] = thrPerGrp;
+          for (uint d = 1; d < sizes.dimensions(); ++d) {
+            if (d != maxDim) {
+              sizes.local()[d] = 1;
+            }
+          }
         }
-        sizes.local()[maxDim] = thrPerGrp;
-        for (uint d = 0; d < sizes.dimensions(); ++d) {
-          if (d != maxDim) {
-            sizes.local()[d] = 1;
+        else {
+          // Check if a local workgroup has the most optimal size
+          if (thrPerGrp > maxSize) {
+            thrPerGrp = maxSize;
+          }
+          sizes.local()[maxDim] = thrPerGrp;
+          for (uint d = 0; d < sizes.dimensions(); ++d) {
+            if (d != maxDim) {
+              sizes.local()[d] = 1;
+            }
           }
         }
       }
-
       dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? sizes.local()[0] : 1;
       dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? sizes.local()[1] : 1;
       dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? sizes.local()[2] : 1;