From 6ac6db7b04f9e3776510de4c79a20ac61c7a170e Mon Sep 17 00:00:00 2001 From: Patryk Wrobel Date: Fri, 11 Feb 2022 15:54:41 +0000 Subject: [PATCH] Store single generic IR in fatbinary when built from SPIR-V input This change prevents embedding identical SPIR-V section for each target requested in fatbinary build. Instead of duplicating SPIR-V, a new file called 'generic_ir' is added to AR archive. It contains SPIR-V, which was used to build fatbinary. Build fallback in runtime has been also adjusted - if 'generic_ir' file is defined in fatbinary and there is no matching binary, then this generic SPIR-V is used to rebuild for the requested target. Additionally, MockOclocArgumentHelper::loadDataFromFile() was adjusted to ensure null-termination of returned strings. This change also removes possible undefined behavior, which was related to reading names of files from AR archive. Previously, if filename was shorter than requested target name, we tried to read more memory than allowed. Related-To: NEO-6490 Signed-off-by: Patryk Wrobel --- .../mock/mock_argument_helper.h | 13 +- .../ocloc_fatbinary_tests.cpp | 184 +++++++++++++++++- .../offline_compiler/ocloc_fatbinary_tests.h | 23 ++- .../offline_compiler/offline_linker_tests.cpp | 26 +-- .../source/ocloc_fatbinary.cpp | 55 +++++- .../offline_compiler/source/ocloc_fatbinary.h | 4 + .../source/offline_linker.cpp | 2 +- .../device_binary_format_ar.cpp | 34 +++- .../device_binary_format_ar_tests.cpp | 152 ++++++++++++++- 9 files changed, 450 insertions(+), 43 deletions(-) diff --git a/opencl/test/unit_test/offline_compiler/mock/mock_argument_helper.h b/opencl/test/unit_test/offline_compiler/mock/mock_argument_helper.h index 9157d2d6b2..b553484cea 100644 --- a/opencl/test/unit_test/offline_compiler/mock/mock_argument_helper.h +++ b/opencl/test/unit_test/offline_compiler/mock/mock_argument_helper.h @@ -24,7 +24,7 @@ class MockOclocArgHelper : public OclocArgHelper { using OclocArgHelper::deviceProductTable; FilesMap &filesMap; bool interceptOutput{false}; - bool shouldReturnReadingError{false}; + bool shouldLoadDataFromFileReturnZeroSize{false}; FilesMap interceptedFiles; std::vector createdFiles{}; bool callBaseFileExists = false; @@ -67,8 +67,9 @@ class MockOclocArgHelper : public OclocArgHelper { return OclocArgHelper::loadDataFromFile(filename, retSize); } - if (shouldReturnReadingError) { - return nullptr; + if (shouldLoadDataFromFileReturnZeroSize) { + retSize = 0; + return {}; } if (!fileExists(filename)) { @@ -76,10 +77,12 @@ class MockOclocArgHelper : public OclocArgHelper { } const auto &file = filesMap[filename]; - std::unique_ptr result{new char[file.size()]}; + std::unique_ptr result{new char[file.size() + 1]}; std::copy(file.begin(), file.end(), result.get()); - retSize = file.size(); + result[file.size()] = '\0'; + + retSize = file.size() + 1; return result; } diff --git a/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.cpp b/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.cpp index b426941c8f..0ffa4a66f8 100644 --- a/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.cpp +++ b/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.cpp @@ -9,6 +9,10 @@ #include "shared/offline_compiler/source/ocloc_arg_helper.h" #include "shared/offline_compiler/source/ocloc_error_code.h" +#include "shared/source/device_binary_format/ar/ar.h" +#include "shared/source/device_binary_format/ar/ar_decoder.h" +#include "shared/source/device_binary_format/elf/elf_decoder.h" +#include "shared/source/device_binary_format/elf/ocl_elf.h" #include "shared/source/helpers/hw_helper.h" #include @@ -16,6 +20,27 @@ namespace NEO { +auto searchInArchiveByFilename(const Ar::Ar &archive, const ConstStringRef &name) { + const auto isSearchedFile = [&name](const auto &file) { + return file.fileName == name; + }; + + const auto &arFiles = archive.files; + return std::find_if(arFiles.begin(), arFiles.end(), isSearchedFile); +} + +std::string prepareTwoDevices(MockOclocArgHelper *argHelper) { + auto allEnabledDeviceConfigs = argHelper->getAllSupportedDeviceConfigs(); + if (allEnabledDeviceConfigs.size() < 2) { + return {}; + } + + const auto cfg1 = argHelper->parseProductConfigFromValue(allEnabledDeviceConfigs[0].config); + const auto cfg2 = argHelper->parseProductConfigFromValue(allEnabledDeviceConfigs[1].config); + + return cfg1 + "," + cfg2; +} + TEST(OclocFatBinaryRequestedFatBinary, WhenDeviceArgMissingThenReturnsFalse) { const char *args[] = {"ocloc", "-aaa", "*", "-device", "*"}; @@ -1049,4 +1074,161 @@ TEST_F(OclocFatBinaryGetTargetConfigsForFatbinary, GivenArgsWhenCorrectDeviceNum EXPECT_FALSE(got.empty()); } -} // namespace NEO +TEST_F(OclocFatBinaryTest, GivenSpirvInputWhenFatBinaryIsRequestedThenArchiveContainsGenericIrFileWithSpirvContent) { + const auto devices = prepareTwoDevices(&mockArgHelper); + if (devices.empty()) { + GTEST_SKIP(); + } + + const std::vector args = { + "ocloc", + "-output", + outputArchiveName, + "-file", + spirvFilename, + "-output_no_suffix", + "-spirv_input", + "-device", + devices}; + + const auto buildResult = buildFatBinary(args, &mockArgHelper); + ASSERT_EQ(OclocErrorCode::SUCCESS, buildResult); + ASSERT_EQ(1u, mockArgHelper.interceptedFiles.count(outputArchiveName)); + + const auto &rawArchive = mockArgHelper.interceptedFiles[outputArchiveName]; + const auto archiveBytes = ArrayRef::fromAny(rawArchive.data(), rawArchive.size()); + + std::string outErrReason{}; + std::string outWarning{}; + const auto decodedArchive = NEO::Ar::decodeAr(archiveBytes, outErrReason, outWarning); + + ASSERT_NE(nullptr, decodedArchive.magic); + ASSERT_TRUE(outErrReason.empty()); + ASSERT_TRUE(outWarning.empty()); + + const auto spirvFileIt = searchInArchiveByFilename(decodedArchive, archiveGenericIrName); + ASSERT_NE(decodedArchive.files.end(), spirvFileIt); + + const auto elf = Elf::decodeElf(spirvFileIt->fileData, outErrReason, outWarning); + ASSERT_NE(nullptr, elf.elfFileHeader); + ASSERT_TRUE(outErrReason.empty()); + ASSERT_TRUE(outWarning.empty()); + + const auto isSpirvSection = [](const auto §ion) { + return section.header && section.header->type == Elf::SHT_OPENCL_SPIRV; + }; + + const auto spirvSectionIt = std::find_if(elf.sectionHeaders.begin(), elf.sectionHeaders.end(), isSpirvSection); + ASSERT_NE(elf.sectionHeaders.end(), spirvSectionIt); + + ASSERT_EQ(spirvFileContent.size() + 1, spirvSectionIt->header->size); + const auto isSpirvDataEqualsInputFileData = std::memcmp(spirvFileContent.data(), spirvSectionIt->data.begin(), spirvFileContent.size()) == 0; + EXPECT_TRUE(isSpirvDataEqualsInputFileData); +} + +TEST_F(OclocFatBinaryTest, GivenSpirvInputAndExcludeIrFlagWhenFatBinaryIsRequestedThenArchiveDoesNotContainGenericIrFile) { + const auto devices = prepareTwoDevices(&mockArgHelper); + if (devices.empty()) { + GTEST_SKIP(); + } + + const std::vector args = { + "ocloc", + "-output", + outputArchiveName, + "-file", + spirvFilename, + "-output_no_suffix", + "-spirv_input", + "-exclude_ir", + "-device", + devices}; + + const auto buildResult = buildFatBinary(args, &mockArgHelper); + ASSERT_EQ(OclocErrorCode::SUCCESS, buildResult); + ASSERT_EQ(1u, mockArgHelper.interceptedFiles.count(outputArchiveName)); + + const auto &rawArchive = mockArgHelper.interceptedFiles[outputArchiveName]; + const auto archiveBytes = ArrayRef::fromAny(rawArchive.data(), rawArchive.size()); + + std::string outErrReason{}; + std::string outWarning{}; + const auto decodedArchive = NEO::Ar::decodeAr(archiveBytes, outErrReason, outWarning); + + ASSERT_NE(nullptr, decodedArchive.magic); + ASSERT_TRUE(outErrReason.empty()); + ASSERT_TRUE(outWarning.empty()); + + const auto spirvFileIt = searchInArchiveByFilename(decodedArchive, archiveGenericIrName); + EXPECT_EQ(decodedArchive.files.end(), spirvFileIt); +} + +TEST_F(OclocFatBinaryTest, GivenClInputFileWhenFatBinaryIsRequestedThenArchiveDoesNotContainGenericIrFile) { + const auto devices = prepareTwoDevices(&mockArgHelper); + if (devices.empty()) { + GTEST_SKIP(); + } + + const std::string clFilename = "some_kernel.cl"; + mockArgHelperFilesMap[clFilename] = "__kernel void some_kernel(){}"; + + const std::vector args = { + "ocloc", + "-output", + outputArchiveName, + "-file", + clFilename, + "-output_no_suffix", + "-device", + devices}; + + const auto buildResult = buildFatBinary(args, &mockArgHelper); + ASSERT_EQ(OclocErrorCode::SUCCESS, buildResult); + ASSERT_EQ(1u, mockArgHelper.interceptedFiles.count(outputArchiveName)); + + const auto &rawArchive = mockArgHelper.interceptedFiles[outputArchiveName]; + const auto archiveBytes = ArrayRef::fromAny(rawArchive.data(), rawArchive.size()); + + std::string outErrReason{}; + std::string outWarning{}; + const auto decodedArchive = NEO::Ar::decodeAr(archiveBytes, outErrReason, outWarning); + + ASSERT_NE(nullptr, decodedArchive.magic); + ASSERT_TRUE(outErrReason.empty()); + ASSERT_TRUE(outWarning.empty()); + + const auto spirvFileIt = searchInArchiveByFilename(decodedArchive, archiveGenericIrName); + EXPECT_EQ(decodedArchive.files.end(), spirvFileIt); +} + +TEST_F(OclocFatBinaryTest, GivenEmptyFileWhenAppendingGenericIrThenInvalidFileIsReturned) { + Ar::ArEncoder ar; + std::string emptyFile{"empty_file.spv"}; + mockArgHelperFilesMap[emptyFile] = ""; + mockArgHelper.shouldLoadDataFromFileReturnZeroSize = true; + + ::testing::internal::CaptureStdout(); + const auto errorCode{appendGenericIr(ar, emptyFile, &mockArgHelper)}; + const auto output{::testing::internal::GetCapturedStdout()}; + + EXPECT_EQ(OclocErrorCode::INVALID_FILE, errorCode); + EXPECT_EQ("Error! Couldn't read input file!\n", output); +} + +TEST_F(OclocFatBinaryTest, GivenInvalidIrFileWhenAppendingGenericIrThenInvalidFileIsReturned) { + Ar::ArEncoder ar; + std::string dummyFile{"dummy_file.spv"}; + mockArgHelperFilesMap[dummyFile] = "This is not IR!"; + + ::testing::internal::CaptureStdout(); + const auto errorCode{appendGenericIr(ar, dummyFile, &mockArgHelper)}; + const auto output{::testing::internal::GetCapturedStdout()}; + + EXPECT_EQ(OclocErrorCode::INVALID_FILE, errorCode); + + const auto expectedErrorMessage{"Error! Input file is not in supported generic IR format! " + "Currently supported format is SPIR-V.\n"}; + EXPECT_EQ(expectedErrorMessage, output); +} + +} // namespace NEO \ No newline at end of file diff --git a/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.h b/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.h index 604dcc5b2a..ec043e206e 100644 --- a/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.h +++ b/opencl/test/unit_test/offline_compiler/ocloc_fatbinary_tests.h @@ -1,5 +1,5 @@ /* - * Copyright (C) 2020-2021 Intel Corporation + * Copyright (C) 2020-2022 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -11,10 +11,12 @@ #include "shared/offline_compiler/source/ocloc_fatbinary.h" #include "gtest/gtest.h" +#include "mock/mock_argument_helper.h" #include namespace NEO { + class OclocFatBinaryGetTargetConfigsForFatbinary : public ::testing::Test { public: OclocFatBinaryGetTargetConfigsForFatbinary() { @@ -23,4 +25,23 @@ class OclocFatBinaryGetTargetConfigsForFatbinary : public ::testing::Test { } std::unique_ptr oclocArgHelperWithoutInput; }; + +class OclocFatBinaryTest : public ::testing::Test { + public: + OclocFatBinaryTest() { + mockArgHelperFilesMap[spirvFilename] = spirvFileContent; + mockArgHelper.interceptOutput = true; + } + + protected: + constexpr static ConstStringRef archiveGenericIrName{"generic_ir"}; + + MockOclocArgHelper::FilesMap mockArgHelperFilesMap{}; + MockOclocArgHelper mockArgHelper{mockArgHelperFilesMap}; + + std::string outputArchiveName{"output_archive"}; + std::string spirvFilename{"input_file.spv"}; + std::string spirvFileContent{"\x07\x23\x02\x03"}; +}; + } // namespace NEO diff --git a/opencl/test/unit_test/offline_compiler/offline_linker_tests.cpp b/opencl/test/unit_test/offline_compiler/offline_linker_tests.cpp index c7ee6ac309..e8a9c43cc5 100644 --- a/opencl/test/unit_test/offline_compiler/offline_linker_tests.cpp +++ b/opencl/test/unit_test/offline_compiler/offline_linker_tests.cpp @@ -339,9 +339,8 @@ TEST_F(OfflineLinkerTest, GivenValidCommandWhenVerificationIsPerformedThenSucces TEST_F(OfflineLinkerTest, GivenEmptyFileWhenLoadingInputFilesThenErrorIsReturned) { const std::string filename{"some_file.spv"}; - - // Empty file is treated as an error. mockArgHelperFilesMap[filename] = ""; + mockArgHelper.shouldLoadDataFromFileReturnZeroSize = true; const std::vector argv = { "ocloc.exe", @@ -380,25 +379,6 @@ TEST_F(OfflineLinkerTest, GivenValidFileWithUnknownFormatWhenLoadingInputFilesTh EXPECT_EQ(expectedErrorMessage, output); } -TEST_F(OfflineLinkerTest, GivenReadingErrorWhenLoadingInputFilesThenErrorIsReturned) { - const std::string filename{"some_file1.spv"}; - mockArgHelperFilesMap[filename] = getEmptySpirvFile(); - - mockArgHelper.shouldReturnReadingError = true; - - MockOfflineLinker mockOfflineLinker{&mockArgHelper}; - mockOfflineLinker.inputFilenames.push_back(filename); - - ::testing::internal::CaptureStdout(); - const auto readingResult = mockOfflineLinker.loadInputFilesContent(); - const auto output{::testing::internal::GetCapturedStdout()}; - - ASSERT_EQ(OclocErrorCode::INVALID_FILE, readingResult); - - const std::string expectedErrorMessage{"Error: Cannot read input file: some_file1.spv\n"}; - EXPECT_EQ(expectedErrorMessage, output); -} - TEST_F(OfflineLinkerTest, GivenValidFilesWithValidFormatsWhenLoadingInputFilesThenFilesAreLoadedAndSuccessIsReturned) { const std::string firstFilename{"some_file1.spv"}; const std::string secondFilename{"some_file2.llvmbc"}; @@ -420,14 +400,14 @@ TEST_F(OfflineLinkerTest, GivenValidFilesWithValidFormatsWhenLoadingInputFilesTh const auto &firstExpectedContent = mockArgHelperFilesMap[firstFilename]; const auto &firstActualContent = mockOfflineLinker.inputFilesContent[0]; - ASSERT_EQ(firstExpectedContent.size(), firstActualContent.size); + ASSERT_EQ(firstExpectedContent.size() + 1, firstActualContent.size); const auto isFirstPairEqual = std::equal(firstExpectedContent.begin(), firstExpectedContent.end(), firstActualContent.bytes.get()); EXPECT_TRUE(isFirstPairEqual); const auto &secondExpectedContent = mockArgHelperFilesMap[secondFilename]; const auto &secondActualContent = mockOfflineLinker.inputFilesContent[1]; - ASSERT_EQ(secondExpectedContent.size(), secondActualContent.size); + ASSERT_EQ(secondExpectedContent.size() + 1, secondActualContent.size); const auto isSecondPairEqual = std::equal(secondExpectedContent.begin(), secondExpectedContent.end(), secondActualContent.bytes.get()); EXPECT_TRUE(isSecondPairEqual); } diff --git a/shared/offline_compiler/source/ocloc_fatbinary.cpp b/shared/offline_compiler/source/ocloc_fatbinary.cpp index 7e39d859f4..1c4dad664a 100644 --- a/shared/offline_compiler/source/ocloc_fatbinary.cpp +++ b/shared/offline_compiler/source/ocloc_fatbinary.cpp @@ -9,6 +9,9 @@ #include "shared/offline_compiler/source/ocloc_error_code.h" #include "shared/offline_compiler/source/utilities/safety_caller.h" +#include "shared/source/compiler_interface/intermediate_representations.h" +#include "shared/source/device_binary_format/elf/elf_encoder.h" +#include "shared/source/device_binary_format/elf/ocl_elf.h" #include "shared/source/helpers/file_io.h" #include "shared/source/helpers/hw_info.h" @@ -341,6 +344,8 @@ int buildFatBinary(const std::vector &args, OclocArgHelper *argHelp std::string inputFileName = ""; std::string outputFileName = ""; std::string outputDirectory = ""; + bool spirvInput = false; + bool excludeIr = false; std::vector argsCopy(args); for (size_t argIndex = 1; argIndex < args.size(); argIndex++) { @@ -362,9 +367,18 @@ int buildFatBinary(const std::vector &args, OclocArgHelper *argHelp } else if ((ConstStringRef("-out_dir") == currArg) && hasMoreArgs) { outputDirectory = args[argIndex + 1]; ++argIndex; + } else if (ConstStringRef("-exclude_ir") == currArg) { + excludeIr = true; + } else if (ConstStringRef("-spirv_input") == currArg) { + spirvInput = true; } } + const bool shouldPreserveGenericIr = spirvInput && !excludeIr; + if (shouldPreserveGenericIr) { + argsCopy.push_back("-exclude_ir"); + } + Ar::ArEncoder fatbinary(true); if (isDeviceWithPlatformAbbreviation(ConstStringRef(args[deviceArgIndex]), argHelper)) { @@ -420,6 +434,14 @@ int buildFatBinary(const std::vector &args, OclocArgHelper *argHelp } } + if (shouldPreserveGenericIr) { + const auto errorCode = appendGenericIr(fatbinary, inputFileName, argHelper); + if (errorCode != OclocErrorCode::SUCCESS) { + argHelper->printf("Error! Couldn't append generic IR file!\n"); + return errorCode; + } + } + auto fatbinaryData = fatbinary.encode(); std::string fatbinaryFileName = outputFileName; if (outputFileName.empty() && (false == inputFileName.empty())) { @@ -433,4 +455,35 @@ int buildFatBinary(const std::vector &args, OclocArgHelper *argHelp return 0; } -} // namespace NEO +int appendGenericIr(Ar::ArEncoder &fatbinary, const std::string &inputFile, OclocArgHelper *argHelper) { + std::size_t fileSize = 0; + std::unique_ptr fileContents = argHelper->loadDataFromFile(inputFile, fileSize); + if (fileSize == 0) { + argHelper->printf("Error! Couldn't read input file!\n"); + return OclocErrorCode::INVALID_FILE; + } + + const auto ir = ArrayRef::fromAny(fileContents.get(), fileSize); + if (!isSpirVBitcode(ir)) { + argHelper->printf("Error! Input file is not in supported generic IR format! " + "Currently supported format is SPIR-V.\n"); + return OclocErrorCode::INVALID_FILE; + } + + const auto encodedElf = createEncodedElfWithSpirv(ir); + ArrayRef genericIrFile{encodedElf.data(), encodedElf.size()}; + + fatbinary.appendFileEntry("generic_ir", genericIrFile); + return OclocErrorCode::SUCCESS; +} + +std::vector createEncodedElfWithSpirv(const ArrayRef &spirv) { + using namespace NEO::Elf; + ElfEncoder elfEncoder; + elfEncoder.getElfFileHeader().type = ET_OPENCL_OBJECTS; + elfEncoder.appendSection(SHT_OPENCL_SPIRV, SectionNamesOpenCl::spirvObject, spirv); + + return elfEncoder.encode(); +} + +} // namespace NEO \ No newline at end of file diff --git a/shared/offline_compiler/source/ocloc_fatbinary.h b/shared/offline_compiler/source/ocloc_fatbinary.h index d5accca016..5b64d802e5 100644 --- a/shared/offline_compiler/source/ocloc_fatbinary.h +++ b/shared/offline_compiler/source/ocloc_fatbinary.h @@ -14,6 +14,7 @@ #include "compiler_options.h" #include "igfxfmid.h" +#include #include #include @@ -48,4 +49,7 @@ std::vector toProductNames(const std::vector &pr PRODUCT_FAMILY asProductId(ConstStringRef product, const std::vector &allSupportedPlatforms); int buildFatBinaryForTarget(int retVal, std::vector argsCopy, std::string pointerSize, Ar::ArEncoder &fatbinary, OfflineCompiler *pCompiler, OclocArgHelper *argHelper, const std::string &deviceConfig); +int appendGenericIr(Ar::ArEncoder &fatbinary, const std::string &inputFile, OclocArgHelper *argHelper); +std::vector createEncodedElfWithSpirv(const ArrayRef &spirv); + } // namespace NEO diff --git a/shared/offline_compiler/source/offline_linker.cpp b/shared/offline_compiler/source/offline_linker.cpp index 4a9595d82f..acc76d5797 100644 --- a/shared/offline_compiler/source/offline_linker.cpp +++ b/shared/offline_compiler/source/offline_linker.cpp @@ -163,7 +163,7 @@ int OfflineLinker::loadInputFilesContent() { for (const auto &filename : inputFilenames) { size = 0; bytes = argHelper->loadDataFromFile(filename, size); - if (bytes == nullptr || size == 0) { + if (size == 0) { argHelper->printf("Error: Cannot read input file: %s\n", filename.c_str()); return OclocErrorCode::INVALID_FILE; } diff --git a/shared/source/device_binary_format/device_binary_format_ar.cpp b/shared/source/device_binary_format/device_binary_format_ar.cpp index daea0a4f5e..212e1e0b13 100644 --- a/shared/source/device_binary_format/device_binary_format_ar.cpp +++ b/shared/source/device_binary_format/device_binary_format_ar.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2020-2021 Intel Corporation + * Copyright (C) 2020-2022 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -27,20 +27,26 @@ SingleDeviceBinary unpackSingleDeviceBinary(co std::string pointerSize = ((requestedTargetDevice.maxPointerSizeInBytes == 8) ? "64" : "32"); std::string filterPointerSizeAndPlatform = pointerSize + "." + requestedProductAbbreviation.str(); std::string filterPointerSizeAndPlatformAndStepping = filterPointerSizeAndPlatform + "." + std::to_string(requestedTargetDevice.stepping); + ConstStringRef genericIrFileName{"generic_ir"}; - Ar::ArFileEntryHeaderAndData *matchedFiles[2] = {}; + Ar::ArFileEntryHeaderAndData *matchedFiles[3] = {}; Ar::ArFileEntryHeaderAndData *&matchedPointerSizeAndPlatformAndStepping = matchedFiles[0]; // best match Ar::ArFileEntryHeaderAndData *&matchedPointerSizeAndPlatform = matchedFiles[1]; - for (auto &f : archiveData.files) { - if (ConstStringRef(f.fileName.begin(), filterPointerSizeAndPlatform.size()) != filterPointerSizeAndPlatform) { - continue; - } + Ar::ArFileEntryHeaderAndData *&matchedGenericIr = matchedFiles[2]; - if (ConstStringRef(f.fileName.begin(), filterPointerSizeAndPlatformAndStepping.size()) != filterPointerSizeAndPlatformAndStepping) { - matchedPointerSizeAndPlatform = &f; - continue; + for (auto &file : archiveData.files) { + const auto &filename = file.fileName; + constexpr std::string::size_type zeroIndex{0}; + + if (filename.size() >= filterPointerSizeAndPlatformAndStepping.size() && + filename.substr(zeroIndex, filterPointerSizeAndPlatformAndStepping.size()) == filterPointerSizeAndPlatformAndStepping) { + matchedPointerSizeAndPlatformAndStepping = &file; + } else if (filename.size() >= filterPointerSizeAndPlatform.size() && + filename.substr(zeroIndex, filterPointerSizeAndPlatform.size()) == filterPointerSizeAndPlatform) { + matchedPointerSizeAndPlatform = &file; + } else if (file.fileName == genericIrFileName) { + matchedGenericIr = &file; } - matchedPointerSizeAndPlatformAndStepping = &f; } std::string unpackErrors; @@ -55,6 +61,14 @@ SingleDeviceBinary unpackSingleDeviceBinary(co if (matchedFile != matchedPointerSizeAndPlatformAndStepping) { outWarning = "Couldn't find perfectly matched binary (right stepping) in AR, using best usable"; } + + if (unpacked.intermediateRepresentation.empty() && matchedGenericIr) { + auto unpackedGenericIr = unpackSingleDeviceBinary(matchedGenericIr->fileData, requestedProductAbbreviation, requestedTargetDevice, unpackErrors, unpackWarnings); + if (!unpackedGenericIr.intermediateRepresentation.empty()) { + unpacked.intermediateRepresentation = unpackedGenericIr.intermediateRepresentation; + } + } + return unpacked; } if (binaryForRecompilation.intermediateRepresentation.empty() && (false == unpacked.intermediateRepresentation.empty())) { diff --git a/shared/test/unit_test/device_binary_format/device_binary_format_ar_tests.cpp b/shared/test/unit_test/device_binary_format/device_binary_format_ar_tests.cpp index c756c08e18..7ed6721521 100644 --- a/shared/test/unit_test/device_binary_format/device_binary_format_ar_tests.cpp +++ b/shared/test/unit_test/device_binary_format/device_binary_format_ar_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2020-2021 Intel Corporation + * Copyright (C) 2020-2022 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -183,6 +183,156 @@ TEST(UnpackSingleDeviceBinaryAr, WhenDeviceBinaryNotMatchedButIrAvailableThenUse EXPECT_FALSE(unpacked.intermediateRepresentation.empty()); } +TEST(UnpackSingleDeviceBinaryAr, WhenDeviceBinaryNotMatchedButGenericIrFileAvailableThenUseGenericIr) { + PatchTokensTestData::ValidEmptyProgram programTokens; + std::string requiredProduct = NEO::hardwarePrefix[productFamily]; + std::string requiredStepping = std::to_string(programTokens.header->SteppingId); + std::string requiredPointerSize = (programTokens.header->GPUPointerSizeInBytes == 4) ? "32" : "64"; + + NEO::Elf::ElfEncoder elfEncoder; + elfEncoder.getElfFileHeader().type = NEO::Elf::ET_OPENCL_OBJECTS; + + const auto spirvFile{ArrayRef::fromAny(NEO::spirvMagic.begin(), NEO::spirvMagic.size())}; + elfEncoder.appendSection(NEO::Elf::SHT_OPENCL_SPIRV, NEO::Elf::SectionNamesOpenCl::spirvObject, spirvFile); + + const auto elfData = elfEncoder.encode(); + NEO::Ar::ArEncoder encoder{true}; + ASSERT_TRUE(encoder.appendFileEntry("generic_ir", ArrayRef(elfData))); + + NEO::TargetDevice target; + target.coreFamily = static_cast(programTokens.header->Device); + target.stepping = programTokens.header->SteppingId; + target.maxPointerSizeInBytes = programTokens.header->GPUPointerSizeInBytes; + + auto arData = encoder.encode(); + std::string unpackErrors; + std::string unpackWarnings; + auto unpacked = NEO::unpackSingleDeviceBinary(arData, requiredProduct, target, unpackErrors, unpackWarnings); + EXPECT_TRUE(unpackErrors.empty()) << unpackErrors; + EXPECT_TRUE(unpackWarnings.empty()) << unpackWarnings; + + EXPECT_FALSE(unpacked.intermediateRepresentation.empty()); +} + +TEST(UnpackSingleDeviceBinaryAr, GivenInvalidGenericIrFileWhenDeviceBinaryNotMatchedButGenericIrFileAvailableThenIrIsEmpty) { + PatchTokensTestData::ValidEmptyProgram programTokens; + std::string requiredProduct = NEO::hardwarePrefix[productFamily]; + std::string requiredStepping = std::to_string(programTokens.header->SteppingId); + std::string requiredPointerSize = (programTokens.header->GPUPointerSizeInBytes == 4) ? "32" : "64"; + + NEO::Ar::ArEncoder encoder{true}; + ASSERT_TRUE(encoder.appendFileEntry(requiredPointerSize + "." + requiredProduct + "." + requiredStepping, programTokens.storage)); + + NEO::Elf::ElfEncoder elfEncoder; + elfEncoder.getElfFileHeader().type = NEO::Elf::ET_OPENCL_OBJECTS; + + const auto elfData = elfEncoder.encode(); + ASSERT_TRUE(encoder.appendFileEntry("generic_ir", ArrayRef(elfData))); + + NEO::TargetDevice target; + target.coreFamily = static_cast(programTokens.header->Device); + target.stepping = programTokens.header->SteppingId; + target.maxPointerSizeInBytes = programTokens.header->GPUPointerSizeInBytes; + + auto arData = encoder.encode(); + std::string unpackErrors; + std::string unpackWarnings; + auto unpacked = NEO::unpackSingleDeviceBinary(arData, requiredProduct, target, unpackErrors, unpackWarnings); + EXPECT_TRUE(unpackErrors.empty()) << unpackErrors; + EXPECT_TRUE(unpackWarnings.empty()) << unpackWarnings; + + EXPECT_TRUE(unpacked.intermediateRepresentation.empty()); +} + +TEST(UnpackSingleDeviceBinaryAr, WhenDeviceBinaryMatchedButHasNoIrAndGenericIrFileAvailableThenUseBinaryWithAssignedGenericIr) { + PatchTokensTestData::ValidEmptyProgram programTokens; + std::string requiredProduct = NEO::hardwarePrefix[productFamily]; + std::string requiredStepping = std::to_string(programTokens.header->SteppingId); + std::string requiredPointerSize = (programTokens.header->GPUPointerSizeInBytes == 4) ? "32" : "64"; + + NEO::Ar::ArEncoder encoder{true}; + ASSERT_TRUE(encoder.appendFileEntry(requiredPointerSize + "." + requiredProduct + "." + requiredStepping, programTokens.storage)); + + NEO::Elf::ElfEncoder elfEncoderIr; + elfEncoderIr.getElfFileHeader().type = NEO::Elf::ET_OPENCL_OBJECTS; + + const std::string customSprivContent{"\x07\x23\x02\x03This is a custom file, with SPIR-V magic!"}; + const auto spirvFile{ArrayRef::fromAny(customSprivContent.data(), customSprivContent.size())}; + elfEncoderIr.appendSection(NEO::Elf::SHT_OPENCL_SPIRV, NEO::Elf::SectionNamesOpenCl::spirvObject, spirvFile); + + const auto elfIrData = elfEncoderIr.encode(); + ASSERT_TRUE(encoder.appendFileEntry("generic_ir", ArrayRef(elfIrData))); + + NEO::TargetDevice target; + target.coreFamily = static_cast(programTokens.header->Device); + target.stepping = programTokens.header->SteppingId; + target.maxPointerSizeInBytes = programTokens.header->GPUPointerSizeInBytes; + + auto arData = encoder.encode(); + std::string unpackErrors; + std::string unpackWarnings; + auto unpacked = NEO::unpackSingleDeviceBinary(arData, requiredProduct, target, unpackErrors, unpackWarnings); + EXPECT_TRUE(unpackErrors.empty()) << unpackErrors; + EXPECT_TRUE(unpackWarnings.empty()) << unpackWarnings; + + ASSERT_FALSE(unpacked.intermediateRepresentation.empty()); + ASSERT_EQ(customSprivContent.size(), unpacked.intermediateRepresentation.size()); + + const auto isSpirvSameAsInGenericIr = std::memcmp(customSprivContent.data(), unpacked.intermediateRepresentation.begin(), customSprivContent.size()) == 0; + EXPECT_TRUE(isSpirvSameAsInGenericIr); + + EXPECT_FALSE(unpacked.deviceBinary.empty()); +} + +TEST(UnpackSingleDeviceBinaryAr, WhenDeviceBinaryMatchedAndHasIrAndGenericIrFileAvailableThenUseBinaryAndItsIr) { + PatchTokensTestData::ValidEmptyProgram programTokens; + std::string requiredProduct = NEO::hardwarePrefix[productFamily]; + std::string requiredStepping = std::to_string(programTokens.header->SteppingId); + std::string requiredPointerSize = (programTokens.header->GPUPointerSizeInBytes == 4) ? "32" : "64"; + + NEO::Elf::ElfEncoder elfEncoderBinary; + elfEncoderBinary.getElfFileHeader().type = NEO::Elf::ET_OPENCL_EXECUTABLE; + + const std::string customSprivContentBinary{"\x07\x23\x02\x03This is a binary's IR!"}; + const auto spirvFileBinary{ArrayRef::fromAny(customSprivContentBinary.data(), customSprivContentBinary.size())}; + elfEncoderBinary.appendSection(NEO::Elf::SHT_OPENCL_SPIRV, NEO::Elf::SectionNamesOpenCl::spirvObject, spirvFileBinary); + elfEncoderBinary.appendSection(NEO::Elf::SHT_OPENCL_DEV_BINARY, NEO::Elf::SectionNamesOpenCl::deviceBinary, programTokens.storage); + + NEO::Ar::ArEncoder encoder{true}; + const auto elfBinaryData = elfEncoderBinary.encode(); + ASSERT_TRUE(encoder.appendFileEntry(requiredPointerSize + "." + requiredProduct + "." + requiredStepping, ArrayRef(elfBinaryData))); + + NEO::Elf::ElfEncoder elfEncoderIr; + elfEncoderIr.getElfFileHeader().type = NEO::Elf::ET_OPENCL_OBJECTS; + + const std::string customSprivContentGenericIr{"\x07\x23\x02\x03This is a generic ir!"}; + const auto spirvFile{ArrayRef::fromAny(customSprivContentGenericIr.data(), customSprivContentGenericIr.size())}; + elfEncoderIr.appendSection(NEO::Elf::SHT_OPENCL_SPIRV, NEO::Elf::SectionNamesOpenCl::spirvObject, spirvFile); + + const auto elfIrData = elfEncoderIr.encode(); + ASSERT_TRUE(encoder.appendFileEntry("generic_ir", ArrayRef(elfIrData))); + + NEO::TargetDevice target; + target.coreFamily = static_cast(programTokens.header->Device); + target.stepping = programTokens.header->SteppingId; + target.maxPointerSizeInBytes = programTokens.header->GPUPointerSizeInBytes; + + auto arData = encoder.encode(); + std::string unpackErrors; + std::string unpackWarnings; + auto unpacked = NEO::unpackSingleDeviceBinary(arData, requiredProduct, target, unpackErrors, unpackWarnings); + EXPECT_TRUE(unpackErrors.empty()) << unpackErrors; + EXPECT_TRUE(unpackWarnings.empty()) << unpackWarnings; + + ASSERT_FALSE(unpacked.intermediateRepresentation.empty()); + ASSERT_EQ(customSprivContentBinary.size(), unpacked.intermediateRepresentation.size()); + + const auto isSpirvSameAsInBinary = std::memcmp(customSprivContentBinary.data(), unpacked.intermediateRepresentation.begin(), customSprivContentBinary.size()) == 0; + EXPECT_TRUE(isSpirvSameAsInBinary); + + EXPECT_FALSE(unpacked.deviceBinary.empty()); +} + TEST(UnpackSingleDeviceBinaryAr, WhenOnlyIrIsAvailableThenUseOneFromBestMatchedBinary) { PatchTokensTestData::ValidEmptyProgram programTokens; std::string requiredProduct = NEO::hardwarePrefix[productFamily];