From a989c7c579017a0e449d52edc39317d7aac65674 Mon Sep 17 00:00:00 2001 From: jiabaxie Date: Sat, 11 Jul 2026 03:29:14 -0400 Subject: [PATCH 1/2] ROCM-27216 - ihipLaunchKernel catch all errors, except InvalidSymbol --- projects/clr/hipamd/src/hip_platform.cpp | 22 ++++++++++------------ 1 file changed, 10 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 8c13231d1fc2..1db05d158fc7 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/projects/clr/hipamd/src/hip_platform.cpp @@ -698,23 +698,21 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi const int deviceId = hip::Stream::DeviceId(stream); const auto [hip_error, func] = [&]() -> std::pair { - hipFunction_t f; + hipFunction_t f = nullptr; const hipError_t err = PlatformState::Instance().StatCO().GetFunc(&f, hostFunction, deviceId); - - // Propagate specific invalid code object errors - if (err == hipErrorInvalidKernelFile || - err == hipErrorInvalidDeviceFunction || - err == hipErrorInvalidImage) { - 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(const_cast(hostFunction))}; + + // Only take fallback for hipErrorInvalidSymbol (not in registered table) + if (err == hipErrorInvalidSymbol) { + return {hipSuccess, reinterpret_cast(const_cast(hostFunction))}; + } + + // Propagate all other errors + return {err != hipSuccess ? err : hipErrorInvalidDeviceFunction, nullptr}; }(); if (hip_error != hipSuccess) { From 59aafdbca3fd6d5de80c4c6b7bc87f4031b14561 Mon Sep 17 00:00:00 2001 From: Jimbo <57198431+jiabaxie@users.noreply.github.com> Date: Mon, 24 Aug 2026 17:05:10 -0400 Subject: [PATCH 2/2] Updated propagate all other errors option --- projects/clr/hipamd/src/hip_platform.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/src/hip_platform.cpp b/projects/clr/hipamd/src/hip_platform.cpp index 7e43bece28f5..cf2cf53486b6 100644 --- a/projects/clr/hipamd/src/hip_platform.cpp +++ b/projects/clr/hipamd/src/hip_platform.cpp @@ -732,7 +732,7 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi } // Propagate all other errors - return {err != hipSuccess ? err : hipErrorInvalidDeviceFunction, nullptr}; + return {err, nullptr}; }(); if (hip_error != hipSuccess) {