[SYCL] Separate host instantiation from HostKernel #18534
Merged
+73
−32
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 HostKernel class is currently used for two different things: 1) To extend the lifetime of the kernel function object/lambda; and 2) To force the kernel to be instantiated on the host. Requiring a single solution to solve both problems prevents optimization. Providing separate solutions will enable us to use the fastest solution for each problem. Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Signed-off-by: John Pennycook <[email protected]>
Apparently the symbol is different for different versions of MSVC... Signed-off-by: John Pennycook <[email protected]>
sergey-semenov
approved these changes
May 28, 2025
|
@intel/llvm-gatekeepers - I think this can be merged now. |
|
When I tried to do this, instantiating wasn't enough, I had to ensure the code isn't DCE'd out (which is why I ultimately kept it). Have you tried to actually run all our tests in preview mode? Update: I see that |
| // This function is needed for host-side compilation to keep kernels | ||
| // instantitated. This is important for debuggers to be able to associate | ||
| // kernel code instructions with source code lines. | ||
| // NOTE: InstatiateKernelOnHost() should not be called. | ||
| void InstantiateKernelOnHost() override { | ||
| using IDBuilder = sycl::detail::Builder; | ||
| constexpr bool HasKernelHandlerArg = |
There was a problem hiding this comment.
Why can't we just call "new" instantiation here?
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 HostKernel class is currently used for two different things:
Requiring a single solution to solve both problems prevents optimization. Providing separate solutions will enable us to use the fastest solution for each problem.
A few notes that might help with review:
I would like us to eventually reach a point where we can remove the
HostKernelclass entirely and avoid extending the lifetime of the kernel function object. We can't do that if it's being relied upon for debugging.I had to change the
runKernelWithoutArgsignature to remove SFINAE because MSVC complained about creating a pointer to an overloaded function. I would have expected that the compiler could figure out the function pointer after SFINAE, but it failed.Removing the old mechanism of instantiating on the host is an ABI break because it removes a function from the vtable.
With the new approach, we can avoid instantiating (and compiling!) the kernel on the host under
NDEBUG. I've verified that-DNDEBUG -fpreview-breaking-changesremoves the kernel symbols from the binary.-DNDEBUGalone isn't currently sufficient, because the old instantiation mechanism is still triggered. I wasn't sure how much of the old implementation to leave in tact.