P4 to Git Change 2019036 by gandryey@gera-win10 on 2019/10/24 17:06:27

SWDEV-208424 - ROCr language runtime should not free code object until executable destroy
	- Keep the code object reader alive until the program destruction. Update HSAIL path only, since LC path already handles it correctly.

Affected files ...

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

rocclr/runtime/device/rocm/rocprogram.cpp

@@ -278,8 +278,7 @@ bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_
   }
 
   // Load the code object.
-  hsa_code_object_reader_t codeObjectReader;
-  status = hsa_code_object_reader_create_from_memory(data, secSize, &codeObjectReader);
+  status = hsa_code_object_reader_create_from_memory(data, secSize, &hsaCodeObjectReader_);
   if (status != HSA_STATUS_SUCCESS) {
     buildLog_ += "Error: AMD HSA Code Object Reader create failed: ";
     buildLog_ += hsa_strerror(status);
@@ -288,7 +287,7 @@ bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_
   }
 
   hsa_agent_t hsaDevice = dev().getBackendDevice();
-  status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, codeObjectReader,
+  status = hsa_executable_load_agent_code_object(hsaExecutable_, hsaDevice, hsaCodeObjectReader_,
                                                  nullptr, nullptr);
   if (status != HSA_STATUS_SUCCESS) {
     buildLog_ += "Error: AMD HSA Code Object loading failed: ";
@@ -297,8 +296,6 @@ bool HSAILProgram::setKernels(amd::option::Options* options, void* binary, size_
     return false;
   }
 
-  hsa_code_object_reader_destroy(codeObjectReader);
-
   // Freeze the executable.
   status = hsa_executable_freeze(hsaExecutable_, nullptr);
   if (status != HSA_STATUS_SUCCESS) {