diff --git a/level_zero/core/test/unit_tests/main.cpp b/level_zero/core/test/unit_tests/main.cpp
index df73ee1b79..ea3607cbc7 100644
--- a/level_zero/core/test/unit_tests/main.cpp
+++ b/level_zero/core/test/unit_tests/main.cpp
@@ -56,7 +56,7 @@ TestEnvironment *environment = nullptr;
 using namespace L0::ult;
 
 PRODUCT_FAMILY productFamily = NEO::DEFAULT_TEST_PLATFORM::hwInfo.platform.eProductFamily;
-GFXCORE_FAMILY renderCoreFamily = NEO::DEFAULT_TEST_PLATFORM::hwInfo.platform.eRenderCoreFamily;
+extern GFXCORE_FAMILY renderCoreFamily;
 int32_t revId = -1;
 uint32_t euPerSubSlice = 0;
 uint32_t sliceCount = 0;
diff --git a/opencl/test/unit_test/gen12lp/rkl/test_hw_info_config_rkl.cpp b/opencl/test/unit_test/gen12lp/rkl/test_hw_info_config_rkl.cpp
index e754f84e81..6008daff39 100644
--- a/opencl/test/unit_test/gen12lp/rkl/test_hw_info_config_rkl.cpp
+++ b/opencl/test/unit_test/gen12lp/rkl/test_hw_info_config_rkl.cpp
@@ -5,6 +5,7 @@
  *
  */
 
+#include "shared/source/helpers/compiler_hw_info_config.h"
 #include "shared/source/os_interface/hw_info_config.h"
 #include "shared/test/common/helpers/default_hw_info.h"
 
@@ -37,3 +38,8 @@ RKLTEST_F(RklHwInfoConfig, givenHwInfoConfigWhenAskedIf3DPipelineSelectWAIsRequi
     const auto &hwInfoConfig = *HwInfoConfig::get(defaultHwInfo->platform.eProductFamily);
     EXPECT_TRUE(hwInfoConfig.is3DPipelineSelectWARequired());
 }
+
+using CompilerHwInfoConfigHelperTestsRkl = ::testing::Test;
+RKLTEST_F(CompilerHwInfoConfigHelperTestsRkl, givenRklWhenIsForceEmuInt32DivRemSPRequiredIsCalledThenReturnsTrue) {
+    EXPECT_TRUE(CompilerHwInfoConfig::get(productFamily)->isForceEmuInt32DivRemSPRequired());
+}
diff --git a/opencl/test/unit_test/gen8/bdw/test_hw_info_config_bdw.cpp b/opencl/test/unit_test/gen8/bdw/test_hw_info_config_bdw.cpp
index 04fc9b746e..ee5961dc11 100644
--- a/opencl/test/unit_test/gen8/bdw/test_hw_info_config_bdw.cpp
+++ b/opencl/test/unit_test/gen8/bdw/test_hw_info_config_bdw.cpp
@@ -5,6 +5,7 @@
  *
  */
 
+#include "shared/source/helpers/compiler_hw_info_config.h"
 #include "shared/test/common/helpers/default_hw_info.h"
 
 #include "test.h"
@@ -77,3 +78,8 @@ BDWTEST_F(BdwHwInfo, givenHwInfoConfigStringThenAfterSetupResultingVmeIsDisabled
     EXPECT_FALSE(hwInfo.capabilityTable.ftrSupportsVmeAvcPreemption);
     EXPECT_FALSE(hwInfo.capabilityTable.supportsVme);
 }
+
+using CompilerHwInfoConfigHelperTestsBdw = ::testing::Test;
+BDWTEST_F(CompilerHwInfoConfigHelperTestsBdw, givenBdwWhenIsStatelessToStatefulBufferOffsetSupportedIsCalledThenReturnsTrue) {
+    EXPECT_FALSE(CompilerHwInfoConfig::get(productFamily)->isStatelessToStatefulBufferOffsetSupported());
+}
diff --git a/opencl/test/unit_test/main.cpp b/opencl/test/unit_test/main.cpp
index adb9c9f0d1..2985dd5015 100644
--- a/opencl/test/unit_test/main.cpp
+++ b/opencl/test/unit_test/main.cpp
@@ -64,7 +64,7 @@ bool disabled = false;
 using namespace NEO;
 
 PRODUCT_FAMILY productFamily = DEFAULT_TEST_PLATFORM::hwInfo.platform.eProductFamily;
-GFXCORE_FAMILY renderCoreFamily = DEFAULT_TEST_PLATFORM::hwInfo.platform.eRenderCoreFamily;
+extern GFXCORE_FAMILY renderCoreFamily;
 
 extern std::string lastTest;
 bool generateRandomInput = false;
diff --git a/opencl/test/unit_test/offline_compiler/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/CMakeLists.txt
index 79ca841e61..0efb40bffc 100644
--- a/opencl/test/unit_test/offline_compiler/CMakeLists.txt
+++ b/opencl/test/unit_test/offline_compiler/CMakeLists.txt
@@ -66,6 +66,7 @@ set(IGDRCL_SRCS_offline_compiler_tests
     ${NEO_SHARED_TEST_DIRECTORY}/common/mocks/mock_compilers.h
     ${NEO_SHARED_TEST_DIRECTORY}/unit_test/device_binary_format/zebin_tests.h
     ${NEO_SHARED_TEST_DIRECTORY}/common/helpers/test_files.cpp
+    ${NEO_SHARED_TEST_DIRECTORY}/common/test_macros/test_excludes.cpp
     ${IGDRCL_SRCS_cloc}
     ${IGDRCL_SRCS_offline_compiler_mock}
     ${NEO_CORE_tests_compiler_mocks}
@@ -90,9 +91,11 @@ link_directories(${CMAKE_RUNTIME_OUTPUT_DIRECTORY})
 add_executable(ocloc_tests ${IGDRCL_SRCS_offline_compiler_tests})
 
 target_include_directories(ocloc_tests PRIVATE
+                           ${NEO_SHARED_TEST_DIRECTORY}/common/test_macros/header${BRANCH_DIR_SUFFIX}
+                           ${NEO_SHARED_TEST_DIRECTORY}/common/test_configuration/unit_tests
                            $<TARGET_PROPERTY:ocloc_lib,INCLUDE_DIRECTORIES>
 )
-target_compile_definitions(ocloc_tests PUBLIC MOCKABLE_VIRTUAL=virtual $<TARGET_PROPERTY:ocloc_lib,INTERFACE_COMPILE_DEFINITIONS>)
+target_compile_definitions(ocloc_tests PUBLIC MOCKABLE_VIRTUAL=virtual $<TARGET_PROPERTY:ocloc_lib,INTERFACE_COMPILE_DEFINITIONS> ${TESTED_GEN_FLAGS_DEFINITONS})
 
 target_link_libraries(ocloc_tests gmock-gtest)
 
diff --git a/opencl/test/unit_test/offline_compiler/gen12lp/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen12lp/CMakeLists.txt
new file mode 100644
index 0000000000..08e63889d3
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen12lp/CMakeLists.txt
@@ -0,0 +1,13 @@
+#
+# Copyright (C) 2021 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(TESTS_GEN12LP)
+  set(IGDRCL_SRCS_offline_compiler_tests_gen12lp
+      ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+  )
+  target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_gen12lp})
+  add_subdirectories()
+endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen12lp/rkl/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen12lp/rkl/CMakeLists.txt
new file mode 100644
index 0000000000..91b25b3ea2
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen12lp/rkl/CMakeLists.txt
@@ -0,0 +1,13 @@
+#
+# Copyright (C) 2021 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(TESTS_RKL)
+  set(IGDRCL_SRCS_offline_compiler_tests_rkl
+      ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+      ${CMAKE_CURRENT_SOURCE_DIR}/offline_compiler_tests_rkl.cpp
+  )
+  target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_rkl})
+endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen12lp/rkl/offline_compiler_tests_rkl.cpp b/opencl/test/unit_test/offline_compiler/gen12lp/rkl/offline_compiler_tests_rkl.cpp
new file mode 100644
index 0000000000..082667efe3
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen12lp/rkl/offline_compiler_tests_rkl.cpp
@@ -0,0 +1,26 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/compiler_interface/compiler_options/compiler_options.h"
+
+#include "opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h"
+#include "opencl/test/unit_test/offline_compiler/offline_compiler_tests.h"
+#include "test.h"
+
+using namespace NEO;
+
+using MockOfflineCompilerRklTests = ::testing::Test;
+RKLTEST_F(MockOfflineCompilerRklTests, givenRklWhenAppendExtraInternalOptionsThenForceEmuInt32DivRemSPIsApplied) {
+
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = " rkl";
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    mockOfflineCompiler.appendExtraInternalOptions(mockOfflineCompiler.hwInfo, internalOptions);
+    size_t found = internalOptions.find(NEO::CompilerOptions::forceEmuInt32DivRemSP.data());
+    EXPECT_NE(std::string::npos, found);
+}
diff --git a/opencl/test/unit_test/offline_compiler/gen8/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen8/CMakeLists.txt
index c6b637f3ec..5357cc856f 100644
--- a/opencl/test/unit_test/offline_compiler/gen8/CMakeLists.txt
+++ b/opencl/test/unit_test/offline_compiler/gen8/CMakeLists.txt
@@ -7,8 +7,9 @@
 if(TESTS_GEN8)
   set(IGDRCL_SRCS_offline_compiler_tests_gen8
       ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
-      ${CMAKE_CURRENT_SOURCE_DIR}/offline_compiler_tests_gen8.cpp
   )
   target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_gen8})
+
   add_subdirectories()
+
 endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen8/bdw/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen8/bdw/CMakeLists.txt
new file mode 100644
index 0000000000..0f94c9e7f5
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen8/bdw/CMakeLists.txt
@@ -0,0 +1,14 @@
+#
+# Copyright (C) 2021 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(TESTS_BDW)
+  set(IGDRCL_SRCS_offline_compiler_tests_bdw
+      ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+      ${CMAKE_CURRENT_SOURCE_DIR}/offline_compiler_tests_bdw.cpp
+  )
+  target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_bdw})
+
+endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen8/bdw/offline_compiler_tests_bdw.cpp b/opencl/test/unit_test/offline_compiler/gen8/bdw/offline_compiler_tests_bdw.cpp
new file mode 100644
index 0000000000..800f8b6d65
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen8/bdw/offline_compiler_tests_bdw.cpp
@@ -0,0 +1,50 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/compiler_interface/compiler_options/compiler_options.h"
+
+#include "opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h"
+#include "opencl/test/unit_test/offline_compiler/offline_compiler_tests.h"
+#include "test.h"
+
+#include "gmock/gmock.h"
+
+using namespace NEO;
+
+using MockOfflineCompilerBdwTests = ::testing::Test;
+
+BDWTEST_F(MockOfflineCompilerBdwTests, givenDebugOptionAndBdwThenInternalOptionShouldNotContainKernelDebugEnable) {
+    std::vector<std::string> argv = {
+        "ocloc",
+        "-q",
+        "-options",
+        "-g",
+        "-file",
+        "test_files/copybuffer.cl",
+        "-device",
+        "bdw"};
+
+    auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
+    ASSERT_NE(nullptr, mockOfflineCompiler);
+    mockOfflineCompiler->initialize(argv.size(), argv);
+
+    std::string internalOptions = mockOfflineCompiler->internalOptions;
+
+    EXPECT_THAT(internalOptions, Not(::testing::HasSubstr("-cl-kernel-debug-enable")));
+}
+
+BDWTEST_F(MockOfflineCompilerBdwTests, GivenBdwWhenParseDebugSettingsThenContainsHasBufferOffsetArg) {
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = "bdw";
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+
+    mockOfflineCompiler.parseDebugSettings();
+
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
+    EXPECT_EQ(std::string::npos, found);
+}
diff --git a/opencl/test/unit_test/offline_compiler/gen8/offline_compiler_tests_gen8.cpp b/opencl/test/unit_test/offline_compiler/gen8/offline_compiler_tests_gen8.cpp
deleted file mode 100644
index 6ae9f6d02e..0000000000
--- a/opencl/test/unit_test/offline_compiler/gen8/offline_compiler_tests_gen8.cpp
+++ /dev/null
@@ -1,33 +0,0 @@
-/*
- * Copyright (C) 2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-#include "opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h"
-#include "opencl/test/unit_test/offline_compiler/offline_compiler_tests.h"
-
-#include "gmock/gmock.h"
-
-namespace NEO {
-
-TEST_F(OfflineCompilerTests, givenDebugOptionWhenPlatformIsBelowGen9ThenInternalOptionShouldNotContainKernelDebugEnable) {
-    std::vector<std::string> argv = {
-        "ocloc",
-        "-options",
-        "-g",
-        "-file",
-        "test_files/copybuffer.cl",
-        "-device",
-        "bdw"};
-
-    auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
-    ASSERT_NE(nullptr, mockOfflineCompiler);
-    mockOfflineCompiler->initialize(argv.size(), argv);
-
-    std::string internalOptions = mockOfflineCompiler->internalOptions;
-
-    EXPECT_THAT(internalOptions, Not(::testing::HasSubstr("-cl-kernel-debug-enable")));
-}
-} // namespace NEO
\ No newline at end of file
diff --git a/opencl/test/unit_test/offline_compiler/gen9/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen9/CMakeLists.txt
new file mode 100644
index 0000000000..c4e2248f26
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen9/CMakeLists.txt
@@ -0,0 +1,14 @@
+#
+# Copyright (C) 2021 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(TESTS_GEN9)
+  set(IGDRCL_SRCS_offline_compiler_tests_gen9
+      ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+  )
+  target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_gen9})
+
+  add_subdirectories()
+endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen9/skl/CMakeLists.txt b/opencl/test/unit_test/offline_compiler/gen9/skl/CMakeLists.txt
new file mode 100644
index 0000000000..e90f4565cc
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen9/skl/CMakeLists.txt
@@ -0,0 +1,14 @@
+#
+# Copyright (C) 2021 Intel Corporation
+#
+# SPDX-License-Identifier: MIT
+#
+
+if(TESTS_SKL)
+  set(IGDRCL_SRCS_offline_compiler_tests_skl
+      ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
+      ${CMAKE_CURRENT_SOURCE_DIR}/offline_compiler_tests_skl.cpp
+  )
+  target_sources(ocloc_tests PRIVATE ${IGDRCL_SRCS_offline_compiler_tests_skl})
+
+endif()
diff --git a/opencl/test/unit_test/offline_compiler/gen9/skl/offline_compiler_tests_skl.cpp b/opencl/test/unit_test/offline_compiler/gen9/skl/offline_compiler_tests_skl.cpp
new file mode 100644
index 0000000000..c676880965
--- /dev/null
+++ b/opencl/test/unit_test/offline_compiler/gen9/skl/offline_compiler_tests_skl.cpp
@@ -0,0 +1,63 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/compiler_interface/compiler_options/compiler_options.h"
+#include "shared/source/debug_settings/debug_settings_manager.h"
+#include "shared/test/common/helpers/debug_manager_state_restore.h"
+
+#include "opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h"
+#include "opencl/test/unit_test/offline_compiler/offline_compiler_tests.h"
+#include "test.h"
+
+namespace NEO {
+
+using MockOfflineCompilerSklTests = ::testing::Test;
+
+SKLTEST_F(MockOfflineCompilerSklTests, GivenSklWhenParseDebugSettingsThenStatelessToStatefullOptimizationIsEnabled) {
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = "skl";
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+    mockOfflineCompiler.parseDebugSettings();
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
+    EXPECT_NE(std::string::npos, found);
+}
+
+SKLTEST_F(MockOfflineCompilerSklTests, GivenSklAndDisabledViaDebugThenStatelessToStatefullOptimizationDisabled) {
+    DebugManagerStateRestore stateRestore;
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = "skl";
+    DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(0);
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+    mockOfflineCompiler.setStatelessToStatefullBufferOffsetFlag();
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
+    EXPECT_EQ(std::string::npos, found);
+}
+
+SKLTEST_F(MockOfflineCompilerSklTests, givenSklWhenAppendExtraInternalOptionsThenForceEmuInt32DivRemSPIsNotApplied) {
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = "skl";
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    mockOfflineCompiler.appendExtraInternalOptions(mockOfflineCompiler.hwInfo, internalOptions);
+    size_t found = internalOptions.find(NEO::CompilerOptions::forceEmuInt32DivRemSP.data());
+    EXPECT_EQ(std::string::npos, found);
+}
+
+SKLTEST_F(MockOfflineCompilerSklTests, givenSklWhenAppendExtraInternalOptionsThenGreaterThan4gbBuffersRequiredIsNotSet) {
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.deviceName = "skl";
+    mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+    std::string internalOptions = mockOfflineCompiler.internalOptions;
+    mockOfflineCompiler.forceStatelessToStatefulOptimization = false;
+    mockOfflineCompiler.appendExtraInternalOptions(mockOfflineCompiler.hwInfo, internalOptions);
+    size_t found = internalOptions.find(NEO::CompilerOptions::greaterThan4gbBuffersRequired.data());
+    EXPECT_EQ(std::string::npos, found);
+}
+
+} // namespace NEO
\ No newline at end of file
diff --git a/opencl/test/unit_test/offline_compiler/main.cpp b/opencl/test/unit_test/offline_compiler/main.cpp
index a24847044d..88cbc36e91 100644
--- a/opencl/test/unit_test/offline_compiler/main.cpp
+++ b/opencl/test/unit_test/offline_compiler/main.cpp
@@ -19,6 +19,7 @@ const char *fSeparator = "/";
 #endif
 
 Environment *gEnvironment;
+extern GFXCORE_FAMILY renderCoreFamily;
 
 std::string getRunPath() {
     char *cwd;
@@ -79,6 +80,15 @@ int main(int argc, char **argv) {
         }
     }
 
+    for (unsigned int productId = 0; productId < IGFX_MAX_PRODUCT; ++productId) {
+        if (NEO::hardwarePrefix[productId] && (0 == strcmp(devicePrefix.c_str(), NEO::hardwarePrefix[productId]))) {
+            if (NEO::hardwareInfoTable[productId]) {
+                renderCoreFamily = NEO::hardwareInfoTable[productId]->platform.eRenderCoreFamily;
+                break;
+            }
+        }
+    }
+
     // we look for test files always relative to binary location
     // this simplifies multi-process execution and using different
     // working directories
diff --git a/opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h b/opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h
index 8294ee1b81..a16d9d6d99 100644
--- a/opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h
+++ b/opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h
@@ -14,18 +14,20 @@ namespace NEO {
 
 class MockOfflineCompiler : public OfflineCompiler {
   public:
+    using OfflineCompiler::appendExtraInternalOptions;
     using OfflineCompiler::argHelper;
     using OfflineCompiler::deviceName;
     using OfflineCompiler::elfBinary;
     using OfflineCompiler::fclDeviceCtx;
+    using OfflineCompiler::forceStatelessToStatefulOptimization;
     using OfflineCompiler::genBinary;
     using OfflineCompiler::genBinarySize;
     using OfflineCompiler::generateFilePathForIr;
     using OfflineCompiler::generateOptsSuffix;
-    using OfflineCompiler::getHardwareInfo;
     using OfflineCompiler::getStringWithinDelimiters;
     using OfflineCompiler::hwInfo;
     using OfflineCompiler::igcDeviceCtx;
+    using OfflineCompiler::initHardwareInfo;
     using OfflineCompiler::inputFileLlvm;
     using OfflineCompiler::inputFileSpirV;
     using OfflineCompiler::internalOptions;
@@ -37,6 +39,7 @@ class MockOfflineCompiler : public OfflineCompiler {
     using OfflineCompiler::outputFile;
     using OfflineCompiler::parseCommandLine;
     using OfflineCompiler::parseDebugSettings;
+    using OfflineCompiler::setStatelessToStatefullBufferOffsetFlag;
     using OfflineCompiler::sourceCode;
     using OfflineCompiler::storeBinary;
     using OfflineCompiler::updateBuildLog;
diff --git a/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp b/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp
index 10ef99091f..919964302d 100644
--- a/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp
+++ b/opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp
@@ -765,43 +765,6 @@ TEST(OfflineCompilerTest, givenStatelessToStatefullOptimizationEnabledWhenDebugS
     EXPECT_NE(std::string::npos, found);
 }
 
-TEST(OfflineCompilerTest, GivenBdwThenStatelessToStatefullOptimizationIsDisabled) {
-    DebugManagerStateRestore stateRestore;
-    MockOfflineCompiler mockOfflineCompiler;
-    mockOfflineCompiler.deviceName = "bdw";
-
-    mockOfflineCompiler.parseDebugSettings();
-
-    std::string internalOptions = mockOfflineCompiler.internalOptions;
-    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
-    EXPECT_EQ(std::string::npos, found);
-}
-
-TEST(OfflineCompilerTest, GivenSklThenStatelessToStatefullOptimizationIsEnabled) {
-    DebugManagerStateRestore stateRestore;
-    MockOfflineCompiler mockOfflineCompiler;
-    mockOfflineCompiler.deviceName = "skl";
-
-    mockOfflineCompiler.parseDebugSettings();
-
-    std::string internalOptions = mockOfflineCompiler.internalOptions;
-    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
-    EXPECT_NE(std::string::npos, found);
-}
-
-TEST(OfflineCompilerTest, GivenSklAndDisabledViaDebugThenStatelessToStatefullOptimizationDisabled) {
-    DebugManagerStateRestore stateRestore;
-    MockOfflineCompiler mockOfflineCompiler;
-    mockOfflineCompiler.deviceName = "skl";
-    DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(0);
-
-    mockOfflineCompiler.parseDebugSettings();
-
-    std::string internalOptions = mockOfflineCompiler.internalOptions;
-    size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
-    EXPECT_EQ(std::string::npos, found);
-}
-
 TEST(OfflineCompilerTest, GivenDelimitersWhenGettingStringThenParseIsCorrect) {
     auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
     ASSERT_NE(nullptr, mockOfflineCompiler);
@@ -851,9 +814,9 @@ TEST(OfflineCompilerTest, GivenValidParamWhenGettingHardwareInfoThenSuccessIsRet
     auto mockOfflineCompiler = std::unique_ptr<MockOfflineCompiler>(new MockOfflineCompiler());
     ASSERT_NE(nullptr, mockOfflineCompiler);
 
-    EXPECT_EQ(CL_INVALID_DEVICE, mockOfflineCompiler->getHardwareInfo("invalid"));
+    EXPECT_EQ(CL_INVALID_DEVICE, mockOfflineCompiler->initHardwareInfo("invalid"));
     EXPECT_EQ(PRODUCT_FAMILY::IGFX_UNKNOWN, mockOfflineCompiler->getHardwareInfo().platform.eProductFamily);
-    EXPECT_EQ(CL_SUCCESS, mockOfflineCompiler->getHardwareInfo(gEnvironment->devicePrefix.c_str()));
+    EXPECT_EQ(CL_SUCCESS, mockOfflineCompiler->initHardwareInfo(gEnvironment->devicePrefix.c_str()));
     EXPECT_NE(PRODUCT_FAMILY::IGFX_UNKNOWN, mockOfflineCompiler->getHardwareInfo().platform.eProductFamily);
 }
 
@@ -1472,7 +1435,7 @@ TEST(OfflineCompilerTest, givenNoRevisionIdWhenCompilerIsInitializedThenHwInfoHa
         "-device",
         gEnvironment->devicePrefix.c_str()};
 
-    mockOfflineCompiler->getHardwareInfo(gEnvironment->devicePrefix.c_str());
+    mockOfflineCompiler->initHardwareInfo(gEnvironment->devicePrefix.c_str());
     auto revId = mockOfflineCompiler->hwInfo.platform.usRevId;
     int retVal = mockOfflineCompiler->initialize(argv.size(), argv);
     EXPECT_EQ(OfflineCompiler::ErrorCode::SUCCESS, retVal);
@@ -1512,36 +1475,6 @@ TEST(OfflineCompilerTest, whenDeviceIsSpecifiedThenDefaultConfigFromTheDeviceIsU
     EXPECT_EQ(euPerSubSliceCount * subSlicePerSliceCount * sliceCount, hwInfo.gtSystemInfo.EUCount);
 }
 
-struct WorkaroundApplicableForDevice {
-    const char *deviceName;
-    bool applicable;
-};
-
-using OfflineCompilerTestWithParams = testing::TestWithParam<WorkaroundApplicableForDevice>;
-
-TEST_P(OfflineCompilerTestWithParams, givenRklWhenExtraSettingsResolvedThenForceEmuInt32DivRemSPIsApplied) {
-    WorkaroundApplicableForDevice params = GetParam();
-    MockOfflineCompiler mockOfflineCompiler;
-    mockOfflineCompiler.deviceName = params.deviceName;
-
-    mockOfflineCompiler.parseDebugSettings();
-
-    std::string internalOptions = mockOfflineCompiler.internalOptions;
-    size_t found = internalOptions.find(NEO::CompilerOptions::forceEmuInt32DivRemSP.data());
-    if (params.applicable) {
-        EXPECT_NE(std::string::npos, found);
-    } else {
-        EXPECT_EQ(std::string::npos, found);
-    }
-}
-
-WorkaroundApplicableForDevice workaroundApplicableForDeviceArray[] = {{"rkl", true}, {"dg1", false}, {"tgllp", false}};
-
-INSTANTIATE_TEST_CASE_P(
-    WorkaroundApplicable,
-    OfflineCompilerTestWithParams,
-    testing::ValuesIn(workaroundApplicableForDeviceArray));
-
 TEST(OclocCompile, whenDetectedPotentialInputTypeMismatchThenEmitsWarning) {
     std::string sourceOclC = "__kernel void k() { }";
     std::string sourceLlvmBc = NEO::llvmBcMagic.str();
@@ -1589,7 +1522,7 @@ TEST(OclocCompile, whenDetectedPotentialInputTypeMismatchThenEmitsWarning) {
             "-device",
             gEnvironment->devicePrefix.c_str()};
 
-        ocloc.getHardwareInfo(gEnvironment->devicePrefix.c_str());
+        ocloc.initHardwareInfo(gEnvironment->devicePrefix.c_str());
         int retVal = ocloc.initialize(argv.size(), argv);
         ASSERT_EQ(0, retVal);
 
@@ -1713,4 +1646,26 @@ TEST(OclocCompile, givenSpirvInputThenDontGenerateSpirvFile) {
     EXPECT_TRUE(compilerOutputExists("offline_compiler_test/binary_with_zeroes", "bin"));
     EXPECT_FALSE(compilerOutputExists("offline_compiler_test/binary_with_zeroes", "spv"));
 }
+
+TEST(OfflineCompilerTest, GivenDebugFlagWhenSetStatelessToStatefullBufferOffsetFlagThenStatelessToStatefullOptimizationIsSetCorrectly) {
+    DebugManagerStateRestore stateRestore;
+    MockOfflineCompiler mockOfflineCompiler;
+    {
+        DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(0);
+        mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+        mockOfflineCompiler.setStatelessToStatefullBufferOffsetFlag();
+        std::string internalOptions = mockOfflineCompiler.internalOptions;
+        size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
+        EXPECT_EQ(std::string::npos, found);
+    }
+    {
+        DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(1);
+        mockOfflineCompiler.initHardwareInfo(mockOfflineCompiler.deviceName);
+        mockOfflineCompiler.setStatelessToStatefullBufferOffsetFlag();
+        std::string internalOptions = mockOfflineCompiler.internalOptions;
+        size_t found = internalOptions.find(NEO::CompilerOptions::hasBufferOffsetArg.data());
+        EXPECT_NE(std::string::npos, found);
+    }
+}
+
 } // namespace NEO
diff --git a/shared/offline_compiler/source/CMakeLists.txt b/shared/offline_compiler/source/CMakeLists.txt
index 3c7927e093..a06f9a6504 100644
--- a/shared/offline_compiler/source/CMakeLists.txt
+++ b/shared/offline_compiler/source/CMakeLists.txt
@@ -29,6 +29,7 @@ set(CLOC_LIB_SRCS_LIB
     ${NEO_SHARED_DIRECTORY}/helpers/abort.cpp
     ${NEO_SHARED_DIRECTORY}/helpers/compiler_hw_info_config.h
     ${NEO_SHARED_DIRECTORY}/helpers/compiler_hw_info_config.cpp
+    ${NEO_SHARED_DIRECTORY}/helpers/compiler_hw_info_config_base.inl
     ${NEO_SHARED_DIRECTORY}/helpers/compiler_hw_info_config_bdw_and_later.inl
     ${NEO_SHARED_DIRECTORY}/helpers/compiler_options_parser.cpp
     ${NEO_SHARED_DIRECTORY}/helpers/compiler_options_parser.h
@@ -50,7 +51,6 @@ set(CLOC_LIB_SRCS_LIB
     ${OCLOC_DIRECTORY}/source/decoder/helper.h
     ${OCLOC_DIRECTORY}/source/decoder/iga_wrapper.h
     ${OCLOC_DIRECTORY}/source/decoder/translate_platform_base.h
-    ${OCLOC_DIRECTORY}/source${BRANCH_DIR_SUFFIX}extra_settings.cpp
     ${OCLOC_DIRECTORY}/source/multi_command.cpp
     ${OCLOC_DIRECTORY}/source/multi_command.h
     ${OCLOC_DIRECTORY}/source/ocloc_api.cpp
@@ -116,6 +116,7 @@ macro(macro_for_each_platform)
     foreach(BRANCH ${BRANCH_DIR_LIST})
       foreach(SRC_FILE ${NEO_SOURCE_DIR}/shared/source${BRANCH}${GEN_TYPE_LOWER}/definitions${BRANCH_DIR_SUFFIX}hw_info_setup_${PLATFORM_IT_LOWER}.inl
                 ${NEO_SOURCE_DIR}/shared/source${BRANCH}${GEN_TYPE_LOWER}${BRANCH_DIR}hw_info_${PLATFORM_IT_LOWER}.cpp
+                ${NEO_SOURCE_DIR}/shared/source${BRANCH}${GEN_TYPE_LOWER}/compiler_hw_info_config_${PLATFORM_IT_LOWER}.inl
         )
         if(EXISTS ${SRC_FILE})
           list(APPEND CLOC_LIB_SRCS_LIB ${SRC_FILE})
diff --git a/shared/offline_compiler/source/extra_settings.cpp b/shared/offline_compiler/source/extra_settings.cpp
deleted file mode 100644
index a52b0e78f9..0000000000
--- a/shared/offline_compiler/source/extra_settings.cpp
+++ /dev/null
@@ -1,20 +0,0 @@
-/*
- * Copyright (C) 2020-2021 Intel Corporation
- *
- * SPDX-License-Identifier: MIT
- *
- */
-
-#include "shared/offline_compiler/source/offline_compiler.h"
-#include "shared/source/os_interface/os_library.h"
-
-#include "compiler_options.h"
-
-namespace NEO {
-void OfflineCompiler::resolveExtraSettings() {
-    if (deviceName == "rkl") {
-        CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::forceEmuInt32DivRemSP);
-    }
-}
-
-} // namespace NEO
diff --git a/shared/offline_compiler/source/offline_compiler.cpp b/shared/offline_compiler/source/offline_compiler.cpp
index 6d040b41ed..ce0b402f6e 100644
--- a/shared/offline_compiler/source/offline_compiler.cpp
+++ b/shared/offline_compiler/source/offline_compiler.cpp
@@ -316,7 +316,7 @@ std::string &OfflineCompiler::getBuildLog() {
     return buildLog;
 }
 
-int OfflineCompiler::getHardwareInfo(std::string deviceName) {
+int OfflineCompiler::initHardwareInfo(std::string deviceName) {
     int retVal = INVALID_DEVICE;
 
     overridePlatformName(deviceName);
@@ -433,7 +433,7 @@ int OfflineCompiler::initialize(size_t numArgs, const std::vector<std::string> &
         }
     }
 
-    retVal = deviceName.empty() ? SUCCESS : getHardwareInfo(deviceName.c_str());
+    retVal = deviceName.empty() ? SUCCESS : initHardwareInfo(deviceName.c_str());
     if (retVal != SUCCESS) {
         argHelper->printf("Error: Cannot get HW Info for device %s.\n", deviceName.c_str());
         return retVal;
@@ -450,6 +450,7 @@ int OfflineCompiler::initialize(size_t numArgs, const std::vector<std::string> &
         CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::enableImageSupport);
     } else {
         appendExtensionsToInternalOptions(hwInfo, options, internalOptions);
+        appendExtraInternalOptions(hwInfo, internalOptions);
     }
 
     parseDebugSettings();
@@ -698,8 +699,9 @@ int OfflineCompiler::parseCommandLine(size_t numArgs, const std::vector<std::str
 
 void OfflineCompiler::setStatelessToStatefullBufferOffsetFlag() {
     bool isStatelessToStatefulBufferOffsetSupported = true;
-    if (deviceName == "bdw") {
-        isStatelessToStatefulBufferOffsetSupported = false;
+    if (!deviceName.empty()) {
+        const auto &compilerHwInfoConfig = *CompilerHwInfoConfig::get(hwInfo.platform.eProductFamily);
+        isStatelessToStatefulBufferOffsetSupported = compilerHwInfoConfig.isStatelessToStatefulBufferOffsetSupported();
     }
     if (DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.get() != -1) {
         isStatelessToStatefulBufferOffsetSupported = DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.get() != 0;
@@ -709,9 +711,18 @@ void OfflineCompiler::setStatelessToStatefullBufferOffsetFlag() {
     }
 }
 
+void OfflineCompiler::appendExtraInternalOptions(const HardwareInfo &hwInfo, std::string &internalOptions) {
+    const auto &compilerHwInfoConfig = *CompilerHwInfoConfig::get(hwInfo.platform.eProductFamily);
+    if (compilerHwInfoConfig.isForceToStatelessRequired() && !forceStatelessToStatefulOptimization) {
+        CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::greaterThan4gbBuffersRequired);
+    }
+    if (compilerHwInfoConfig.isForceEmuInt32DivRemSPRequired()) {
+        CompilerOptions::concatenateAppend(internalOptions, CompilerOptions::forceEmuInt32DivRemSP);
+    }
+}
+
 void OfflineCompiler::parseDebugSettings() {
     setStatelessToStatefullBufferOffsetFlag();
-    resolveExtraSettings();
 }
 
 std::string OfflineCompiler::parseBinAsCharArray(uint8_t *binary, size_t size, std::string &fileName) {
diff --git a/shared/offline_compiler/source/offline_compiler.h b/shared/offline_compiler/source/offline_compiler.h
index 2ea1cb91ca..3506ebeb3b 100644
--- a/shared/offline_compiler/source/offline_compiler.h
+++ b/shared/offline_compiler/source/offline_compiler.h
@@ -90,12 +90,12 @@ class OfflineCompiler {
   protected:
     OfflineCompiler();
 
-    int getHardwareInfo(std::string deviceName);
+    int initHardwareInfo(std::string deviceName);
     std::string getStringWithinDelimiters(const std::string &src);
     int initialize(size_t numArgs, const std::vector<std::string> &allArgs, bool dumpFiles);
     int parseCommandLine(size_t numArgs, const std::vector<std::string> &allArgs);
     void setStatelessToStatefullBufferOffsetFlag();
-    void resolveExtraSettings();
+    void appendExtraInternalOptions(const HardwareInfo &hwInfo, std::string &internalOptions);
     void parseDebugSettings();
     void storeBinary(char *&pDst, size_t &dstSize, const void *pSrc, const size_t srcSize);
     MOCKABLE_VIRTUAL int buildSourceCode();
diff --git a/shared/source/enable_gens.cmake b/shared/source/enable_gens.cmake
index 43df419b6d..c0e36c01bf 100644
--- a/shared/source/enable_gens.cmake
+++ b/shared/source/enable_gens.cmake
@@ -68,6 +68,10 @@ macro(macro_for_each_platform)
       if(EXISTS ${SRC_FILE})
         list(APPEND ${GEN_TYPE}_SRC_LINK_BASE ${SRC_FILE})
       endif()
+      set(SRC_FILE ${CMAKE_CURRENT_SOURCE_DIR}${BRANCH_DIR}${GEN_TYPE_LOWER}/compiler_hw_info_config_${PLATFORM_IT_LOWER}.inl)
+      if(EXISTS ${SRC_FILE})
+        list(APPEND ${GEN_TYPE}_SRC_LINK_BASE ${SRC_FILE})
+      endif()
       set(SRC_FILE ${CMAKE_CURRENT_SOURCE_DIR}/os_interface/linux/local${BRANCH_DIR_SUFFIX}${PLATFORM_IT_LOWER}/enable_local_memory_helper_${PLATFORM_IT_LOWER}.cpp)
       if(EXISTS ${SRC_FILE})
         list(APPEND ${GEN_TYPE}_SRC_LINK_LINUX ${SRC_FILE})
diff --git a/shared/source/gen11/enable_gen11.cpp b/shared/source/gen11/enable_gen11.cpp
index 18ec0eb741..63249afe4d 100644
--- a/shared/source/gen11/enable_gen11.cpp
+++ b/shared/source/gen11/enable_gen11.cpp
@@ -6,6 +6,7 @@
  */
 
 #include "shared/source/gen11/hw_cmds.h"
+#include "shared/source/helpers/compiler_hw_info_config_base.inl"
 #include "shared/source/helpers/compiler_hw_info_config_bdw_and_later.inl"
 #include "shared/source/helpers/enable_product.inl"
 #include "shared/source/os_interface/hw_info_config.h"
diff --git a/shared/source/gen12lp/compiler_hw_info_config_rkl.inl b/shared/source/gen12lp/compiler_hw_info_config_rkl.inl
new file mode 100644
index 0000000000..09e60da493
--- /dev/null
+++ b/shared/source/gen12lp/compiler_hw_info_config_rkl.inl
@@ -0,0 +1,11 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+template <>
+bool CompilerHwInfoConfigHw<IGFX_ROCKETLAKE>::isForceEmuInt32DivRemSPRequired() const {
+    return true;
+}
diff --git a/shared/source/gen12lp/enable_gen12lp.cpp b/shared/source/gen12lp/enable_gen12lp.cpp
index 3e91aaf283..ca7608b3e2 100644
--- a/shared/source/gen12lp/enable_gen12lp.cpp
+++ b/shared/source/gen12lp/enable_gen12lp.cpp
@@ -6,6 +6,7 @@
  */
 
 #include "shared/source/gen12lp/hw_cmds.h"
+#include "shared/source/helpers/compiler_hw_info_config_base.inl"
 #include "shared/source/helpers/compiler_hw_info_config_bdw_and_later.inl"
 #include "shared/source/helpers/enable_product.inl"
 #include "shared/source/os_interface/hw_info_config.h"
@@ -22,6 +23,8 @@ static EnableCompilerHwInfoConfig<IGFX_DG1> enableCompilerHwInfoConfigDG1;
 #endif
 #ifdef SUPPORT_RKL
 static EnableGfxProductHw<IGFX_ROCKETLAKE> enableGfxProductHwRKL;
+
+#include "shared/source/gen12lp/compiler_hw_info_config_rkl.inl"
 static EnableCompilerHwInfoConfig<IGFX_ROCKETLAKE> enableCompilerHwInfoConfigRKL;
 #endif
 #ifdef SUPPORT_ADLS
diff --git a/shared/source/gen8/compiler_hw_info_config_bdw.inl b/shared/source/gen8/compiler_hw_info_config_bdw.inl
new file mode 100644
index 0000000000..1e72b405a9
--- /dev/null
+++ b/shared/source/gen8/compiler_hw_info_config_bdw.inl
@@ -0,0 +1,11 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+template <>
+bool CompilerHwInfoConfigHw<IGFX_BROADWELL>::isStatelessToStatefulBufferOffsetSupported() const {
+    return false;
+}
diff --git a/shared/source/gen8/enable_gen8.cpp b/shared/source/gen8/enable_gen8.cpp
index b71949716b..132fb4cc3a 100644
--- a/shared/source/gen8/enable_gen8.cpp
+++ b/shared/source/gen8/enable_gen8.cpp
@@ -6,6 +6,7 @@
  */
 
 #include "shared/source/gen8/hw_cmds.h"
+#include "shared/source/helpers/compiler_hw_info_config_base.inl"
 #include "shared/source/helpers/compiler_hw_info_config_bdw_and_later.inl"
 #include "shared/source/helpers/enable_product.inl"
 #include "shared/source/os_interface/hw_info_config.h"
@@ -14,6 +15,8 @@ namespace NEO {
 
 #ifdef SUPPORT_BDW
 static EnableGfxProductHw<IGFX_BROADWELL> enableGfxProductHwBDW;
+
+#include "shared/source/gen8/compiler_hw_info_config_bdw.inl"
 static EnableCompilerHwInfoConfig<IGFX_BROADWELL> enableCompilerHwInfoConfigBDW;
 #endif
 
diff --git a/shared/source/gen9/enable_gen9.cpp b/shared/source/gen9/enable_gen9.cpp
index 71a0cbf057..effb663de5 100644
--- a/shared/source/gen9/enable_gen9.cpp
+++ b/shared/source/gen9/enable_gen9.cpp
@@ -6,6 +6,7 @@
  */
 
 #include "shared/source/gen9/hw_cmds.h"
+#include "shared/source/helpers/compiler_hw_info_config_base.inl"
 #include "shared/source/helpers/compiler_hw_info_config_bdw_and_later.inl"
 #include "shared/source/helpers/enable_product.inl"
 #include "shared/source/os_interface/hw_info_config.h"
diff --git a/shared/source/helpers/CMakeLists.txt b/shared/source/helpers/CMakeLists.txt
index 36a7b752be..b9c94c3a14 100644
--- a/shared/source/helpers/CMakeLists.txt
+++ b/shared/source/helpers/CMakeLists.txt
@@ -31,6 +31,7 @@ set(NEO_CORE_HELPERS
     ${CMAKE_CURRENT_SOURCE_DIR}/common_types.h
     ${CMAKE_CURRENT_SOURCE_DIR}/compiler_hw_info_config.h
     ${CMAKE_CURRENT_SOURCE_DIR}/compiler_hw_info_config.cpp
+    ${CMAKE_CURRENT_SOURCE_DIR}/compiler_hw_info_config_base.inl
     ${CMAKE_CURRENT_SOURCE_DIR}/compiler_hw_info_config_bdw_and_later.inl
     ${CMAKE_CURRENT_SOURCE_DIR}/compiler_options_parser.cpp
     ${CMAKE_CURRENT_SOURCE_DIR}/compiler_options_parser.h
diff --git a/shared/source/helpers/compiler_hw_info_config.h b/shared/source/helpers/compiler_hw_info_config.h
index dd451b396e..3b429d5812 100644
--- a/shared/source/helpers/compiler_hw_info_config.h
+++ b/shared/source/helpers/compiler_hw_info_config.h
@@ -24,6 +24,9 @@ class CompilerHwInfoConfig {
     }
 
     virtual bool isMidThreadPreemptionSupported(const HardwareInfo &hwInfo) const = 0;
+    virtual bool isForceToStatelessRequired() const = 0;
+    virtual bool isForceEmuInt32DivRemSPRequired() const = 0;
+    virtual bool isStatelessToStatefulBufferOffsetSupported() const = 0;
 };
 
 template <PRODUCT_FAMILY gfxProduct>
@@ -35,6 +38,9 @@ class CompilerHwInfoConfigHw : public CompilerHwInfoConfig {
     }
 
     bool isMidThreadPreemptionSupported(const HardwareInfo &hwInfo) const override;
+    bool isForceToStatelessRequired() const override;
+    bool isForceEmuInt32DivRemSPRequired() const override;
+    bool isStatelessToStatefulBufferOffsetSupported() const override;
 
   protected:
     CompilerHwInfoConfigHw() = default;
diff --git a/shared/source/helpers/compiler_hw_info_config_base.inl b/shared/source/helpers/compiler_hw_info_config_base.inl
new file mode 100644
index 0000000000..1f93f7ac5f
--- /dev/null
+++ b/shared/source/helpers/compiler_hw_info_config_base.inl
@@ -0,0 +1,27 @@
+/*
+ * Copyright (C) 2021 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma once
+
+#include "shared/source/helpers/compiler_hw_info_config.h"
+
+namespace NEO {
+template <PRODUCT_FAMILY gfxProduct>
+bool CompilerHwInfoConfigHw<gfxProduct>::isForceToStatelessRequired() const {
+    return false;
+}
+
+template <PRODUCT_FAMILY gfxProduct>
+bool CompilerHwInfoConfigHw<gfxProduct>::isForceEmuInt32DivRemSPRequired() const {
+    return false;
+}
+template <PRODUCT_FAMILY gfxProduct>
+bool CompilerHwInfoConfigHw<gfxProduct>::isStatelessToStatefulBufferOffsetSupported() const {
+    return true;
+}
+
+} // namespace NEO
diff --git a/shared/source/xe_hp_core/enable_xe_hp_core.cpp b/shared/source/xe_hp_core/enable_xe_hp_core.cpp
index 8dea7798ff..9fd1463eb2 100644
--- a/shared/source/xe_hp_core/enable_xe_hp_core.cpp
+++ b/shared/source/xe_hp_core/enable_xe_hp_core.cpp
@@ -5,6 +5,7 @@
  *
  */
 
+#include "shared/source/helpers/compiler_hw_info_config_base.inl"
 #include "shared/source/helpers/compiler_hw_info_config_bdw_and_later.inl"
 #include "shared/source/helpers/enable_product.inl"
 #include "shared/source/os_interface/hw_info_config.h"
diff --git a/shared/test/common/test_macros/test_excludes.cpp b/shared/test/common/test_macros/test_excludes.cpp
index f41756ba64..df412e6c44 100644
--- a/shared/test/common/test_macros/test_excludes.cpp
+++ b/shared/test/common/test_macros/test_excludes.cpp
@@ -16,6 +16,8 @@
 
 using namespace NEO;
 
+GFXCORE_FAMILY renderCoreFamily = {};
+
 static std::unique_ptr<std::map<std::string, std::unordered_set<uint32_t>>> pProductExcludesPerTest;
 static std::unique_ptr<std::map<std::string, std::unordered_set<uint32_t>>> pGfxExcludesPerTest;
 
diff --git a/shared/test/unit_test/main.cpp b/shared/test/unit_test/main.cpp
index 5124235709..c58ec57a54 100644
--- a/shared/test/unit_test/main.cpp
+++ b/shared/test/unit_test/main.cpp
@@ -66,7 +66,7 @@ bool disabled = false;
 using namespace NEO;
 
 PRODUCT_FAMILY productFamily = DEFAULT_TEST_PLATFORM::hwInfo.platform.eProductFamily;
-GFXCORE_FAMILY renderCoreFamily = DEFAULT_TEST_PLATFORM::hwInfo.platform.eRenderCoreFamily;
+extern GFXCORE_FAMILY renderCoreFamily;
 
 extern std::string lastTest;
 bool generateRandomInput = false;