[SYCL] Static linking support in Level Zero plugin #5266
Merged
Conversation
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
The Level Zero driver now has "static linking" support, the ability to create a single Level Zero module from a list of SPIR-V modules. Use this feature to implement `piProgramLink`, replacing the previous implementation that relied on "dynamic linking" of Level Zero modules. There are two main improvements: * Static linking provides more optimizations because the compiler sees the SPIR-V for all modules, which enables cross-module optimizations. * The previous implementation was mostly using mock driver APIs, and it didn't work in the general case when `piProgramLink` tried to link more than one input module together. This is mostly an infrastructure improvement, which does not provide any new features or bug fixes. Since there is no Level Zero interop API that can create a `kernel_bundle` in `object` or `input` state, there is no need currently to link multiple Level Zero modules together with `piProgramLink`. However, this commit provides the infrastructure we need to enable online linking of the "device library" into device code. That feature will be enabled in a future commit.
The previous commit adds new logic to the Level Zero implementation of `piProgramLink`, but this logic is only used currently if the SYCL application explicitly calls `sycl::link`. When an application merely executes a kernel, the SYCL runtime generally calls `piProgramBuild` instead of `piProgramCompile` / `piProgramLink`. In order to enable wider testing of the new `piProgramLink` logic, this commit adds a temporary environment variable which forces the SYCL runtime to build all programs by calling `piProgramCompile` / `piProgramLink`. This is usually safe, though there are a few cases where it is unsafe and results in false errors. Therefore, use this environment variable with caution. Usage is like: ``` SYCL_FORCE_LINK=[0|1] ``` When running tests with `llvm-lit`, you can set this environment variable by passing `--param extra_environment=SYCL_FORCE_LINK=1`: ``` $ llvm-lit --param extra_environment=SYCL_FORCE_LINK=1 \ --param sycl_be=level_zero --param target_devices=gpu \ -j12 -v -s <test> ``` Here are the known cases when specifying `SYCL_FORCE_LINK=1` results in a false error: * Tests using AOT (ahead-of-time compilation) generally fail because a `pi_program` that is created from native code cannot be compiled with `piProgramCompile`. * Some tests enable tracing with `SYCL_PI_TRACE` and explicitly check for a call to `piProgramBuild`. These tests fail because the runtime calls `piProgramCompile` / `piProgramLink` instead of calling `piProgramBuild`.
|
@againull: @smaslov-intel thought you would be interested in reviewing the thread locking logic in this PR. See the use of |
smaslov-intel
approved these changes
Jan 6, 2022
There was a problem hiding this comment.
LGTM
@againull : please give your explicit OK for program mutex usage here
againull
approved these changes
Jan 7, 2022
bader
reviewed
Jan 10, 2022
| @@ -963,8 +963,11 @@ ProgramManager::ProgramPtr ProgramManager::build( | |||
| DeviceLibReqMask); | |||
| } | |||
|
|
|||
| static const char *ForceLinkEnv = std::getenv("SYCL_FORCE_LINK"); | |||
There was a problem hiding this comment.
@againull, @smaslov-intel, shouldn't this variable be documented in https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md#debugging-variables-for-dpc-runtime section?
There was a problem hiding this comment.
I added this environment variable to help the device driver team reproduce two bugs that I reported. Since I expect to remove the environment variable after those bugs are fixed, I didn't bother adding documentation. If people want, I can add documentation now, but we will just remove it again once the bugs are fixed.
There was a problem hiding this comment.
I am OK to add this to experimental section
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
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.
The Level Zero driver now has "static linking" support, the ability to
create a single Level Zero module from a list of SPIR-V modules. Use
this feature to implement
piProgramLink, replacing the previousimplementation that relied on "dynamic linking" of Level Zero modules.
There are two main improvements:
Static linking provides more optimizations because the compiler sees
the SPIR-V for all modules, which enables cross-module optimizations.
The previous implementation was mostly using mock driver APIs, and it
didn't work in the general case when
piProgramLinktried to linkmore than one input module together.
This is mostly an infrastructure improvement, which does not provide
any new features or bug fixes. Since there is no Level Zero interop
API that can create a
kernel_bundleinobjectorinputstate,there is no need currently to link multiple Level Zero modules together
with
piProgramLink.However, this commit provides the infrastructure we need to enable
online linking of the "device library" into device code. That feature
will be enabled in a future commit.
Note that a temporary environment variable has also been added
(
SYCL_FORCE_LINK=[0|1]), which was added to help test this change.Defining this variable to
1forces the runtime to callpiProgramCompile/piProgramLinkinstead ofpiProgramBuildevenif the program is created from a single SPIR-V module. This is usually
safe, but it sometimes results in false errors, so only enable this
variable if you know what you are doing. We expect support for this
environment variable to removed soon.