-
Notifications
You must be signed in to change notification settings - Fork 795
[SYCL] Fix kernel program cache for multiple devices and refactor some unit tests #5017
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
20 commits
Select commit
Hold shift + click to select a range
7bf586c
Fix kernel program cache for multiple devices
maximdimakov a042881
Clang-format fix
maximdimakov 8458b5e
Merge remote-tracking branch 'origin/sycl' into fix_CP_cache
maximdimakov deaa97c
Add unittest for the fix
maximdimakov ab22f7c
Clang-format fix
maximdimakov c248f7d
Make test only for opencl backend
maximdimakov ca3c7dd
Clang-format fix
maximdimakov 1aa4a39
Fixed problem with kernel release
maximdimakov 5674a12
Use size_t for reinterpret_cast
maximdimakov a718e3f
Fix fail on windows
maximdimakov 30db943
Clang-format
maximdimakov 004c9d5
Addressed review comments
maximdimakov eb5abfb
Clang-format fix
maximdimakov bc4be30
Add more clarifying comment
maximdimakov a530172
Clang-format fix
maximdimakov 034d92c
Merge remote-tracking branch 'origin/sycl' into fix_CP_cache
maximdimakov b3f60a1
[SYCL][NFC] Don't include sycl.hpp from headers
bader cb440dc
[SYCL][NFC] Factor out empty kernel creation boilerplate
bader ed4811e
Refactor some tests
maximdimakov ad007d9
Clang-format fix
maximdimakov 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
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,55 @@ | ||
//===----------------------------------------------------------------------===// | ||
// | ||
// 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 | ||
// | ||
//===----------------------------------------------------------------------===// | ||
|
||
#pragma once | ||
|
||
#include "PiImage.hpp" | ||
|
||
class TestKernel; | ||
|
||
__SYCL_INLINE_NAMESPACE(cl) { | ||
namespace sycl { | ||
namespace detail { | ||
template <> struct KernelInfo<TestKernel> { | ||
static constexpr unsigned getNumParams() { return 0; } | ||
static const kernel_param_desc_t &getParamDesc(int) { | ||
static kernel_param_desc_t Dummy; | ||
return Dummy; | ||
} | ||
static constexpr const char *getName() { return "TestKernel"; } | ||
static constexpr bool isESIMD() { return false; } | ||
static constexpr bool callsThisItem() { return false; } | ||
static constexpr bool callsAnyThisFreeFunction() { return false; } | ||
}; | ||
|
||
} // namespace detail | ||
} // namespace sycl | ||
} // __SYCL_INLINE_NAMESPACE(cl) | ||
|
||
static sycl::unittest::PiImage generateDefaultImage() { | ||
using namespace sycl::unittest; | ||
|
||
PiPropertySet PropSet; | ||
|
||
std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data | ||
|
||
PiArray<PiOffloadEntry> Entries = makeEmptyKernels({"TestKernel"}); | ||
|
||
PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format | ||
__SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec | ||
"", // Compile options | ||
"", // Link options | ||
std::move(Bin), | ||
std::move(Entries), | ||
std::move(PropSet)}; | ||
|
||
return Img; | ||
} | ||
|
||
static sycl::unittest::PiImage Img = generateDefaultImage(); | ||
static sycl::unittest::PiImageArray<1> ImgArray{&Img}; |
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
184 changes: 184 additions & 0 deletions
184
sycl/unittests/kernel-and-program/MultipleDevsCache.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,184 @@ | ||
//==--- KPCache.cpp --- KernelProgramCache for multiple devices unit test --==// | ||
// | ||
// 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 | ||
// | ||
//===----------------------------------------------------------------------===// | ||
|
||
#define SYCL2020_DISABLE_DEPRECATION_WARNINGS | ||
|
||
#include "detail/context_impl.hpp" | ||
#include "detail/kernel_bundle_impl.hpp" | ||
#include "detail/kernel_program_cache.hpp" | ||
#include <helpers/CommonRedefinitions.hpp> | ||
#include <helpers/PiImage.hpp> | ||
#include <helpers/PiMock.hpp> | ||
#include <helpers/TestKernel.hpp> | ||
|
||
#include <gtest/gtest.h> | ||
|
||
#include <iostream> | ||
|
||
using namespace sycl; | ||
|
||
static pi_result redefinedContextCreate( | ||
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) { | ||
*ret_context = reinterpret_cast<pi_context>(123); | ||
return PI_SUCCESS; | ||
} | ||
|
||
static pi_result redefinedContextRelease(pi_context context) { | ||
return PI_SUCCESS; | ||
} | ||
|
||
static pi_result redefinedDevicesGet(pi_platform platform, | ||
pi_device_type device_type, | ||
pi_uint32 num_entries, pi_device *devices, | ||
pi_uint32 *num_devices) { | ||
if (num_devices) { | ||
*num_devices = static_cast<pi_uint32>(2); | ||
return PI_SUCCESS; | ||
} | ||
|
||
if (num_entries == 2 && devices) { | ||
devices[0] = reinterpret_cast<pi_device>(1111); | ||
devices[1] = reinterpret_cast<pi_device>(2222); | ||
} | ||
return PI_SUCCESS; | ||
} | ||
|
||
static pi_result redefinedDeviceGetInfo(pi_device device, | ||
pi_device_info param_name, | ||
size_t param_value_size, | ||
void *param_value, | ||
size_t *param_value_size_ret) { | ||
if (param_name == PI_DEVICE_INFO_TYPE) { | ||
auto *Result = reinterpret_cast<_pi_device_type *>(param_value); | ||
*Result = PI_DEVICE_TYPE_GPU; | ||
} | ||
if (param_name == PI_DEVICE_INFO_COMPILER_AVAILABLE) { | ||
auto *Result = reinterpret_cast<pi_bool *>(param_value); | ||
*Result = true; | ||
} | ||
return PI_SUCCESS; | ||
} | ||
|
||
static pi_result redefinedDeviceRetain(pi_device device) { return PI_SUCCESS; } | ||
|
||
static pi_result redefinedDeviceRelease(pi_device device) { return PI_SUCCESS; } | ||
|
||
static pi_result redefinedQueueCreate(pi_context context, pi_device device, | ||
pi_queue_properties properties, | ||
pi_queue *queue) { | ||
*queue = reinterpret_cast<pi_queue>(1234); | ||
return PI_SUCCESS; | ||
} | ||
|
||
static pi_result redefinedQueueRelease(pi_queue command_queue) { | ||
return PI_SUCCESS; | ||
} | ||
|
||
static size_t ProgramNum = 12345; | ||
static pi_result redefinedProgramCreate(pi_context context, const void *il, | ||
size_t length, | ||
pi_program *res_program) { | ||
size_t CurrentProgram = ProgramNum; | ||
*res_program = reinterpret_cast<pi_program>(CurrentProgram); | ||
++ProgramNum; | ||
return PI_SUCCESS; | ||
} | ||
|
||
static int RetainCounter = 0; | ||
static pi_result redefinedProgramRetain(pi_program program) { | ||
++RetainCounter; | ||
return PI_SUCCESS; | ||
} | ||
|
||
static int KernelReleaseCounter = 0; | ||
static pi_result redefinedKernelRelease(pi_kernel kernel) { | ||
++KernelReleaseCounter; | ||
return PI_SUCCESS; | ||
} | ||
|
||
class MultipleDeviceCacheTest : public ::testing::Test { | ||
public: | ||
MultipleDeviceCacheTest() : Plt{default_selector()} {} | ||
|
||
protected: | ||
void SetUp() override { | ||
if (Plt.is_host() || Plt.get_backend() != backend::opencl) { | ||
return; | ||
} | ||
|
||
Mock = std::make_unique<unittest::PiMock>(Plt); | ||
|
||
setupDefaultMockAPIs(*Mock); | ||
Mock->redefine<detail::PiApiKind::piDevicesGet>(redefinedDevicesGet); | ||
Mock->redefine<detail::PiApiKind::piDeviceGetInfo>(redefinedDeviceGetInfo); | ||
Mock->redefine<detail::PiApiKind::piDeviceRetain>(redefinedDeviceRetain); | ||
Mock->redefine<detail::PiApiKind::piDeviceRelease>(redefinedDeviceRelease); | ||
Mock->redefine<detail::PiApiKind::piContextCreate>(redefinedContextCreate); | ||
Mock->redefine<detail::PiApiKind::piContextRelease>( | ||
redefinedContextRelease); | ||
Mock->redefine<detail::PiApiKind::piQueueCreate>(redefinedQueueCreate); | ||
Mock->redefine<detail::PiApiKind::piQueueRelease>(redefinedQueueRelease); | ||
Mock->redefine<detail::PiApiKind::piProgramRetain>(redefinedProgramRetain); | ||
Mock->redefine<detail::PiApiKind::piProgramCreate>(redefinedProgramCreate); | ||
Mock->redefine<detail::PiApiKind::piKernelRelease>(redefinedKernelRelease); | ||
} | ||
|
||
protected: | ||
std::unique_ptr<unittest::PiMock> Mock; | ||
platform Plt; | ||
}; | ||
|
||
// Test that program is retained for each device and each kernel is released | ||
// once | ||
TEST_F(MultipleDeviceCacheTest, ProgramRetain) { | ||
if (Plt.is_host() || Plt.get_backend() != backend::opencl) { | ||
return; | ||
} | ||
{ | ||
std::vector<sycl::device> Devices = Plt.get_devices(info::device_type::gpu); | ||
sycl::context Context(Devices); | ||
sycl::queue Queue(Context, Devices[0]); | ||
assert(Devices.size() == 2); | ||
|
||
auto Bundle = cl::sycl::get_kernel_bundle<sycl::bundle_state::input>( | ||
Queue.get_context()); | ||
Queue.submit( | ||
[&](cl::sycl::handler &cgh) { cgh.single_task<TestKernel>([]() {}); }); | ||
|
||
auto BundleObject = cl::sycl::build(Bundle, Bundle.get_devices()); | ||
auto KernelID = cl::sycl::get_kernel_id<TestKernel>(); | ||
auto Kernel = BundleObject.get_kernel(KernelID); | ||
|
||
// Because of emulating 2 devices program is retained for each one in | ||
// build(). It is also depends on number of device images. This test has one | ||
// image, but other tests can create other images. Additional variable is | ||
// added to control count of piProgramRetain calls | ||
auto BundleImpl = getSyclObjImpl(Bundle); | ||
int NumRetains = BundleImpl->size() * 2; | ||
|
||
EXPECT_EQ(RetainCounter, NumRetains) | ||
<< "Expect " << NumRetains << " piProgramRetain calls"; | ||
|
||
auto CtxImpl = detail::getSyclObjImpl(Context); | ||
detail::KernelProgramCache::KernelCacheT &KernelCache = | ||
CtxImpl->getKernelProgramCache().acquireKernelsPerProgramCache().get(); | ||
|
||
EXPECT_EQ(KernelCache.size(), (size_t)2) << "Expect 2 kernels in cache"; | ||
} | ||
// First kernel creating is called in handler::single_task(). | ||
// kernel_bundle::get_kernel() creates a kernel and shares it with created | ||
// programs. Also the kernel is retained in kernel_bundle::get_kernel(). A | ||
// kernel is removed from cache if piKernelRelease was called for it, so it | ||
// will not be removed twice for the other programs. As a result we must | ||
// expect 3 piKernelRelease calls. | ||
EXPECT_EQ(KernelReleaseCounter, 3) << "Expect 3 piKernelRelease calls"; | ||
} |
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
Oops, something went wrong.
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.
Please, add comment describing how number
3
was retrieved. Why is it expected and not 2 or four?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.
And why not test to see if KernelReleaseCounter is equal to the RetainCounter?
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.
I added clarifying comments
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.
AFAIR, there's some mem-leak test which checks for retain-release parity.
Although, I'm not sure it checks for kernels.