Skip to content

fix(clr) - elegantly return err when kernels are unsupported #10150

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
agunashe merged 4 commits into from
Aug 17, 2026
Merged

fix(clr) - elegantly return err when kernels are unsupported #10150

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
4ae913e
fix(clr) - elegently return err when kernels are unsupported
agunashe Aug 13, 2026
58bcfe0
check hipErrorNotSupported specifically
agunashe Aug 14, 2026
8530cc0
formatting errors
agunashe Aug 14, 2026
22ff79c
fix missing braces
agunashe Aug 14, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
12 changes: 4 additions & 8 deletions projects/clr/hipamd/src/hip_code_object.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -537,11 +537,7 @@ hipError_t StatCO::GetFunc(hipFunction_t* hfunc, const void* hostFunction, int d
if (module != nullptr) {
std::scoped_lock lock(sclock_);
if (*(module) == nullptr) {
hipError_t err = DigestFatBinary(module_to_hostModule_[module], *module);

if (err != hipSuccess) {
return err;
}
IHIP_RETURN_ONFAIL(DigestFatBinary(module_to_hostModule_[module], *module));
}
} else {
// Module was nullptr
Expand All @@ -563,7 +559,7 @@ hipError_t StatCO::GetFuncAttr(hipFuncAttributes* func_attr, const void* hostFun
// Lazy load
FatBinaryInfo** module = it->second->ModuleInfo();
if (*(module) == nullptr) {
std::ignore = DigestFatBinary(module_to_hostModule_[module], *module);
IHIP_RETURN_ONFAIL(DigestFatBinary(module_to_hostModule_[module], *module));
}

return it->second->GetStatFuncAttr(func_attr, deviceId);
Expand Down Expand Up @@ -594,7 +590,7 @@ hipError_t StatCO::GetGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_
// Lazy load
FatBinaryInfo** module = it->second->ModuleInfo();
if (*(module) == nullptr) {
std::ignore = DigestFatBinary(module_to_hostModule_[module], *module);
IHIP_RETURN_ONFAIL(DigestFatBinary(module_to_hostModule_[module], *module));
}

amd::Memory* mem = nullptr;
Expand Down Expand Up @@ -643,7 +639,7 @@ hipError_t StatCO::InitManagedVarDevicePtr(int deviceId) {
// Lazy load
FatBinaryInfo** module = var->ModuleInfo();
if (*(module) == nullptr) {
std::ignore = DigestFatBinary(module_to_hostModule_[module], *module);
IHIP_RETURN_ONFAIL(DigestFatBinary(module_to_hostModule_[module], *module));
}
hip::Stream* stream = g_devices.at(deviceId)->NullStream();
if (stream == nullptr) {
Expand Down
8 changes: 4 additions & 4 deletions projects/clr/hipamd/src/hip_platform.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -720,19 +720,19 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi
const auto [hip_error, func] = [&]() -> std::pair<hipError_t, hipFunction_t> {
hipFunction_t f;
const hipError_t err = PlatformState::Instance().StatCO().GetFunc(&f, hostFunction, deviceId);

// Propagate specific invalid code object errors
if (err == hipErrorInvalidKernelFile ||
err == hipErrorInvalidDeviceFunction ||
err == hipErrorInvalidImage) {
err == hipErrorInvalidImage ||
err == hipErrorNotSupported) { // ROCM_KPACK_ENABLED=OFF
return {err, nullptr};
}

// If successful lookup with valid function, use it
if (err == hipSuccess && f) {
return {hipSuccess, f};
}

// Fallback: assume it's a hip function type
return {hipSuccess, reinterpret_cast<hipFunction_t>(const_cast<void*>(hostFunction))};
}();
Expand Down
Loading