diff --git a/shared/source/compiler_interface/compiler_interface.cpp b/shared/source/compiler_interface/compiler_interface.cpp index 7bed73249d..1f5551c6f4 100644 --- a/shared/source/compiler_interface/compiler_interface.cpp +++ b/shared/source/compiler_interface/compiler_interface.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2018-2024 Intel Corporation + * Copyright (C) 2018-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -33,12 +33,6 @@ namespace NEO { SpinLock CompilerInterface::spinlock; -enum CachingMode { - None, - Direct, - PreProcess -}; - void TranslationOutput::makeCopy(MemAndSize &dst, CIF::Builtins::BufferSimple *src) { if ((nullptr == src) || (src->GetSizeRaw() == 0)) { dst.mem.reset(); @@ -73,19 +67,11 @@ TranslationOutput::ErrorCode CompilerInterface::build( intermediateCodeType = input.preferredIntermediateType; } - CachingMode cachingMode = None; - - if (cache != nullptr && cache->getConfig().enabled) { - if ((srcCodeType == IGC::CodeType::oclC) && (std::strstr(input.src.begin(), "#include") == nullptr)) { - cachingMode = CachingMode::Direct; - } else { - cachingMode = CachingMode::PreProcess; - } - } + CachingMode cachingMode = CompilerCacheHelper::getCachingMode(cache.get(), srcCodeType, input.src); std::string kernelFileHash; const auto &igc = *getIgc(&device); - if (cachingMode == CachingMode::Direct) { + if (cachingMode == CachingMode::direct) { kernelFileHash = cache->getCachedFileName(device.getHardwareInfo(), input.src, input.apiOptions, @@ -141,7 +127,7 @@ TranslationOutput::ErrorCode CompilerInterface::build( intermediateCodeType = srcCodeType; } - if (cachingMode == CachingMode::PreProcess) { + if (cachingMode == CachingMode::preProcess) { const ArrayRef irRef(intermediateRepresentation->GetMemory(), intermediateRepresentation->GetSize()); const ArrayRef specIdsRef(idsBuffer->GetMemory(), idsBuffer->GetSize()); const ArrayRef specValuesRef(valuesBuffer->GetMemory(), valuesBuffer->GetSize()); @@ -736,4 +722,48 @@ bool CompilerCacheHelper::processPackedCacheBinary(ArrayRef archi return false; } + +CachingMode CompilerCacheHelper::getCachingMode(CompilerCache *compilerCache, IGC::CodeType::CodeType_t srcCodeType, const ArrayRef source) { + if (compilerCache == nullptr || !compilerCache->getConfig().enabled) { + return CachingMode::none; + } + + if (srcCodeType == IGC::CodeType::oclC && + validateIncludes(source, CompilerCacheHelper::whitelistedIncludes)) { + return CachingMode::direct; + } + + return CachingMode::preProcess; +} + +bool CompilerCacheHelper::validateIncludes(const ArrayRef source, const WhitelistedIncludesVec &whitelistedIncludes) { + const char *sourcePtr = source.begin(); + const char *sourceEnd = source.end(); + + while (sourcePtr < sourceEnd) { + const char *includePos = std::strstr(sourcePtr, "#include"); + if (includePos == nullptr) { + break; + } + + bool isKnownInclude = false; + for (const auto &knownInclude : whitelistedIncludes) { + if (std::strncmp(includePos, knownInclude.data(), knownInclude.size()) == 0) { + isKnownInclude = true; + break; + } + } + + if (!isKnownInclude) { + return false; + } + + sourcePtr = includePos + 1; + } + + return true; +} + +CompilerCacheHelper::WhitelistedIncludesVec CompilerCacheHelper::whitelistedIncludes{}; + } // namespace NEO \ No newline at end of file diff --git a/shared/source/compiler_interface/compiler_interface.h b/shared/source/compiler_interface/compiler_interface.h index f2bc773428..4c198e8e1b 100644 --- a/shared/source/compiler_interface/compiler_interface.h +++ b/shared/source/compiler_interface/compiler_interface.h @@ -1,5 +1,5 @@ /* - * Copyright (C) 2018-2024 Intel Corporation + * Copyright (C) 2018-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -8,6 +8,7 @@ #pragma once #include "shared/source/utilities/arrayref.h" #include "shared/source/utilities/spinlock.h" +#include "shared/source/utilities/stackvec.h" #include "cif/common/cif_main.h" #include "ocl_igc_interface/code_type.h" @@ -103,6 +104,12 @@ struct SpecConstantInfo { CIF::RAII::UPtr_t sizesBuffer; }; +enum class CachingMode { + none, + direct, + preProcess +}; + class CompilerInterface { public: CompilerInterface(); @@ -245,9 +252,14 @@ class CompilerCacheHelper { public: static void packAndCacheBinary(CompilerCache &compilerCache, const std::string &kernelFileHash, const NEO::TargetDevice &targetDevice, const NEO::TranslationOutput &translationOutput); static bool loadCacheAndSetOutput(CompilerCache &compilerCache, const std::string &kernelFileHash, NEO::TranslationOutput &output, const NEO::Device &device); + static CachingMode getCachingMode(CompilerCache *compilerCache, IGC::CodeType::CodeType_t srcCodeType, const ArrayRef source); protected: static bool processPackedCacheBinary(ArrayRef archive, TranslationOutput &output, const NEO::Device &device); + + using WhitelistedIncludesVec = StackVec; + static bool validateIncludes(const ArrayRef source, const WhitelistedIncludesVec &whitelistedIncludes); + static WhitelistedIncludesVec whitelistedIncludes; }; } // namespace NEO diff --git a/shared/test/unit_test/compiler_interface/compiler_cache_tests.cpp b/shared/test/unit_test/compiler_interface/compiler_cache_tests.cpp index 059fcec0a8..f0807a365a 100644 --- a/shared/test/unit_test/compiler_interface/compiler_cache_tests.cpp +++ b/shared/test/unit_test/compiler_interface/compiler_cache_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2019-2024 Intel Corporation + * Copyright (C) 2019-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -827,3 +827,142 @@ TEST_F(CompilerInterfaceOclElfCacheTest, GivenBinaryAndDebugDataWhenBuildingThen gEnvironment->fclPopDebugVars(); } + +class CompilerCacheHelperWhitelistedTest : public ::testing::Test, public CompilerCacheHelper { + public: + using CompilerCacheHelper::whitelistedIncludes; + + bool isValidIncludeFormat(const std::string_view &entry) { + size_t spacePos = entry.find(' '); + if (spacePos == std::string_view::npos) { + return false; + } + + std::string_view firstWord = entry.substr(0, spacePos); + if (firstWord != "#include") { + return false; + } + + std::string_view secondPart = entry.substr(spacePos + 1); + if (secondPart.empty()) { + return false; + } + + return true; + } +}; + +TEST_F(CompilerCacheHelperWhitelistedTest, GivenWhitelistedIncludesWhenCheckingFormatThenExpectValidIncludeFormat) { + for (const auto &entry : whitelistedIncludes) { + EXPECT_TRUE(isValidIncludeFormat(entry)) + << "Invalid format in whitelisted include: " << entry; + } +} + +TEST_F(CompilerCacheHelperWhitelistedTest, GivenWhitelistedIncludesWhenCheckingContentsThenExpectCorrectEntries) { + std::vector expectedEntries; + + for (const auto &expectedEntry : expectedEntries) { + auto it = std::find(whitelistedIncludes.begin(), whitelistedIncludes.end(), expectedEntry); + EXPECT_NE(it, whitelistedIncludes.end()) + << "Expected entry '" << expectedEntry << "' not found in whitelistedIncludes."; + } + + EXPECT_EQ(whitelistedIncludes.size(), expectedEntries.size()) + << "whitelistedIncludes contains unexpected entries."; +} + +class CompilerCacheHelperMockedWhitelistedIncludesTests : public ::testing::Test, public CompilerCacheHelper { + public: + using CompilerCacheHelper::whitelistedIncludes; + + CompilerCacheHelperMockedWhitelistedIncludesTests() : backup(&whitelistedIncludes, {"#include "}) {} + + void SetUp() override {} + void TearDown() override {} + + VariableBackup backup; +}; + +TEST_F(CompilerCacheHelperMockedWhitelistedIncludesTests, GivenSourceWithWhitelistedIncludeWhenCheckingWhitelistedIncludesThenReturnsTrue) { + const StackVec sources = { + "__kernel void k() {}", // no includes + "#include \n__kernel void k() {}", // whitelisted include + "", // empty source + " #include \n__kernel void k() {}", // leading spaces + "\t#include \n__kernel void k() {}", // leading tabs + "#include \n#include \n__kernel void k() {}", // multiple whitelisted + "\n#include \n__kernel void k() {}", // leading newline + "\r\n#include \n__kernel void k() {}" // leading carriage return + newline + }; + + for (const auto &source : sources) { + ArrayRef sourceRef(source, strlen(source)); + EXPECT_TRUE(CompilerCacheHelper::validateIncludes(sourceRef, CompilerCacheHelper::whitelistedIncludes)) + << "Failed for source: " << source; + } +} + +TEST_F(CompilerCacheHelperMockedWhitelistedIncludesTests, GivenSourceWithNonWhitelistedIncludeWhenCheckingWhitelistedIncludesThenReturnsFalse) { + const StackVec sources = { + "#include \n__kernel void k() {}", // non-whitelisted include + "#include \n#include \n__kernel void k() {}", // mixed includes - first whitelisted, second non-whitelisted + "#include \n#include \n__kernel void k() {}" // mixed includes - first non-whitelisted, second whitelisted + }; + + for (const auto &source : sources) { + ArrayRef sourceRef(source, strlen(source)); + EXPECT_FALSE(CompilerCacheHelper::validateIncludes(sourceRef, CompilerCacheHelper::whitelistedIncludes)) + << "Failed for source: " << source; + } +} + +TEST_F(CompilerCacheHelperMockedWhitelistedIncludesTests, GivenDisabledOrNullCacheWhenGettingCachingModeThenReturnsCachingModeNone) { + CompilerCache cache(CompilerCacheConfig{false, "", "", 0}); + const StackVec, 3> testCases = { + {"#include \n__kernel void k() {}", IGC::CodeType::oclC}, // disabled cache + {"__kernel void k() {}", IGC::CodeType::oclC}, // disabled cache, no includes + {"#include \n__kernel void k() {}", IGC::CodeType::oclC} // disabled cache, non-whitelisted include + }; + + for (const auto &testCase : testCases) { + ArrayRef sourceRef(testCase.first, strlen(testCase.first)); + EXPECT_EQ(CachingMode::none, CompilerCacheHelper::getCachingMode(&cache, testCase.second, sourceRef)) + << "Failed for source: " << testCase.first; + } + + EXPECT_EQ(CachingMode::none, CompilerCacheHelper::getCachingMode(nullptr, IGC::CodeType::oclC, ArrayRef())) + << "Failed for nullptr CompilerCache"; +} + +TEST_F(CompilerCacheHelperMockedWhitelistedIncludesTests, GivenEnabledCacheAndOclCWithWhitelistedIncludesWhenGettingCachingModeThenReturnsCachingModeDirect) { + CompilerCache cache(CompilerCacheConfig{true, "", "", 0}); + const StackVec, 2> testCases = { + {"__kernel void k() {}", IGC::CodeType::oclC}, // no includes + {"#include \n__kernel void k() {}", IGC::CodeType::oclC} // only whitelisted include + }; + + for (const auto &testCase : testCases) { + ArrayRef sourceRef(testCase.first, strlen(testCase.first)); + EXPECT_EQ(CachingMode::direct, CompilerCacheHelper::getCachingMode(&cache, testCase.second, sourceRef)) + << "Failed for source: " << testCase.first; + } +} + +TEST_F(CompilerCacheHelperMockedWhitelistedIncludesTests, GivenEnabledCacheAndNonWhitelistedIncludesOrNonOclCWhenGettingCachingModeThenReturnsCachingModePreProcess) { + CompilerCache cache(CompilerCacheConfig{true, "", "", 0}); + const StackVec, 6> testCases = { + {"#include \n__kernel void k() {}", IGC::CodeType::oclC}, // non-whitelisted include + {"#include \n#include \n__kernel void k() {}", IGC::CodeType::oclC}, // mixed includes + {"#include \n#include \n__kernel void k() {}", IGC::CodeType::oclC}, // mixed includes different order + {"__kernel void k() {}", IGC::CodeType::spirV}, // non-oclC type + {"#include \n__kernel void k() {}", IGC::CodeType::spirV}, // non-oclC type with whitelisted include + {"#include \n__kernel void k() {}", IGC::CodeType::spirV} // non-oclC type with non-whitelisted include + }; + + for (const auto &testCase : testCases) { + ArrayRef sourceRef(testCase.first, strlen(testCase.first)); + EXPECT_EQ(CachingMode::preProcess, CompilerCacheHelper::getCachingMode(&cache, testCase.second, sourceRef)) + << "Failed for source: " << testCase.first; + } +}