fix(zebin): Enforce fallback to CTNI on TGL/ICL for nGEN dummy kernel

For TGL and ICL platforms - if on clCreateProgramWithSource()
call we detect a nGen dummy kernel usage - enforce fallback to the
patchtokens format (only for this kernel).
- corrected naming
- minor ULTs refactor (less dependencies).
Signed-off-by: Kacper Nowak <kacper.nowak@intel.com>

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<const char>(options.c_str(), options.length());
             inputArgs.internalOptions = ArrayRef<const char>(internalOptions.c_str(), internalOptions.length());
             inputArgs.GTPinInput = gtpinGetIgcInit();

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;
     }

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<uint32_t>::max();
     std::mutex lockMutex;
     uint32_t exposedKernels = 0;

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;

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<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
     ailConfigurationTable[productFamily] = nullptr;
@@ -644,6 +656,53 @@ TEST_F(ProgramFromSourceTest, givenEmptyAilWhenCreateProgramWithSourcesThenSourc
     pProgram->release();
 }
 
+HWTEST2_F(MinimumProgramFixture, givenEmptyAilWhenCreateProgramWithSourcesAndWithDummyKernelThenDoNotSetFallbackRequired, IsICLLPOrTGLLP) {
+
+    VariableBackup<AILConfiguration *> 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<MockProgram>(
+        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<SucceedingGenBinaryProgram>(
+        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);

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<uint32_t>::max();
+};

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 {
+bool AILConfiguration::sourcesContain(const std::string &sources, std::string_view contentToFind) const {
-    return (kernelsSources.find(kernelName) != std::string::npos);
+    return (sources.find(contentToFind) != std::string::npos);
 }
 
 } // namespace NEO

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 <PRODUCT_FAMILY product>

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 <PRODUCT_FAMILY Product>
 void AILConfigurationHw<Product>::modifyKernelIfRequired(std::string &kernel) {
 }
 
+template <PRODUCT_FAMILY Product>
+inline void AILConfigurationHw<Product>::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) {
+}
+
 template <PRODUCT_FAMILY Product>
 inline void AILConfigurationHw<Product>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
 }

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<IGFX_ICELAKE_LP>::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<IGFX_ICELAKE_LP>::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &setFallback) {
+    std::string_view dummyKernelSource{"kernel void _(){}"};
+    if (sourcesContain(kernelSources, dummyKernelSource)) {
+        setFallback = true;
+    }
+}
+
 template class AILConfigurationHw<IGFX_ICELAKE_LP>;
 } // namespace NEO

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<IGFX_TIGERLAKE_LP> enableAILTGLLP;
 
 std::map<std::string_view, std::vector<AILEnumeration>> 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<IGFX_TIGERLAKE_LP>::forceFallbackToPatchtokensIfRequired(const std::string &kernelSources, bool &requiresFallback) {
+    std::string_view dummyKernelSource{"kernel void _(){}"};
+    if (sourcesContain(kernelSources, dummyKernelSource)) {
+        requiresFallback = true;
+    }
+}
+
 template class AILConfigurationHw<IGFX_TIGERLAKE_LP>;
 
 } // namespace NEO

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 <algorithm>
@@ -62,7 +62,7 @@ void AILConfigurationHw<IGFX_DG2>::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);
         }
     }

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<IGFX_ALDERLAKE_S>;
 using IsADLP = IsProduct<IGFX_ALDERLAKE_P>;
 using IsRKL = IsProduct<IGFX_ROCKETLAKE>;
 
+using IsICLLPOrTGLLP = IsAnyProducts<IGFX_ICELAKE_LP, IGFX_TIGERLAKE_LP>;
+
 using IsXEHP = IsProduct<IGFX_XE_HP_SDV>;
 using IsNotXEHP = IsNotWithinProducts<IGFX_XE_HP_SDV, IGFX_XE_HP_SDV>;
 

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<productFamily> {
   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_TRUE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddle"));
-    EXPECT_FALSE(ail.sourcesContainKernel(kernelSources, "CopyBufferToBufferMiddleStateless"));
+    EXPECT_FALSE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddleStateless"));
 }
 
 HWTEST2_F(AILTests, whenCheckingIsKernelHashCorrectThenCorrectResultIsReturned, IsAtLeastGen12lp) {

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()

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<IGFX_ICELAKE_LP> 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<IGFX_ICELAKE_LP> ail;
+    ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI);
+
+    EXPECT_FALSE(enforceRebuildToCTNI);
+}
+
+} // namespace NEO

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()

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<IGFX_TIGERLAKE_LP> 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<IGFX_TIGERLAKE_LP> ail;
+    ail.forceFallbackToPatchtokensIfRequired(dummyKernelSource, enforceRebuildToCTNI);
+
+    EXPECT_FALSE(enforceRebuildToCTNI);
+}
+
+} // namespace NEO

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<productFamily> {
       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);