-
Notifications
You must be signed in to change notification settings - Fork 787
[SYCL] Fix a race condition when enqueueing an interop kernel #8111
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Merged
Merged
Changes from all commits
Commits
Show all changes
3 commits
Select commit
Hold shift + click to select a range
File filter
Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
There are no files selected for viewing
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
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
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,126 @@ | ||
| //==-- KernelInteropCommon.hpp --- Common kernel interop redefinitions -----==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include <helpers/PiMock.hpp> | ||
sergey-semenov marked this conversation as resolved.
Show resolved
Hide resolved
|
||
|
|
||
| struct TestContext { | ||
|
|
||
| // SYCL RT has number of checks that all devices and contexts are consistent | ||
| // between kernel, kernel_bundle and other objects. | ||
| // | ||
| // To ensure that those checks pass, we intercept some PI calls to extract | ||
| // the exact PI handles of device and context used in queue creation to later | ||
| // return them when program/context/kernel info is requested. | ||
| pi_device deviceHandle; | ||
| pi_context contextHandle; | ||
|
|
||
| pi_program programHandle = createDummyHandle<pi_program>(); | ||
|
|
||
| ~TestContext() { releaseDummyHandle<pi_program>(programHandle); } | ||
| }; | ||
|
|
||
| TestContext GlobalContext; | ||
|
|
||
| pi_result after_piContextGetInfo(pi_context context, pi_context_info param_name, | ||
| size_t param_value_size, void *param_value, | ||
| size_t *param_value_size_ret) { | ||
| switch (param_name) { | ||
| case PI_CONTEXT_INFO_DEVICES: | ||
| if (param_value) | ||
| *static_cast<pi_device *>(param_value) = GlobalContext.deviceHandle; | ||
| if (param_value_size_ret) | ||
| *param_value_size_ret = sizeof(GlobalContext.deviceHandle); | ||
| break; | ||
| default:; | ||
| } | ||
|
|
||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| pi_result after_piProgramGetInfo(pi_program program, pi_program_info param_name, | ||
| size_t param_value_size, void *param_value, | ||
| size_t *param_value_size_ret) { | ||
|
|
||
| switch (param_name) { | ||
| case PI_PROGRAM_INFO_DEVICES: | ||
| if (param_value_size_ret) | ||
| *param_value_size_ret = sizeof(GlobalContext.deviceHandle); | ||
| if (param_value) | ||
| *static_cast<pi_device *>(param_value) = GlobalContext.deviceHandle; | ||
| break; | ||
| default:; | ||
| } | ||
|
|
||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| pi_result redefined_piProgramGetBuildInfo(pi_program program, pi_device device, | ||
| _pi_program_build_info param_name, | ||
| size_t param_value_size, | ||
| void *param_value, | ||
| size_t *param_value_size_ret) { | ||
| switch (param_name) { | ||
| case PI_PROGRAM_BUILD_INFO_BINARY_TYPE: | ||
| if (param_value_size_ret) | ||
| *param_value_size_ret = sizeof(pi_program_binary_type); | ||
| if (param_value) | ||
| *static_cast<pi_program_binary_type *>(param_value) = | ||
| PI_PROGRAM_BINARY_TYPE_EXECUTABLE; | ||
| break; | ||
| default:; | ||
| } | ||
|
|
||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| pi_result after_piContextCreate(const pi_context_properties *properties, | ||
| pi_uint32 num_devices, const pi_device *devices, | ||
| void (*pfn_notify)(const char *errinfo, | ||
| const void *private_info, | ||
| size_t cb, void *user_data), | ||
| void *user_data, pi_context *ret_context) { | ||
| if (ret_context) | ||
| GlobalContext.contextHandle = *ret_context; | ||
| GlobalContext.deviceHandle = *devices; | ||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| pi_result after_piKernelGetInfo(pi_kernel kernel, pi_kernel_info param_name, | ||
| size_t param_value_size, void *param_value, | ||
| size_t *param_value_size_ret) { | ||
| switch (param_name) { | ||
| case PI_KERNEL_INFO_CONTEXT: | ||
| if (param_value_size_ret) | ||
| *param_value_size_ret = sizeof(GlobalContext.contextHandle); | ||
| if (param_value) | ||
| *static_cast<pi_context *>(param_value) = GlobalContext.contextHandle; | ||
| break; | ||
| case PI_KERNEL_INFO_PROGRAM: | ||
| if (param_value_size_ret) | ||
| *param_value_size_ret = sizeof(GlobalContext.programHandle); | ||
| if (param_value) | ||
| *(pi_program *)param_value = GlobalContext.programHandle; | ||
| break; | ||
| default:; | ||
| } | ||
|
|
||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| void redefineMockForKernelInterop(sycl::unittest::PiMock &Mock) { | ||
| Mock.redefineAfter<sycl::detail::PiApiKind::piContextCreate>( | ||
| after_piContextCreate); | ||
| Mock.redefineAfter<sycl::detail::PiApiKind::piProgramGetInfo>( | ||
| after_piProgramGetInfo); | ||
| Mock.redefineAfter<sycl::detail::PiApiKind::piContextGetInfo>( | ||
| after_piContextGetInfo); | ||
| Mock.redefineAfter<sycl::detail::PiApiKind::piKernelGetInfo>( | ||
| after_piKernelGetInfo); | ||
| Mock.redefine<sycl::detail::PiApiKind::piProgramGetBuildInfo>( | ||
| redefined_piProgramGetBuildInfo); | ||
| } | ||
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,3 +1,4 @@ | ||
| add_sycl_unittest(ThreadSafetyTests OBJECT | ||
| HostAccessorDeadLock.cpp | ||
| InteropKernelEnqueue.cpp | ||
| ) |
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
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,65 @@ | ||
| //==-------- InteropKernelEnqueue.cpp --- Thread safety unit tests ---------==// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include <cstddef> | ||
| #include <gtest/gtest.h> | ||
| #include <helpers/KernelInteropCommon.hpp> | ||
| #include <helpers/PiMock.hpp> | ||
| #include <sycl/sycl.hpp> | ||
|
|
||
| #include "ThreadUtils.h" | ||
|
|
||
| namespace { | ||
| using namespace sycl; | ||
|
|
||
| constexpr std::size_t NArgs = 16; | ||
| constexpr std::size_t ThreadCount = 4; | ||
| constexpr std::size_t LaunchCount = 8; | ||
|
|
||
| pi_uint32 LastArgSet = -1; | ||
| std::size_t LastThread = -1; | ||
| pi_result redefined_piKernelSetArg(pi_kernel kernel, pi_uint32 arg_index, | ||
| size_t arg_size, const void *arg_value) { | ||
| EXPECT_EQ((LastArgSet + 1) % NArgs, arg_index); | ||
| LastArgSet = arg_index; | ||
| std::size_t ArgValue = *static_cast<const std::size_t *>(arg_value); | ||
| if (arg_index == 0) | ||
| LastThread = ArgValue; | ||
| else | ||
| EXPECT_EQ(LastThread, ArgValue); | ||
| return PI_SUCCESS; | ||
| } | ||
|
|
||
| TEST(KernelEnqueue, InteropKernel) { | ||
| unittest::PiMock Mock; | ||
| redefineMockForKernelInterop(Mock); | ||
| Mock.redefine<sycl::detail::PiApiKind::piKernelSetArg>( | ||
| redefined_piKernelSetArg); | ||
|
|
||
| platform Plt = Mock.getPlatform(); | ||
| queue Q; | ||
|
|
||
| DummyHandleT Handle; | ||
| auto KernelCL = reinterpret_cast<typename sycl::backend_traits< | ||
| sycl::backend::opencl>::template input_type<sycl::kernel>>(&Handle); | ||
| auto Kernel = | ||
| sycl::make_kernel<sycl::backend::opencl>(KernelCL, Q.get_context()); | ||
|
|
||
| auto TestLambda = [&](std::size_t ThreadId) { | ||
| Q.submit([&](sycl::handler &CGH) { | ||
| for (std::size_t I = 0; I < NArgs; ++I) | ||
| CGH.set_arg(I, ThreadId); | ||
| CGH.single_task(Kernel); | ||
| }).wait(); | ||
| }; | ||
|
|
||
| for (std::size_t I = 0; I < LaunchCount; ++I) { | ||
| ThreadPool Pool(ThreadCount, TestLambda); | ||
| } | ||
| } | ||
| } // namespace |
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.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
It would be probably better to have an
std::mutexstored inkernel_impl. So, we do not serialize submission of differentsycl::kernels.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
@romanovvlad I tend to agree, although wouldn't this mutex need to be tied to
pi_kernellike described in the comment above? Otherwise we will run into issues if two interopsycl::kernelare created with the same native kernel.There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Sounds reasonable.