Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Refactor per platform extra settings in ocloc

Browse Source 
Signed-off-by: Kamil Kopryk <kamil.kopryk@intel.com>
Related-To: NEO-6382
This commit is contained in:
Kamil Kopryk
2021-10-26 15:25:20 +00:00
committed by Compute-Runtime-Automation
parent 515130dbe8
commit 03540d5301
35 changed files with 359 additions and 138 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
level_zero/core/test/unit_tests/main.cpp
View File

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

6
opencl/test/unit_test/gen12lp/rkl/test_hw_info_config_rkl.cpp
View File

@ -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());
}

6
opencl/test/unit_test/gen8/bdw/test_hw_info_config_bdw.cpp
View File

@ -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());
}

2
opencl/test/unit_test/main.cpp
View File

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

5
opencl/test/unit_test/offline_compiler/CMakeLists.txt
View File

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

13
opencl/test/unit_test/offline_compiler/gen12lp/CMakeLists.txt Normal file
View File

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

13
opencl/test/unit_test/offline_compiler/gen12lp/rkl/CMakeLists.txt Normal file
View File

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

26
opencl/test/unit_test/offline_compiler/gen12lp/rkl/offline_compiler_tests_rkl.cpp Normal file
View File

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

3
opencl/test/unit_test/offline_compiler/gen8/CMakeLists.txt
View File

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

14
opencl/test/unit_test/offline_compiler/gen8/bdw/CMakeLists.txt Normal file
View File

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

50
opencl/test/unit_test/offline_compiler/gen8/bdw/offline_compiler_tests_bdw.cpp Normal file
View File

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

33
opencl/test/unit_test/offline_compiler/gen8/offline_compiler_tests_gen8.cpp
View File

@ -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

14
opencl/test/unit_test/offline_compiler/gen9/CMakeLists.txt Normal file
View File

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

14
opencl/test/unit_test/offline_compiler/gen9/skl/CMakeLists.txt Normal file
View File

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

63
opencl/test/unit_test/offline_compiler/gen9/skl/offline_compiler_tests_skl.cpp Normal file
View File

@ -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

10
opencl/test/unit_test/offline_compiler/main.cpp
View File

@ -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

5
opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h
View File

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

97
opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp
View File

@ -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

3
shared/offline_compiler/source/CMakeLists.txt
View File

@ -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})

20
shared/offline_compiler/source/extra_settings.cpp
View File

@ -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

21
shared/offline_compiler/source/offline_compiler.cpp
View File

@ -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) {

4
shared/offline_compiler/source/offline_compiler.h
View File

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

4
shared/source/enable_gens.cmake
View File

@ -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})

1
shared/source/gen11/enable_gen11.cpp
View File

@ -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"

11
shared/source/gen12lp/compiler_hw_info_config_rkl.inl Normal file
View File

@ -0,0 +1,11 @@
/*
* Copyright (C) 2021 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
template <>
bool CompilerHwInfoConfigHw<IGFX_ROCKETLAKE>::isForceEmuInt32DivRemSPRequired() const {
return true;
}

3
shared/source/gen12lp/enable_gen12lp.cpp
View File

@ -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

11
shared/source/gen8/compiler_hw_info_config_bdw.inl Normal file
View File

@ -0,0 +1,11 @@
/*
* Copyright (C) 2021 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
template <>
bool CompilerHwInfoConfigHw<IGFX_BROADWELL>::isStatelessToStatefulBufferOffsetSupported() const {
return false;
}

3
shared/source/gen8/enable_gen8.cpp
View File

@ -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

1
shared/source/gen9/enable_gen9.cpp
View File

@ -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"

1
shared/source/helpers/CMakeLists.txt
View File

@ -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

6
shared/source/helpers/compiler_hw_info_config.h
View File

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

27
shared/source/helpers/compiler_hw_info_config_base.inl Normal file
View File

@ -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

1
shared/source/xe_hp_core/enable_xe_hp_core.cpp
View File

@ -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"

2
shared/test/common/test_macros/test_excludes.cpp
View File

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

2
shared/test/unit_test/main.cpp
View File

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