SWDEV-506793 - segfault in ihipMallocManaged when no devices are available #122
Conversation
benrichard-amd
requested review from
chrispaquot,
gandryey,
saleelk,
mangupta and
rakesroy
as code owners
hipamd/src/hip_hmm.cpp
Outdated
| @@ -287,6 +287,10 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) { | |||
| assert((hip::host_context != nullptr) && "Current host context must be valid"); | |||
| amd::Context& ctx = *hip::host_context; | |||
|
|
|||
| if (ctx.devices().size() == 0) { | |||
There was a problem hiding this comment.
HIP_INIT_API macro already validates the active devices. I don't see why we need extra validation, unless something else is broken.
#define HIP_INIT_API(cid, ...)
HIP_INIT_API_INTERNAL(0, cid, VA_ARGS)
if (hip::g_devices.size() == 0) {
HIP_RETURN(hipErrorNoDevice);
}
There was a problem hiding this comment.
Just validate the devices in __hipRegisterManagedVar() itself.
void __hipRegisterManagedVar(
. . .
HIP_INIT_VOID();
hipError_t status = ihipMallocManaged(pointer, size, align);
There was a problem hiding this comment.
Moved device count check to __hipRegisterManagedVar().
A segfault was occurring immediately in HIP programs that uses static managed memory (i.e. globals with
__managed__modifier).This avoids a segfault in
ihipMallocManaged()when no devices are available. And because this happens during initialization before enteringmain(), we do not abort. Instead, the managed pointer is set to nullptr.