diff --git a/opencl/source/program/build.cpp b/opencl/source/program/build.cpp index 4f4c11ca1a..cb83832d4d 100644 --- a/opencl/source/program/build.cpp +++ b/opencl/source/program/build.cpp @@ -117,6 +117,10 @@ cl_int Program::build( NEO::CompilerOptions::concatenateAppend(internalOptions, NEO::DebugManager.flags.InjectInternalBuildOptions.get()); } + if (this->enforceFallbackToPatchtokens) { + CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::disableZebin); + } + inputArgs.apiOptions = ArrayRef(options.c_str(), options.length()); inputArgs.internalOptions = ArrayRef(internalOptions.c_str(), internalOptions.length()); inputArgs.GTPinInput = gtpinGetIgcInit(); diff --git a/opencl/source/program/create.inl b/opencl/source/program/create.inl index 8579f05f0a..28bab2219a 100644 --- a/opencl/source/program/create.inl +++ b/opencl/source/program/create.inl @@ -1,5 +1,5 @@ /* - * Copyright (C) 2018-2022 Intel Corporation + * Copyright (C) 2018-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -78,6 +78,9 @@ T *Program::create( } program = new T(pContext, false, pContext->getDevices()); + if (ail) { + ail->forceFallbackToPatchtokensIfRequired(combinedString, program->enforceFallbackToPatchtokens); + } program->sourceCode.swap(combinedString); program->createdFrom = CreatedFrom::SOURCE; } diff --git a/opencl/source/program/program.h b/opencl/source/program/program.h index 9a7a9cd864..3fcb498cf4 100644 --- a/opencl/source/program/program.h +++ b/opencl/source/program/program.h @@ -371,6 +371,7 @@ class Program : public BaseObject<_cl_program> { bool isBuiltIn = false; bool kernelDebugEnabled = false; + bool enforceFallbackToPatchtokens = false; uint32_t maxRootDeviceIndex = std::numeric_limits::max(); std::mutex lockMutex; uint32_t exposedKernels = 0; diff --git a/opencl/test/unit_test/mocks/mock_program.h b/opencl/test/unit_test/mocks/mock_program.h index 30b47f62d1..5ecd1ee45c 100644 --- a/opencl/test/unit_test/mocks/mock_program.h +++ b/opencl/test/unit_test/mocks/mock_program.h @@ -54,6 +54,7 @@ class MockProgram : public Program { using Program::createProgramFromBinary; using Program::deviceBuildInfos; using Program::disableZebinIfVmeEnabled; + using Program::enforceFallbackToPatchtokens; using Program::extractInternalOptions; using Program::getKernelInfo; using Program::internalOptionsToExtract; diff --git a/opencl/test/unit_test/program/program_tests.cpp b/opencl/test/unit_test/program/program_tests.cpp index 023616e514..75ed9b2f1d 100644 --- a/opencl/test/unit_test/program/program_tests.cpp +++ b/opencl/test/unit_test/program/program_tests.cpp @@ -623,7 +623,19 @@ HWTEST_F(ProgramFromBinaryTest, givenIsaAllocationUsedByMultipleCsrsWhenItIsDele EXPECT_TRUE(csr1.requiresInstructionCacheFlush); } -TEST_F(ProgramFromSourceTest, givenEmptyAilWhenCreateProgramWithSourcesThenSourcesDoNotChange) { +void MinimumProgramFixture::SetUp() { + PlatformFixture::setUp(); + cl_device_id device = pPlatform->getClDevice(0); + rootDeviceIndex = pPlatform->getClDevice(0)->getRootDeviceIndex(); + NEO::ContextFixture::setUp(1, &device); +} + +void MinimumProgramFixture::TearDown() { + NEO::ContextFixture::tearDown(); + NEO::PlatformFixture::tearDown(); +} + +TEST_F(MinimumProgramFixture, givenEmptyAilWhenCreateProgramWithSourcesThenSourcesDoNotChange) { VariableBackup ailConfigurationBackup(&ailConfigurationTable[productFamily]); ailConfigurationTable[productFamily] = nullptr; @@ -644,6 +656,53 @@ TEST_F(ProgramFromSourceTest, givenEmptyAilWhenCreateProgramWithSourcesThenSourc pProgram->release(); } +HWTEST2_F(MinimumProgramFixture, givenEmptyAilWhenCreateProgramWithSourcesAndWithDummyKernelThenDoNotSetFallbackRequired, IsICLLPOrTGLLP) { + + VariableBackup ailConfigurationBackup(&ailConfigurationTable[productFamily]); + ailConfigurationTable[productFamily] = nullptr; + const char *dummyKernelSources[] = {"kernel void _(){}"}; // if detected - should trigger fallback to CTNI + size_t knownSourceSize = strlen(dummyKernelSources[0]); + + auto pProgram = Program::create( + pContext, + 1, + dummyKernelSources, + &knownSourceSize, + retVal); + + ASSERT_NE(nullptr, pProgram); + ASSERT_EQ(CL_SUCCESS, retVal); + + EXPECT_FALSE(pProgram->enforceFallbackToPatchtokens); + pProgram->release(); +} + +TEST_F(MinimumProgramFixture, givenEnforceLegacyBinaryFormatFlagSetWhenBuildingProgramThenInternalOptionsShouldContainDisableZebinOption) { + const char *kernelSources[] = {"some source code"}; + size_t knownSourceSize = strlen(kernelSources[0]); + + auto cip = new MockCompilerInterfaceCaptureBuildOptions(); + auto pDevice = pContext->getDevice(0); + pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip); + + auto pProgram = Program::create( + pContext, + 1, + kernelSources, + &knownSourceSize, + retVal); + + ASSERT_NE(nullptr, pProgram); + ASSERT_EQ(CL_SUCCESS, retVal); + + pProgram->enforceFallbackToPatchtokens = true; + retVal = pProgram->build(pProgram->getDevices(), "", false); + EXPECT_EQ(CL_SUCCESS, retVal); + + EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::disableZebin)); + pProgram->release(); +} + TEST_F(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSuccessOrCorrectErrorCodeIsReturned) { KernelBinaryHelper kbHelper(binaryFileName, true); auto device = pPlatform->getClDevice(0); diff --git a/opencl/test/unit_test/program/program_tests.h b/opencl/test/unit_test/program/program_tests.h index 34418c7329..2230c8b925 100644 --- a/opencl/test/unit_test/program/program_tests.h +++ b/opencl/test/unit_test/program/program_tests.h @@ -1,5 +1,5 @@ /* - * Copyright (C) 2018-2022 Intel Corporation + * Copyright (C) 2018-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -8,6 +8,8 @@ #pragma once #include "opencl/test/unit_test/fixtures/cl_device_fixture.h" #include "opencl/test/unit_test/fixtures/context_fixture.h" +#include "opencl/test/unit_test/fixtures/platform_fixture.h" +#include "opencl/test/unit_test/mocks/mock_program.h" #include "gtest/gtest.h" @@ -23,3 +25,18 @@ class ProgramTests : public NEO::ClDeviceFixture, void SetUp() override; void TearDown() override; }; + +class MinimumProgramFixture : public NEO::ContextFixture, + public NEO::PlatformFixture, + public ::testing::Test { + + using NEO::ContextFixture::setUp; + using NEO::PlatformFixture::setUp; + + protected: + void SetUp() override; + void TearDown() override; + + cl_int retVal = CL_SUCCESS; + uint32_t rootDeviceIndex = std::numeric_limits::max(); +}; \ No newline at end of file diff --git a/shared/source/ail/ail_configuration.cpp b/shared/source/ail/ail_configuration.cpp index 5899a755c4..02311327ae 100644 --- a/shared/source/ail/ail_configuration.cpp +++ b/shared/source/ail/ail_configuration.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -20,8 +20,8 @@ bool AILConfiguration::isKernelHashCorrect(const std::string &kernelsSources, ui return hash == expectedHash; } -bool AILConfiguration::sourcesContainKernel(const std::string &kernelsSources, std::string_view kernelName) const { - return (kernelsSources.find(kernelName) != std::string::npos); +bool AILConfiguration::sourcesContain(const std::string &sources, std::string_view contentToFind) const { + return (sources.find(contentToFind) != std::string::npos); } } // namespace NEO diff --git a/shared/source/ail/ail_configuration.h b/shared/source/ail/ail_configuration.h index 287ab7b2b9..95381289d2 100644 --- a/shared/source/ail/ail_configuration.h +++ b/shared/source/ail/ail_configuration.h @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -49,11 +49,13 @@ class AILConfiguration { virtual void modifyKernelIfRequired(std::string &kernel) = 0; + virtual void forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) = 0; + protected: virtual void applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) = 0; std::string processName; - bool sourcesContainKernel(const std::string &kernelSources, std::string_view kernelName) const; + bool sourcesContain(const std::string &sources, std::string_view contentToFind) const; MOCKABLE_VIRTUAL bool isKernelHashCorrect(const std::string &kernelSources, uint64_t expectedHash) const; }; @@ -70,6 +72,7 @@ class AILConfigurationHw : public AILConfiguration { void applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) override; void modifyKernelIfRequired(std::string &kernel) override; + void forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) override; }; template diff --git a/shared/source/ail/ail_configuration_base.inl b/shared/source/ail/ail_configuration_base.inl index ba8ebd592b..7fb4195505 100644 --- a/shared/source/ail/ail_configuration_base.inl +++ b/shared/source/ail/ail_configuration_base.inl @@ -1,5 +1,5 @@ /* - * Copyright (C) 2022 Intel Corporation + * Copyright (C) 2022-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -15,6 +15,10 @@ template void AILConfigurationHw::modifyKernelIfRequired(std::string &kernel) { } +template +inline void AILConfigurationHw::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) { +} + template inline void AILConfigurationHw::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) { } diff --git a/shared/source/ail/gen11/icllp/ail_configuration_icllp.cpp b/shared/source/ail/gen11/icllp/ail_configuration_icllp.cpp index c49538f2c1..26f69f6550 100644 --- a/shared/source/ail/gen11/icllp/ail_configuration_icllp.cpp +++ b/shared/source/ail/gen11/icllp/ail_configuration_icllp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2022 Intel Corporation + * Copyright (C) 2022-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -34,5 +34,17 @@ inline void AILConfigurationHw::applyExt(RuntimeCapabilityTable } } +// To avoid a known oneDNN issue in ZEBin handling, affecting ICL and TGL platforms, +// fall back to legacy (patchtoken) format when dummy kernel used by nGen is detected. +// Only this specific kernel with that exact source code will be affected. + +template <> +inline void AILConfigurationHw::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &setFallback) { + std::string_view dummyKernelSource{"kernel void _(){}"}; + if (sourcesContain(kernelSources, dummyKernelSource)) { + setFallback = true; + } +} + template class AILConfigurationHw; } // namespace NEO diff --git a/shared/source/ail/gen12lp/tgllp/ail_configuration_tgllp.cpp b/shared/source/ail/gen12lp/tgllp/ail_configuration_tgllp.cpp index e2f3833ff5..0bc96a0095 100644 --- a/shared/source/ail/gen12lp/tgllp/ail_configuration_tgllp.cpp +++ b/shared/source/ail/gen12lp/tgllp/ail_configuration_tgllp.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -15,6 +15,18 @@ static EnableAIL enableAILTGLLP; std::map> applicationMapTGLLP = {}; +// To avoid a known oneDNN issue in ZEBin handling, affecting ICL and TGL platforms, +// fall back to legacy (patchtoken) format when dummy kernel used by nGen is detected. +// Only this specific kernel with that exact source code will be affected. + +template <> +inline void AILConfigurationHw::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) { + std::string_view dummyKernelSource{"kernel void _(){}"}; + if (sourcesContain(kernelSources, dummyKernelSource)) { + requiresFallback = true; + } +} + template class AILConfigurationHw; } // namespace NEO diff --git a/shared/source/ail/xe_hpg_core/dg2/ail_configuration_dg2.cpp b/shared/source/ail/xe_hpg_core/dg2/ail_configuration_dg2.cpp index a69fecdd80..7ca7b2378a 100644 --- a/shared/source/ail/xe_hpg_core/dg2/ail_configuration_dg2.cpp +++ b/shared/source/ail/xe_hpg_core/dg2/ail_configuration_dg2.cpp @@ -1,11 +1,11 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * */ -#include "shared/source/ail/ail_configuration.h" +#include "shared/source/ail/ail_configuration_base.inl" #include "shared/source/helpers/hw_info.h" #include @@ -62,7 +62,7 @@ void AILConfigurationHw::modifyKernelIfRequired(std::string &kernelsSo if (it != applicationsKernelFixesDG2.end()) { - if (sourcesContainKernel(kernelsSources, it->kernelName) && isKernelHashCorrect(kernelsSources, it->kernelHash)) { + if (sourcesContain(kernelsSources, it->kernelName) && isKernelHashCorrect(kernelsSources, it->kernelHash)) { kernelsSources.insert(it->fixStartPosition, it->fixCode); } } diff --git a/shared/test/common/test_macros/header/common_matchers.h b/shared/test/common/test_macros/header/common_matchers.h index 7a2ba8cd96..0432ad9229 100644 --- a/shared/test/common/test_macros/header/common_matchers.h +++ b/shared/test/common/test_macros/header/common_matchers.h @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -60,6 +60,8 @@ using IsADLS = IsProduct; using IsADLP = IsProduct; using IsRKL = IsProduct; +using IsICLLPOrTGLLP = IsAnyProducts; + using IsXEHP = IsProduct; using IsNotXEHP = IsNotWithinProducts; diff --git a/shared/test/unit_test/ail/ail_tests.cpp b/shared/test/unit_test/ail/ail_tests.cpp index c7f5f50811..17ce8749db 100644 --- a/shared/test/unit_test/ail/ail_tests.cpp +++ b/shared/test/unit_test/ail/ail_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2022 Intel Corporation + * Copyright (C) 2021-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -21,7 +21,7 @@ class AILMock : public AILConfigurationHw { public: using AILConfiguration::isKernelHashCorrect; using AILConfiguration::processName; - using AILConfiguration::sourcesContainKernel; + using AILConfiguration::sourcesContain; }; HWTEST2_F(AILTests, givenInitializedTemplateWhenGetAILConfigurationThenNullptrIsNotReturned, IsSKL) { @@ -131,8 +131,8 @@ __kernel void CopyBufferToBufferMiddle( uint4 loaded = vload4(gid, pSrc); vstore4(loaded, gid, pDst);)"; - EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "CopyBufferToBufferMiddle")); - EXPECT_FALSE(ail.sourcesContainKernel(kernelSources, "CopyBufferToBufferMiddleStateless")); + EXPECT_TRUE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddle")); + EXPECT_FALSE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddleStateless")); } HWTEST2_F(AILTests, whenCheckingIsKernelHashCorrectThenCorrectResultIsReturned, IsAtLeastGen12lp) { diff --git a/shared/test/unit_test/gen11/icllp/CMakeLists.txt b/shared/test/unit_test/gen11/icllp/CMakeLists.txt index ed3d5c6fb0..4569ed09dd 100644 --- a/shared/test/unit_test/gen11/icllp/CMakeLists.txt +++ b/shared/test/unit_test/gen11/icllp/CMakeLists.txt @@ -1,5 +1,5 @@ # -# Copyright (C) 2022 Intel Corporation +# Copyright (C) 2022-2023 Intel Corporation # # SPDX-License-Identifier: MIT # @@ -14,6 +14,7 @@ if(TESTS_ICLLP) ${NEO_SHARED_tests_gen11_icllp} ${CMAKE_CURRENT_SOURCE_DIR}/test_hw_info_config_icllp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/test_program_media_sampler_icllp.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ail_tests_icllp.cpp ) add_subdirectories() diff --git a/shared/test/unit_test/gen11/icllp/ail_tests_icllp.cpp b/shared/test/unit_test/gen11/icllp/ail_tests_icllp.cpp new file mode 100644 index 0000000000..3ba9ac94bb --- /dev/null +++ b/shared/test/unit_test/gen11/icllp/ail_tests_icllp.cpp @@ -0,0 +1,35 @@ +/* + * Copyright (C) 2023 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include "shared/source/ail/ail_configuration.h" +#include "shared/test/common/test_macros/hw_test.h" + +namespace NEO { + +using AILTestsIcllp = ::testing::Test; + +HWTEST2_F(AILTestsIcllp, whenKernelSourceIsANGenDummyKernelThenDoEnforcePatchtokensFormat, IsICLLP) { + std::string dummyKernelSource{"kernel void _(){}"}; + bool enforceRebuildToCTNI = false; + + AILConfigurationHw ail; + ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI); + + EXPECT_TRUE(enforceRebuildToCTNI); +} + +HWTEST2_F(AILTestsIcllp, whenKernelSourceIsNotANGenDummyKernelThenDoNotEnforcePatchtokensFormat, IsICLLP) { + std::string dummyKernelSource{"kernel void copybuffer(__global int* a, __global int* b){ //some code }"}; + bool enforceRebuildToCTNI = false; + + AILConfigurationHw ail; + ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI); + + EXPECT_FALSE(enforceRebuildToCTNI); +} + +} // namespace NEO \ No newline at end of file diff --git a/shared/test/unit_test/gen12lp/tgllp/CMakeLists.txt b/shared/test/unit_test/gen12lp/tgllp/CMakeLists.txt index 54757aeafa..405a4a9a69 100644 --- a/shared/test/unit_test/gen12lp/tgllp/CMakeLists.txt +++ b/shared/test/unit_test/gen12lp/tgllp/CMakeLists.txt @@ -1,5 +1,5 @@ # -# Copyright (C) 2022 Intel Corporation +# Copyright (C) 2022-2023 Intel Corporation # # SPDX-License-Identifier: MIT # @@ -11,6 +11,7 @@ if(TESTS_TGLLP) ${NEO_SHARED_tests_genlp12_tgllp} ${CMAKE_CURRENT_SOURCE_DIR}/test_hw_helper_tgllp.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hw_info_config_tests_tgllp.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/ail_tests_tgllp.cpp ) add_subdirectories() diff --git a/shared/test/unit_test/gen12lp/tgllp/ail_tests_tgllp.cpp b/shared/test/unit_test/gen12lp/tgllp/ail_tests_tgllp.cpp new file mode 100644 index 0000000000..81cb7c6d41 --- /dev/null +++ b/shared/test/unit_test/gen12lp/tgllp/ail_tests_tgllp.cpp @@ -0,0 +1,35 @@ +/* + * Copyright (C) 2023 Intel Corporation + * + * SPDX-License-Identifier: MIT + * + */ + +#include "shared/source/ail/ail_configuration.h" +#include "shared/test/common/test_macros/hw_test.h" + +namespace NEO { + +using AILTestsTgllp = ::testing::Test; + +HWTEST2_F(AILTestsTgllp, whenKernelSourceIsANGenDummyKernelThenDoEnforcePatchtokensFormat, IsTGLLP) { + std::string dummyKernelSource{"kernel void _(){}"}; + bool enforceRebuildToCTNI = false; + + AILConfigurationHw ail; + ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI); + + EXPECT_TRUE(enforceRebuildToCTNI); +} + +HWTEST2_F(AILTestsTgllp, whenKernelSourceIsNotANGenDummyKernelThenDoNotEnforcePatchtokensFormat, IsTGLLP) { + std::string dummyKernelSource{"kernel void copybuffer(__global int* a, __global int* b){ //some code }"}; + bool enforceRebuildToCTNI = false; + + AILConfigurationHw ail; + ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI); + + EXPECT_FALSE(enforceRebuildToCTNI); +} + +} // namespace NEO \ No newline at end of file diff --git a/shared/test/unit_test/xe_hpg_core/dg2/ail_tests_dg2.cpp b/shared/test/unit_test/xe_hpg_core/dg2/ail_tests_dg2.cpp index 91c54dc3a3..1fd7c5f66c 100644 --- a/shared/test/unit_test/xe_hpg_core/dg2/ail_tests_dg2.cpp +++ b/shared/test/unit_test/xe_hpg_core/dg2/ail_tests_dg2.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2022 Intel Corporation + * Copyright (C) 2022-2023 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -19,7 +19,7 @@ HWTEST2_F(AILTestsDg2, givenFixesForApplicationsWhenModifyKernelIfRequiredIsCall class AILMock : public AILConfigurationHw { public: using AILConfiguration::processName; - using AILConfiguration::sourcesContainKernel; + using AILConfiguration::sourcesContain; bool isKernelHashCorrect(const std::string &kernelSources, uint64_t expectedHash) const override { return hashCorrect; @@ -46,7 +46,7 @@ HWTEST2_F(AILTestsDg2, givenFixesForApplicationsWhenModifyKernelIfRequiredIsCall kernelSources.resize(16480u, 'a'); auto copyKernelSources = kernelSources; - EXPECT_FALSE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions")); + EXPECT_FALSE(ail.sourcesContain(kernelSources, "findBlocksWithInteractions")); // sources should not change ail.modifyKernelIfRequired(kernelSources); @@ -62,7 +62,7 @@ HWTEST2_F(AILTestsDg2, givenFixesForApplicationsWhenModifyKernelIfRequiredIsCall kernelSources.resize(16480u, 'a'); kernelSources.insert(1024u, "findBlocksWithInteractions"); auto copyKernelSources = kernelSources; - EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions")); + EXPECT_TRUE(ail.sourcesContain(kernelSources, "findBlocksWithInteractions")); // sources should change ail.modifyKernelIfRequired(kernelSources); @@ -85,7 +85,7 @@ HWTEST2_F(AILTestsDg2, givenFixesForApplicationsWhenModifyKernelIfRequiredIsCall kernelSources.insert(1024u, "findBlocksWithInteractions"); auto copyKernelSources = kernelSources; - EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions")); + EXPECT_TRUE(ail.sourcesContain(kernelSources, "findBlocksWithInteractions")); // sources should not change ail.modifyKernelIfRequired(kernelSources);