diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index 570ab35d46..fbd4ee6366 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -2334,8 +2334,10 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam, address parameters = captureArguments(kernels_[Scheduler]); - bool result = false; - result = gpu().submitKernelInternal(ndrange, *kernels_[Scheduler], parameters, nullptr); + if (!gpu().submitKernelInternal(ndrange, *kernels_[Scheduler], + parameters, nullptr)) { + return false; + } releaseArguments(parameters); if (hsa_signal_wait_acquire(schedulerSignal, HSA_SIGNAL_CONDITION_LT, 1, (-1), diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index e7b63a0c27..0222aaee18 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -1840,7 +1840,6 @@ amd::Memory *Device::IpcAttach(const void* handle, size_t mem_size, unsigned int /* FIX_ME (SWDEV-215976): Mapping for all devices is not optimal. In future map only for devices on need basis */ /* Retrieve the devPtr from the handle */ - hsa_agent_t hsa_agent = getBackendDevice(); hsa_status = hsa_amd_ipc_memory_attach(reinterpret_cast(handle), mem_size, (1 + p2p_agents_.size()), p2p_agents_list_, dev_ptr); diff --git a/projects/clr/rocclr/device/rocm/rockernel.cpp b/projects/clr/rocclr/device/rocm/rockernel.cpp index 62ee61afae..e7ff95adf1 100644 --- a/projects/clr/rocclr/device/rocm/rockernel.cpp +++ b/projects/clr/rocclr/device/rocm/rockernel.cpp @@ -45,9 +45,6 @@ Kernel::Kernel(std::string name, Program* prog) #if defined(USE_COMGR_LIBRARY) bool LightningKernel::init() { - - hsa_agent_t hsaDevice = program()->hsaDevice(); - if (!GetAttrCodePropMetadata()) { LogError("[ROC][Kernel] Could not get Code Prop Meta Data \n"); return false; diff --git a/projects/clr/rocclr/platform/program.cpp b/projects/clr/rocclr/platform/program.cpp index 24c621d10a..de33be9fe1 100755 --- a/projects/clr/rocclr/platform/program.cpp +++ b/projects/clr/rocclr/platform/program.cpp @@ -110,11 +110,13 @@ int32_t Program::addDeviceProgram(Device& device, const void* image, size_t leng if (devicePrograms_[&rootDev] != NULL) { return CL_SUCCESS; } - bool emptyOptions = false; + +#if defined(WITH_COMPILER_LIB) + bool emptyOptions = (options == nulltr); +#endif amd::option::Options emptyOpts; if (options == NULL) { options = &emptyOpts; - emptyOptions = true; } #if defined(WITH_COMPILER_LIB) @@ -333,11 +335,11 @@ int32_t Program::link(const std::vector& devices, size_t numInputs, continue; } inputDevPrograms[i] = findIt->second; - device::Program::binary_t binary = inputDevPrograms[i]->binary(); // Check the binary's target for the first found device program. // TODO: Revise these binary's target checks // and possibly remove them after switching to HSAIL by default. #if defined(WITH_COMPILER_LIB) + device::Program::binary_t binary = inputDevPrograms[i]->binary(); if (!found && binary.first != NULL && binary.second > 0 && aclValidateBinaryImage(binary.first, binary.second, BINARY_TYPE_ELF)) { acl_error errorCode = ACL_SUCCESS;