From a1c86d10b514ec11dd34fe8136e7a8df6422ca58 Mon Sep 17 00:00:00 2001 From: foreman Date: Thu, 24 Oct 2019 17:15:52 -0400 Subject: [PATCH] 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 | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/rocclr/runtime/device/rocm/rocprogram.cpp b/rocclr/runtime/device/rocm/rocprogram.cpp index 53f68dd7d8..f4a3d3b656 100644 --- a/rocclr/runtime/device/rocm/rocprogram.cpp +++ b/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) {