Add ail for FAHBench on DG2

Add missing synchronization in the FAHBench kernel
which caused hang on DG2.

Signed-off-by: Kamil Kopryk <kamil.kopryk@intel.com>
Related-To: NEO-6946

opencl/source/program/create.inl

@@ -5,6 +5,7 @@
  *
  */
 
+#include "shared/source/ail/ail_configuration.h"
 #include "shared/source/debug_settings/debug_settings_manager.h"
 #include "shared/source/device/device.h"
 #include "shared/source/helpers/constants.h"
@@ -73,6 +74,13 @@ T *Program::create(
         lengths);
 
     if (CL_SUCCESS == retVal) {
+
+        auto &hwInfo = pContext->getDevice(0)->getHardwareInfo();
+        auto ail = AILConfiguration::get(hwInfo.platform.eProductFamily);
+        if (ail) {
+            ail->modifyKernelIfRequired(combinedString);
+        }
+
         program = new T(pContext, false, pContext->getDevices());
         program->sourceCode.swap(combinedString);
         program->createdFrom = CreatedFrom::SOURCE;

opencl/test/unit_test/program/program_tests.cpp

@@ -7,6 +7,7 @@
 
 #include "opencl/test/unit_test/program/program_tests.h"
 
+#include "shared/source/ail/ail_configuration.h"
 #include "shared/source/command_stream/command_stream_receiver_hw.h"
 #include "shared/source/compiler_interface/compiler_warnings/compiler_warnings.h"
 #include "shared/source/compiler_interface/intermediate_representations.h"
@@ -696,6 +697,27 @@ HWTEST_F(ProgramFromBinaryTest, givenIsaAllocationUsedByMultipleCsrsWhenItIsDele
     EXPECT_TRUE(csr1.requiresInstructionCacheFlush);
 }
 
+TEST_F(ProgramFromSourceTest, givenEmptyAilWhenCreateProgramWithSourcesThenSourcesDoNotChange) {
+
+    VariableBackup<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
+    ailConfigurationTable[productFamily] = nullptr;
+    const char *sources[] = {"kernel() {}"};
+    size_t knownSourceSize = strlen(sources[0]);
+
+    auto pProgram = Program::create<MockProgram>(
+        pContext,
+        1,
+        sources,
+        &knownSourceSize,
+        retVal);
+
+    ASSERT_NE(nullptr, pProgram);
+    ASSERT_EQ(CL_SUCCESS, retVal);
+
+    EXPECT_STREQ(sources[0], pProgram->sourceCode.c_str());
+    pProgram->release();
+}
+
 TEST_F(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
     KernelBinaryHelper kbHelper(binaryFileName, true);
     auto device = pPlatform->getClDevice(0);

shared/source/ail/CMakeLists.txt

@@ -1,5 +1,5 @@
 #
-# Copyright (C) 2021 Intel Corporation
+# Copyright (C) 2021-2022 Intel Corporation
 #
 # SPDX-License-Identifier: MIT
 #
@@ -7,7 +7,9 @@
 set(NEO_CORE_AIL
     ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
     ${CMAKE_CURRENT_SOURCE_DIR}/ail_configuration.h
-    ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}ail_configuration.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/ail_configuration.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}ail_configuration_extra.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/ail_configuration_base.inl
 )
 
 set_property(GLOBAL PROPERTY NEO_CORE_AIL ${NEO_CORE_AIL})

shared/source/ail/ail_configuration.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -7,41 +7,18 @@
 
 #include "shared/source/ail/ail_configuration.h"
 
-#include <map>
+#include "shared/source/helpers/hash.h"
+
+#include <string>
 #include <string_view>
 
 namespace NEO {
-/*
- * fp64 support is unavailable on some Intel GPUs, and the SW emulation in IGC should not be enabled by default.
- * For Blender, fp64 is not performance-critical - SW emulation is good enough for the application to be usable
- * (some versions would not function correctly without it).
- *
- */
-
-std::map<std::string_view, std::vector<AILEnumeration>> applicationMap = {{"blender", {AILEnumeration::ENABLE_FP64}}};
-
-AILConfiguration *ailConfigurationTable[IGFX_MAX_PRODUCT] = {};
-
-AILConfiguration *AILConfiguration::get(PRODUCT_FAMILY productFamily) {
-    return ailConfigurationTable[productFamily];
+bool AILConfiguration::isKernelHashCorrect(const std::string &kernelsSources, uint64_t expectedHash) const {
+    const auto hash = Hash::hash(kernelsSources.c_str(), kernelsSources.length());
+    return hash == expectedHash;
 }
 
-void AILConfiguration::apply(RuntimeCapabilityTable &runtimeCapabilityTable) {
-    auto search = applicationMap.find(processName);
-
-    if (search != applicationMap.end()) {
-        for (size_t i = 0; i < search->second.size(); ++i) {
-            switch (search->second[i]) {
-            case AILEnumeration::ENABLE_FP64:
-                runtimeCapabilityTable.ftrSupportsFP64 = true;
-                break;
-            default:
-                break;
-            }
-        }
-    }
-
-    applyExt(runtimeCapabilityTable);
+bool AILConfiguration::sourcesContainKernel(const std::string &kernelsSources, std::string_view kernelName) const {
+    return (kernelsSources.find(kernelName) != std::string::npos);
 }
-
 } // namespace NEO

shared/source/ail/ail_configuration.h

@@ -44,9 +44,14 @@ class AILConfiguration {
 
     virtual void apply(RuntimeCapabilityTable &runtimeCapabilityTable);
 
+    virtual void modifyKernelIfRequired(std::string &kernel) = 0;
+
   protected:
     virtual void applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) = 0;
     std::string processName;
+
+    bool sourcesContainKernel(const std::string &kernelSources, std::string_view kernelName) const;
+    MOCKABLE_VIRTUAL bool isKernelHashCorrect(const std::string &kernelSources, uint64_t expectedHash) const;
 };
 
 extern AILConfiguration *ailConfigurationTable[IGFX_MAX_PRODUCT];
@@ -60,6 +65,8 @@ class AILConfigurationHw : public AILConfiguration {
     }
 
     void applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) override;
+
+    void modifyKernelIfRequired(std::string &kernel) override;
 };
 
 template <PRODUCT_FAMILY product>
@@ -68,4 +75,5 @@ struct EnableAIL {
         ailConfigurationTable[product] = &AILConfigurationHw<product>::get();
     }
 };
+
 } // namespace NEO

shared/source/ail/ail_configuration_base.inl

@@ -0,0 +1,20 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include <string>
+
+namespace NEO {
+
+template <PRODUCT_FAMILY Product>
+void AILConfigurationHw<Product>::modifyKernelIfRequired(std::string &kernel) {
+}
+
+template <PRODUCT_FAMILY Product>
+inline void AILConfigurationHw<Product>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
+}
+
+} // namespace NEO

shared/source/ail/ail_configuration_extra.cpp

@@ -0,0 +1,48 @@
+/*
+ * Copyright (C) 2021-2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/ail/ail_configuration.h"
+#include "shared/source/helpers/hash.h"
+
+#include <map>
+#include <string_view>
+
+namespace NEO {
+/*
+ * fp64 support is unavailable on some Intel GPUs, and the SW emulation in IGC should not be enabled by default.
+ * For Blender, fp64 is not performance-critical - SW emulation is good enough for the application to be usable
+ * (some versions would not function correctly without it).
+ *
+ */
+
+std::map<std::string_view, std::vector<AILEnumeration>> applicationMap = {{"blender", {AILEnumeration::ENABLE_FP64}}};
+
+AILConfiguration *ailConfigurationTable[IGFX_MAX_PRODUCT] = {};
+
+AILConfiguration *AILConfiguration::get(PRODUCT_FAMILY productFamily) {
+    return ailConfigurationTable[productFamily];
+}
+
+void AILConfiguration::apply(RuntimeCapabilityTable &runtimeCapabilityTable) {
+    auto search = applicationMap.find(processName);
+
+    if (search != applicationMap.end()) {
+        for (size_t i = 0; i < search->second.size(); ++i) {
+            switch (search->second[i]) {
+            case AILEnumeration::ENABLE_FP64:
+                runtimeCapabilityTable.ftrSupportsFP64 = true;
+                break;
+            default:
+                break;
+            }
+        }
+    }
+
+    applyExt(runtimeCapabilityTable);
+}
+
+} // namespace NEO

shared/source/ail/gen12lp/adln/ail_configuration_adln.cpp

@@ -6,14 +6,15 @@
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_ALDERLAKE_N> enableAILADLN;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapADLN = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_ALDERLAKE_N>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_ALDERLAKE_N>;
+
 } // namespace NEO

shared/source/ail/gen12lp/adlp/ail_configuration_adlp.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_ALDERLAKE_P> enableAILADLP;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapADLP = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_ALDERLAKE_P>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_ALDERLAKE_P>;
+
 } // namespace NEO

shared/source/ail/gen12lp/adls/ail_configuration_adls.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_ALDERLAKE_S> enableAILADLS;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapADLS = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_ALDERLAKE_S>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_ALDERLAKE_S>;
+
 } // namespace NEO

shared/source/ail/gen12lp/dg1/ail_configuration_dg1.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_DG1> enableAILDG1;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapDG1 = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_DG1>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_DG1>;
+
 } // namespace NEO

shared/source/ail/gen12lp/rkl/ail_configuration_rkl.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_ROCKETLAKE> enableAILRKL;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapRKL = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_ROCKETLAKE>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_ROCKETLAKE>;
+
 } // namespace NEO

shared/source/ail/gen12lp/tgllp/ail_configuration_tgllp.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_TIGERLAKE_LP> enableAILTGLLP;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapTGLLP = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_TIGERLAKE_LP>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_TIGERLAKE_LP>;
+
 } // namespace NEO

shared/source/ail/xe_hp_core/xe_hp_sdv/ail_configuration_xe_hp_sdv.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_XE_HP_SDV> enableAILXEHPSDV;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapXEHPSDV = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_XE_HP_SDV>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_XE_HP_SDV>;
+
 } // namespace NEO

shared/source/ail/xe_hpc_core/pvc/ail_configuration_pvc.cpp

@@ -1,19 +1,20 @@
 /*
- * Copyright (C) 2021 Intel Corporation
+ * Copyright (C) 2021-2022 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
 #include <map>
+
 namespace NEO {
 static EnableAIL<IGFX_PVC> enableAILPVC;
 
 std::map<std::string_view, std::vector<AILEnumeration>> applicationMapPVC = {};
 
-template <>
-inline void AILConfigurationHw<IGFX_PVC>::applyExt(RuntimeCapabilityTable &runtimeCapabilityTable) {
-}
+template class AILConfigurationHw<IGFX_PVC>;
+
 } // namespace NEO

shared/source/ail/xe_hpg_core/dg2/ail_configuration_dg2.cpp

@@ -6,8 +6,13 @@
  */
 
 #include "shared/source/ail/ail_configuration.h"
+#include "shared/source/ail/ail_configuration_base.inl"
 
+#include <algorithm>
 #include <map>
+#include <string>
+#include <vector>
+
 namespace NEO {
 static EnableAIL<IGFX_DG2> enableAILDG2;
 
@@ -29,4 +34,37 @@ inline void AILConfigurationHw<IGFX_DG2>::applyExt(RuntimeCapabilityTable &runti
         }
     }
 }
+
+struct ApplicationKernelFixDg2 {
+    std::string_view applicationName;
+    std::string_view kernelName;
+    uint64_t kernelHash;
+    size_t fixStartPosition;
+    std::string_view fixCode;
+};
+
+// There is a known functional bug in OpenMM that was recently fixed (https://github.com/openmm/openmm/commit/7af08783e08d3219c1a5f5aa3eff18f8421a9d83)
+// FAHbench is known to use older version of OpenMM (containing a bug) - we patch this kernel by injecting the missing barrier to ensure it's functionally correct on DG2.
+
+const std::vector<ApplicationKernelFixDg2> applicationsKernelFixesDG2 =
+    {{"FAHBench-gui", "findBlocksWithInteractions", 0xa39732fc26656899, 12651u, "else { SYNC_WARPS; }"},
+     {"FAHBench-cmd", "findBlocksWithInteractions", 0xa39732fc26656899, 12651u, "else { SYNC_WARPS; }"}};
+
+template <>
+void AILConfigurationHw<IGFX_DG2>::modifyKernelIfRequired(std::string &kernelsSources) {
+
+    auto it = std::find_if(applicationsKernelFixesDG2.begin(), applicationsKernelFixesDG2.end(), [this](const auto &param) {
+        return this->processName == param.applicationName;
+    });
+
+    if (it != applicationsKernelFixesDG2.end()) {
+
+        if (sourcesContainKernel(kernelsSources, it->kernelName) && isKernelHashCorrect(kernelsSources, it->kernelHash)) {
+            kernelsSources.insert(it->fixStartPosition, it->fixCode);
+        }
+    }
+}
+
+template class AILConfigurationHw<IGFX_DG2>;
+
 } // namespace NEO

shared/test/unit_test/ail/ail_tests.cpp

@@ -18,7 +18,9 @@ using AILTests = ::testing::Test;
 template <PRODUCT_FAMILY productFamily>
 class AILMock : public AILConfigurationHw<productFamily> {
   public:
+    using AILConfiguration::isKernelHashCorrect;
     using AILConfiguration::processName;
+    using AILConfiguration::sourcesContainKernel;
 };
 
 HWTEST2_F(AILTests, givenUninitializedTemplateWhenGetAILConfigurationThenNullptrIsReturned, IsSKL) {
@@ -63,4 +65,82 @@ HWTEST2_F(AILTests, givenInitilizedTemplateWhenApplyWithWondershareFilmora11IsCa
     EXPECT_EQ(rtTable.blitterOperationsSupported, false);
 }
 
+HWTEST2_F(AILTests, whenCheckingIfSourcesContainKernelThenCorrectResultIsReturned, IsAtLeastGen12lp) {
+    VariableBackup<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
+    AILMock<productFamily> ail;
+    ailConfigurationTable[productFamily] = &ail;
+    auto ailConfiguration = AILConfiguration::get(productFamily);
+    ASSERT_NE(nullptr, ailConfiguration);
+
+    std::string kernelSources = R"( 
+__kernel void CopyBufferToBufferLeftLeftover(
+    const __global uchar* pSrc,
+    __global uchar* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
+}
+
+__kernel void CopyBufferToBufferMiddle(
+    const __global uint* pSrc,
+    __global uint* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst += dstOffsetInBytes >> 2;
+    pSrc += srcOffsetInBytes >> 2;
+    uint4 loaded = vload4(gid, pSrc);
+    vstore4(loaded, gid, pDst);)";
+
+    EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "CopyBufferToBufferMiddle"));
+    EXPECT_FALSE(ail.sourcesContainKernel(kernelSources, "CopyBufferToBufferMiddleStateless"));
+}
+
+HWTEST2_F(AILTests, whenCheckingIsKernelHashCorrectThenCorrectResultIsReturned, IsAtLeastGen12lp) {
+    VariableBackup<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
+    AILMock<productFamily> ail;
+    ailConfigurationTable[productFamily] = &ail;
+    auto ailConfiguration = AILConfiguration::get(productFamily);
+    ASSERT_NE(nullptr, ailConfiguration);
+
+    std::string kernelSources = R"( 
+__kernel void CopyBufferToBufferLeftLeftover(
+    const __global uchar* pSrc,
+    __global uchar* pDst,
+    uint srcOffsetInBytes,
+    uint dstOffsetInBytes)
+{
+    unsigned int gid = get_global_id(0);
+    pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
+}
+)";
+
+    auto expectedHash = 0xafeba928e880fd89;
+
+    // If this check fails, probably hash algorithm has been changed.
+    // In this case we must regenerate hashes in AIL applications kernels
+    EXPECT_TRUE(ail.isKernelHashCorrect(kernelSources, expectedHash));
+
+    kernelSources.insert(0, "text");
+    EXPECT_FALSE(ail.isKernelHashCorrect(kernelSources, expectedHash));
+}
+
+HWTEST2_F(AILTests, whenModifyKernelIfRequiredIsCalledThenDontChangeKernelSources, IsAtLeastGen12lp) {
+    VariableBackup<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
+    AILMock<productFamily> ail;
+    ailConfigurationTable[productFamily] = &ail;
+    auto ailConfiguration = AILConfiguration::get(productFamily);
+    ASSERT_NE(nullptr, ailConfiguration);
+
+    std::string kernelSources = "example_kernel(){}";
+    auto copyKernel = kernelSources;
+
+    ail.modifyKernelIfRequired(kernelSources);
+
+    EXPECT_STREQ(copyKernel.c_str(), kernelSources.c_str());
+}
+
 } // namespace NEO

shared/test/unit_test/xe_hpg_core/dg2/CMakeLists.txt

@@ -7,12 +7,13 @@
 if(TESTS_DG2)
   target_sources(neo_shared_tests PRIVATE
                  ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
-                 ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}product_config_tests_dg2.cpp
+                 ${CMAKE_CURRENT_SOURCE_DIR}/ail_tests_dg2.cpp
                  ${CMAKE_CURRENT_SOURCE_DIR}/compute_mode_tests_dg2.cpp
                  ${CMAKE_CURRENT_SOURCE_DIR}/device_binary_format_ar_tests_dg2.cpp
                  ${CMAKE_CURRENT_SOURCE_DIR}/excludes_xe_hpg_core_dg2.cpp
                  ${CMAKE_CURRENT_SOURCE_DIR}/hw_helper_tests_dg2.cpp
                  ${CMAKE_CURRENT_SOURCE_DIR}/hw_info_config_tests_dg2.cpp
+                 ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR_SUFFIX}product_config_tests_dg2.cpp
   )
   add_subdirectories()
 endif()

shared/test/unit_test/xe_hpg_core/dg2/ail_tests_dg2.cpp

@@ -0,0 +1,111 @@
+/*
+ * Copyright (C) 2022 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/ail/ail_configuration.h"
+#include "shared/test/common/helpers/default_hw_info.h"
+#include "shared/test/common/helpers/variable_backup.h"
+#include "shared/test/common/test_macros/test.h"
+
+namespace NEO {
+
+using AILTestsDg2 = ::testing::Test;
+
+HWTEST2_F(AILTestsDg2, givenFixesForApplicationsWhenModifyKernelIfRequiredIsCalledThenReturnCorrectResults, IsDG2) {
+
+    class AILMock : public AILConfigurationHw<productFamily> {
+      public:
+        using AILConfiguration::processName;
+        using AILConfiguration::sourcesContainKernel;
+
+        bool isKernelHashCorrect(const std::string &kernelSources, uint64_t expectedHash) const override {
+            return hashCorrect;
+        }
+
+        bool hashCorrect = {true};
+    };
+
+    VariableBackup<AILConfiguration *> ailConfigurationBackup(&ailConfigurationTable[productFamily]);
+    AILMock ail;
+    ailConfigurationTable[productFamily] = &ail;
+    auto ailConfiguration = AILConfiguration::get(defaultHwInfo->platform.eProductFamily);
+    ASSERT_NE(nullptr, ailConfiguration);
+
+    std::string_view fixCode = "else { SYNC_WARPS; }";
+
+    for (auto name : {"FAHBench-gui", "FAHBench-cmd"}) {
+        {
+            ail.processName = name;
+            ail.hashCorrect = true;
+
+            // sources don't contain kernel name
+            std::string kernelSources;
+            kernelSources.resize(16480u, 'a');
+            auto copyKernelSources = kernelSources;
+
+            EXPECT_FALSE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions"));
+
+            // sources should not change
+            ail.modifyKernelIfRequired(kernelSources);
+            EXPECT_STREQ(kernelSources.c_str(), copyKernelSources.c_str());
+
+            // sources should not contain extra synchronization
+            auto it = kernelSources.find(fixCode);
+            EXPECT_EQ(it, std::string::npos);
+        }
+        {
+            // sources contain kernel name
+            std::string kernelSources;
+            kernelSources.resize(16480u, 'a');
+            kernelSources.insert(1024u, "findBlocksWithInteractions");
+            auto copyKernelSources = kernelSources;
+            EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions"));
+
+            // sources should change
+            ail.modifyKernelIfRequired(kernelSources);
+            EXPECT_STRNE(kernelSources.c_str(), copyKernelSources.c_str());
+
+            // sources should contain extra synchronization
+            auto it = kernelSources.find(fixCode);
+            EXPECT_NE(it, std::string::npos);
+
+            constexpr auto expectedFixStartPosition = 12651u;
+            EXPECT_EQ(expectedFixStartPosition, it);
+        }
+        {
+            // hash doesn't match
+            ail.hashCorrect = false;
+
+            // sources  contain kernel name
+            std::string kernelSources;
+            kernelSources.resize(16480u, 'a');
+            kernelSources.insert(1024u, "findBlocksWithInteractions");
+            auto copyKernelSources = kernelSources;
+
+            EXPECT_TRUE(ail.sourcesContainKernel(kernelSources, "findBlocksWithInteractions"));
+
+            // sources should not change
+            ail.modifyKernelIfRequired(kernelSources);
+            EXPECT_STREQ(kernelSources.c_str(), copyKernelSources.c_str());
+
+            // sources should not contain extra synchronization
+            auto it = kernelSources.find(fixCode);
+            EXPECT_EQ(it, std::string::npos);
+        }
+    }
+
+    {
+        // sources should not change for non-existent application
+        ail.processName = "nonExistentApplication";
+        ail.hashCorrect = true;
+        std::string kernelSources = "example_kernel(){}";
+        auto copyKernelSources = kernelSources;
+        ail.modifyKernelIfRequired(kernelSources);
+
+        EXPECT_STREQ(copyKernelSources.c_str(), kernelSources.c_str());
+    }
+}
+} // namespace NEO

shared/test/unit_test/xe_hpg_core/dg2/excludes_xe_hpg_core_dg2.cpp

@@ -18,3 +18,4 @@ HWTEST_EXCLUDE_PRODUCT(HwInfoConfigTest, givenHwInfoConfigWhenAskedIfTile64With3
 HWTEST_EXCLUDE_PRODUCT(HwInfoConfigTest, WhenAllowRenderCompressionIsCalledThenTrueIsReturned, IGFX_DG2);
 HWTEST_EXCLUDE_PRODUCT(HwInfoConfigTest, whenConvertingTimestampsToCsDomainThenNothingIsChanged, IGFX_DG2);
 HWTEST_EXCLUDE_PRODUCT(HwInfoConfigTest, givenHwInfoConfigWhenAskedIfStorageInfoAdjustmentIsRequiredThenFalseIsReturned, IGFX_DG2);
+HWTEST_EXCLUDE_PRODUCT(AILTests, whenModifyKernelIfRequiredIsCalledThenDontChangeKernelSources, IGFX_DG2);