From c130d4ca6b0ef6293dac27cf3c3dcdddca72eae2 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 18 Oct 2017 11:39:33 -0400
Subject: [PATCH] P4 to Git Change 1471997 by lmoriche@lmoriche_opencl_dev2 on
 2017/10/18 11:30:15

	SWDEV-132888 - OCL kernel execution time is long at the first run on Linux, Vega, HSAIL (ProRender unit test)
	Don't finalize ISA binaries if recompilation is not requested.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#74 edit


[ROCm/clr commit: d53d7be4896abeb31495e5e57be53c562120c9da]
---
 .../rocclr/runtime/device/rocm/rocprogram.cpp | 29 ++++++++++---------
 1 file changed, 15 insertions(+), 14 deletions(-)

diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
index 7dcf218bb5..a0d60b24e1 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
@@ -1092,20 +1092,20 @@ bool HSAILProgram::linkImpl(amd::option::Options* options) {
   }
 
 #if !defined(WITH_LIGHTNING_COMPILER)
-  hsa_agent_t hsaDevice = dev().getBackendDevice();
-
-  std::string fin_options(options->origOptionStr);
-  // Append an option so that we can selectively enable a SCOption on CZ
-  // whenever IOMMUv2 is enabled.
-  if (dev().isFineGrainedSystem(true)) {
-    fin_options.append(" -sc-xnack-iommu");
-  }
-  errorCode = aclCompile(dev().compiler(), binaryElf_, fin_options.c_str(), ACL_TYPE_CG,
-                         ACL_TYPE_ISA, logFunction);
-  buildLog_ += aclGetCompilerLog(dev().compiler());
-  if (errorCode != ACL_SUCCESS) {
-    buildLog_ += "Error: BRIG finalization to ISA failed.\n";
-    return false;
+  if (finalize) {
+    std::string fin_options(options->origOptionStr);
+    // Append an option so that we can selectively enable a SCOption on CZ
+    // whenever IOMMUv2 is enabled.
+    if (dev().isFineGrainedSystem(true)) {
+      fin_options.append(" -sc-xnack-iommu");
+    }
+    errorCode = aclCompile(dev().compiler(), binaryElf_, fin_options.c_str(), ACL_TYPE_CG,
+                           ACL_TYPE_ISA, logFunction);
+    buildLog_ += aclGetCompilerLog(dev().compiler());
+    if (errorCode != ACL_SUCCESS) {
+      buildLog_ += "Error: BRIG finalization to ISA failed.\n";
+      return false;
+    }
   }
   size_t secSize;
   void* data =
@@ -1135,6 +1135,7 @@ bool HSAILProgram::linkImpl(amd::option::Options* options) {
     return false;
   }
 
+  hsa_agent_t hsaDevice = dev().getBackendDevice();
   status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, codeObjectReader,
                                                  nullptr, nullptr);
   if (status != HSA_STATUS_SUCCESS) {