diff --git a/rocclr/device/rocm/rocprogram.cpp b/rocclr/device/rocm/rocprogram.cpp index ef9156d798..977ce3442f 100644 --- a/rocclr/device/rocm/rocprogram.cpp +++ b/rocclr/device/rocm/rocprogram.cpp @@ -197,32 +197,35 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr, return false; } - /* Find HSA Symbol Address */ - status = hsa_executable_symbol_get_info(global_symbol, + // Handle size 0 symbols + if (*bytes != 0) { + // Find HSA Symbol Address + status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, device_pptr); - if (status != HSA_STATUS_SUCCESS) { - buildLog_ += "Error: Failed to find the Symbol Address : "; - buildLog_ += hsa_strerror(status); - buildLog_ += "\n"; - return false; + if (status != HSA_STATUS_SUCCESS) { + buildLog_ += "Error: Failed to find the Symbol Address : "; + buildLog_ += hsa_strerror(status); + buildLog_ += "\n"; + return false; + } + + roc_device = static_cast(&dev()); + *amd_mem_obj = new(roc_device->context()) amd::Buffer(roc_device->context(), 0, *bytes, + *device_pptr); + + if (*amd_mem_obj == nullptr) { + buildLog_ += "[OCL] Failed to create a mem object!"; + buildLog_ += "\n"; + return false; + } + + if (!((*amd_mem_obj)->create(nullptr))) { + buildLog_ += "[OCL] failed to create a svm hidden buffer!"; + buildLog_ += "\n"; + (*amd_mem_obj)->release(); + return false; + } } - - roc_device = static_cast(&dev()); - *amd_mem_obj = new(roc_device->context()) amd::Buffer(roc_device->context(), 0, *bytes, *device_pptr); - - if (*amd_mem_obj == nullptr) { - buildLog_ += "[OCL] Failed to create a mem object!"; - buildLog_ += "\n"; - return false; - } - - if (!((*amd_mem_obj)->create(nullptr))) { - buildLog_ += "[OCL] failed to create a svm hidden buffer!"; - buildLog_ += "\n"; - (*amd_mem_obj)->release(); - return false; - } - return true; }