OpenCL: fail the module when the built program lacks a source kernel - #1541
Draft
pvelesko wants to merge 2 commits into
Draft
OpenCL: fail the module when the built program lacks a source kernel#1541pvelesko wants to merge 2 commits into
pvelesko wants to merge 2 commits into
Conversation
…ernel On Mali the OpenCL driver accepts a SPIR-V module with an unresolved __device__ function import, returns CL_SUCCESS from clBuildProgram with an empty log, and leaves the kernel that calls the import out of the program. The runtime then fails in the kernel binding of Device::getOrCreateModule with hipErrorLaunchFailure and the module is never remembered as failed. #1539
The Arm Mali driver accepts a SPIR-V module with an unresolved function import: clBuildProgram returns CL_SUCCESS with an empty build log, but clCreateKernelsInProgram returns a program without the kernels that call the import. The module was registered as compiled and the kernel binding in Device::getOrCreateModule then failed with hipErrorLaunchFailure on every launch, and since compile() had not thrown, the module could never be remembered as failed. After clCreateKernelsInProgram, compare the program's kernels with the source module's registered kernels and throw hipErrorSharedObjectInitFailed naming the missing ones, the same path a rejected build takes. The program is stored in the module cache only after that check passes. #1539
Collaborator
Author
|
/run-aurora-ci |
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.
Fixes #1539
On Mali the driver accepts a SPIR-V module with an unresolved
__device__function import:clBuildProgramreturnsCL_SUCCESSwith an empty log but the kernel calling the import is left out of the program, so the runtime registered the module as compiled and then failed in the kernel binding withhipErrorLaunchFailureon every launch. AfterclCreateKernelsInProgram,CHIPModuleOpenCL::compilenow compares the program's kernels with the source module's registered kernels and throwshipErrorSharedObjectInitFailednaming the missing ones, the same path a rejected build takes, so the module gets remembered as failed by #1525 (this complements that PR). The program is stored in the module cache only after the check passes.Validated on Mali G52 (salami: both launches now return
hipErrorSharedObjectInitFailed, warm and cold module cache) and on the Intel CPU OpenCL runtime (unchanged,clBuildProgramrejects the module). New test:tests/runtime/TestFix1539MissingKernelAfterBuild.hip.