runtime: remember source modules that failed to compile - #1525
Open
pvelesko wants to merge 5 commits into
Open
Conversation
Collaborator
Author
|
/run-aurora-ci |
… every use A source module whose device program fails to build is compiled again by every later API call touching one of its symbols, and the failure is reported as hipErrorNotInitialized (OpenCL) or hipErrorTbd (Level Zero). The test drives a launch and hipMemcpyToSymbol at a module with an unresolved extern __device__ function and expects hipErrorSharedObjectInitFailed from both, a prompt second return, and a single module-cache lookup in the runtime log. Fails today with two lookups and hipErrorNotInitialized.
chipstar::Device::getOrCreateModule only recorded a source module on a successful compile, so a module whose device program failed to build or link was compiled again from scratch by every later API call touching one of its symbols. With Kokkos_CoreUnitTest_HIP on Aurora PVC that was an 11 minute compile-and-fail cycle per hipMemcpyToSymbol until the test timed out. Keep the CHIPError thrown by the first attempt in FailedSrcMods_ next to SrcModToCompiledMod_ and throw it again on every later request for the same source module. Report a failed program build or link as hipErrorSharedObjectInitFailed in both backends instead of hipErrorInitializationError (OpenCL) and hipErrorTbd (Level Zero). The Level Zero relink fallback no longer goes through the error table, whose missing zeModuleDynamicLink entry turned a link failure into abort(), and zeModuleCreate's ZE_RESULT_ERROR_MODULE_LINK_FAILURE gets a mapping. Fixes #1507
…Recompiled The test required hipErrorSharedObjectInitFailed from both calls, which is what OpenCL reports when clLinkProgram rejects the unresolved symbol. Level Zero builds the module inside zeModuleCreate and fails it there with ZE_RESULT_ERROR_MODULE_BUILD_FAILURE, which the runtime reports as hipErrorInvalidImage, so on Aurora PVC the test failed although the runtime fix worked: first call: hipErrorInvalidImage in 20.2 ms second call: hipErrorInvalidImage in 0.0 ms compile attempts logged: 1 FAILED: first call returned hipErrorInvalidImage, expected hipErrorSharedObjectInitFailed FAILED: second call returned hipErrorInvalidImage, expected hipErrorSharedObjectInitFailed Accept either code for the first call and require the second call to return the same code, keeping the single compile attempt and prompt second call assertions that are the point of the test.
pvelesko
force-pushed
the
2026-08-29-github-1507-remember-failed-modules
branch
from
August 29, 2026 17:13
12453f9 to
b4464cf
Compare
…edModuleNotRecompiled The error a broken module surfaces as depends on where the device compiler gives up. Intel OpenCL fails clBuildProgram (hipErrorSharedObjectInitFailed), Level Zero fails zeModuleCreate (hipErrorInvalidImage), and Mali builds the program with an empty log, silently dropping the kernel that calls the unresolved function, so the runtime only notices when it binds the source module's kernels to the program and reports hipErrorLaunchFailure from Module::getKernelByName. Both "run on Mali" lanes failed the test on that code alone; the property under test held: one compile attempt, second call back in 0.1 ms with the same error. Require the first call to fail with anything but hipSuccess and keep the checks that carry the fix: the second call returns the first call's error, promptly, and exactly one compile attempt is logged.
…FailedModuleNotRecompiled A device compiler that accepts the broken module (Mali does) stores a binary in the test's private cache directory, so rmdir() fails and every run leaks a /tmp/TestFix1507-cache.XXXXXX directory on the runner.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
When a module's device program failed to build or link,
chipstar::Device::getOrCreateModulerecorded nothing (only successes go intoSrcModToCompiledMod_), so every later API call touching a symbol of that module compiled it again from scratch; found withKokkos_CoreUnitTest_HIPon Aurora PVC, where eachhipMemcpyToSymbolrepeated an 11 minute compile that failed, until the test timed out. The first attempt's error is now kept inFailedSrcMods_next toSrcModToCompiledMod_and thrown again on every later request for the same source module, and a failed build or link is reported ashipErrorSharedObjectInitFailedin both backends instead ofhipErrorInitializationError(OpenCL) andhipErrorTbd(Level Zero). The Level Zero relink fallback no longer routes through the error table, whose missingzeModuleDynamicLinkentry turned a link failure intoabort().TestFix1507FailedModuleNotRecompileddrives a launch andhipMemcpyToSymbolat a module with an unresolvedextern __device__function and checks both error codes, the second call's latency, and that the runtime log holds a single module cache lookup (before the fix: two lookups andhipErrorNotInitialized).Fixes #1507