diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index fca8164e90970..d2c1162eec27f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -547,8 +547,24 @@ static const char *getUrDeviceTarget(const char *URDeviceTarget) { return UR_DEVICE_BINARY_TARGET_UNKNOWN; } +// An image with a compile_target was built for that device architecture only, +// so it only matches a device of that architecture. Images without a +// compile_target are not filtered here. +static bool doesCompileTargetMatchDevice(const RTDeviceBinaryImage &Img, + const device_impl &Dev) { + const RTDeviceBinaryImage::PropertyRange &Range = Img.getDeviceRequirements(); + const bool HasCompileTarget = + std::any_of(Range.begin(), Range.end(), [](const auto &Prop) { + return Prop->Name == std::string_view("compile_target"); + }); + return !HasCompileTarget || doesImageTargetMatchDevice(Img, Dev); +} + static bool compatibleWithDevice(const RTDeviceBinaryImage *BinImage, const device_impl &DeviceImpl) { + if (!doesCompileTargetMatchDevice(*BinImage, DeviceImpl)) + return false; + adapter_impl &Adapter = DeviceImpl.getAdapter(); const ur_device_handle_t &URDeviceHandle = DeviceImpl.getHandleRef(); @@ -791,7 +807,8 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions( // virtual function set and it is incompatible, then we should link its // "dummy" version to avoid link errors about unresolved external // symbols. - if (doesDevSupportDeviceRequirements(Dev, *BinImage)) + if (doesDevSupportDeviceRequirements(Dev, *BinImage) && + doesCompileTargetMatchDevice(*BinImage, Dev)) DeviceImagesToLink.insert(BinImage); } } @@ -1442,11 +1459,20 @@ const RTDeviceBinaryImage &ProgramManager::getDeviceImage( } std::lock_guard ImgMapsGuard(m_ImgMapsMutex); - std::vector RawImgs(ImageSet.size()); - auto ImageIterator = ImageSet.begin(); - for (size_t i = 0; i < ImageSet.size(); i++, ImageIterator++) + // Skip images built for another device architecture. + std::vector Imgs; + for (const RTDeviceBinaryImage *Img : ImageSet) + if (doesCompileTargetMatchDevice(*Img, DeviceImpl)) + Imgs.push_back(Img); + if (Imgs.empty()) + throw exception(make_error_code(errc::runtime), + "No device image is built for the selected device's " + "architecture"); + + std::vector RawImgs(Imgs.size()); + for (size_t i = 0; i < Imgs.size(); i++) RawImgs[i] = reinterpret_cast( - const_cast(&(*ImageIterator)->getRawData())); + const_cast(&Imgs[i]->getRawData())); uint32_t ImgInd = 0; // Ask the native runtime under the given context to choose the device image // it prefers. @@ -1460,15 +1486,14 @@ const RTDeviceBinaryImage &ProgramManager::getDeviceImage( ContextImpl.getAdapter().call( DeviceImpl.getHandleRef(), UrBinaries.data(), UrBinaries.size(), &ImgInd); - ImageIterator = ImageSet.begin(); - std::advance(ImageIterator, ImgInd); + const RTDeviceBinaryImage *SelectedImg = Imgs[ImgInd]; if constexpr (DbgProgMgr > 0) { - std::cerr << "selected device image: " << &(*ImageIterator)->getRawData() + std::cerr << "selected device image: " << &SelectedImg->getRawData() << "\n"; - (*ImageIterator)->print(); + SelectedImg->print(); } - return **ImageIterator; + return *SelectedImg; } // Check if device image is compressed. diff --git a/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp b/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp index 43bfea488c44b..7732a2087af0d 100644 --- a/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp +++ b/sycl/unittests/SYCL2020/SYCLBINSerialize.cpp @@ -62,6 +62,16 @@ sycl::unittest::MockDeviceImage Imgs[] = { makeImageWithCompileTarget("SYCLBINSerializeKernel", "intel_gpu_pvc")}; sycl::unittest::MockDeviceImageArray ImgArray{Imgs}; +// Makes the mock device a PVC GPU, so that the intel_gpu_pvc image is used. +ur_result_t redefinedDeviceGetInfoPVC(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = 0x030f0000; + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} + // Strips the OffloadBinary wrapper that ext_oneapi_get_content emits and // returns a string_view over the contained raw SYCLBIN bytes. Returns // std::nullopt if the wrapper does not parse cleanly; callers should ASSERT @@ -115,6 +125,8 @@ stripOffloadWrapper(const std::vector &Bytes) { // confirm the property survived. TEST(SYCLBINSerialize, CompileTargetRoundTrip) { sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoPVC); const sycl::device Dev = sycl::platform().get_devices()[0]; sycl::context Ctx{Dev}; @@ -222,6 +234,8 @@ TEST(SYCLBINSerialize, BundleStateRoundTripObject) { // inspection alone (e.g. CompileTargetRoundTrip) does not catch. TEST(SYCLBINSerialize, ImageBytesRoundTrip) { sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfoPVC); const sycl::device Dev = sycl::platform().get_devices()[0]; sycl::context Ctx{Dev}; diff --git a/sycl/unittests/program_manager/CMakeLists.txt b/sycl/unittests/program_manager/CMakeLists.txt index a88f3495768c0..d88c60bceea66 100644 --- a/sycl/unittests/program_manager/CMakeLists.txt +++ b/sycl/unittests/program_manager/CMakeLists.txt @@ -21,3 +21,4 @@ endif() add_subdirectory(arg_mask) add_subdirectory(DynamicLinking) +add_subdirectory(CompileTargetSelection) diff --git a/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt b/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt new file mode 100644 index 0000000000000..51cad31e87fe3 --- /dev/null +++ b/sycl/unittests/program_manager/CompileTargetSelection/CMakeLists.txt @@ -0,0 +1,5 @@ +# Uses its own test binary: the registered images decide which devices have a +# compatible image, so other tests' images must not be visible. +add_sycl_unittest(CompileTargetSelectionTests OBJECT + CompileTargetSelection.cpp +) diff --git a/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp b/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp new file mode 100644 index 0000000000000..0607ea026673e --- /dev/null +++ b/sycl/unittests/program_manager/CompileTargetSelection/CompileTargetSelection.cpp @@ -0,0 +1,103 @@ +#include + +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +// Check that an AOT image built for one GPU architecture only makes devices of +// that architecture preferred by the default device selector. + +class CompileTargetSelectionKernel; + +namespace sycl { +inline namespace _V1 { +namespace detail { +template <> +struct KernelInfo + : public unittest::MockKernelInfoBase { + static constexpr const char *getName() { + return "CompileTargetSelectionKernel"; + } +}; +} // namespace detail +} // namespace _V1 +} // namespace sycl + +namespace { +namespace exp = sycl::ext::oneapi::experimental; + +sycl::unittest::MockDeviceImage generatePVCImage() { + std::string CompileTarget = "intel_gpu_pvc"; + std::vector Data(8 + CompileTarget.size()); + std::copy(CompileTarget.begin(), CompileTarget.end(), Data.data() + 8); + sycl::unittest::MockPropertySet PropSet; + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS, + sycl::unittest::MockProperty("compile_target", Data, + SYCL_PROPERTY_TYPE_BYTE_ARRAY)); + return sycl::unittest::MockDeviceImage{ + SYCL_DEVICE_BINARY_TYPE_NATIVE, + __SYCL_DEVICE_BINARY_TARGET_SPIRV64_GEN, + "", // Compile options + "", // Link options + std::vector{1}, + sycl::unittest::makeEmptyKernels({"CompileTargetSelectionKernel"}), + std::move(PropSet)}; +} +sycl::unittest::MockDeviceImage Imgs[] = {generatePVCImage()}; +sycl::unittest::MockDeviceImageArray<1> ImgArray{Imgs}; + +// IP versions of a TGLLP and a PVC GPU. The TGLLP one comes first so that it +// would be chosen if both devices were preferred equally. +uint32_t MockDeviceIPs[] = {0x03000000, 0x030f0000}; + +ur_result_t redefinedDeviceGet(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppNumDevices) { + **Params.ppNumDevices = std::size(MockDeviceIPs); + return UR_RESULT_SUCCESS; + } + if (*Params.pphDevices) + for (uint32_t I = 0; I < *Params.pNumEntries; ++I) + (*Params.pphDevices)[I] = + reinterpret_cast(&MockDeviceIPs[I]); + return UR_RESULT_SUCCESS; +} + +ur_result_t redefinedDeviceGetInfo(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = + *reinterpret_cast(*Params.phDevice); + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} +} // namespace + +TEST(CompileTargetSelection, DefaultSelectorPrefersDeviceWithImage) { + sycl::unittest::UrMock Mock; + mock::getCallbacks().set_replace_callback("urDeviceGet", &redefinedDeviceGet); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + std::vector Devs = sycl::platform().get_devices(); + ASSERT_EQ(Devs.size(), 2u); + ASSERT_EQ(Devs[0].get_info(), + exp::architecture::intel_gpu_tgllp); + ASSERT_EQ(Devs[1].get_info(), + exp::architecture::intel_gpu_pvc); + EXPECT_LT(sycl::default_selector_v(Devs[0]), + sycl::default_selector_v(Devs[1])); + + sycl::device Dev{sycl::default_selector_v}; + EXPECT_EQ(Dev.get_info(), + exp::architecture::intel_gpu_pvc); +} diff --git a/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt b/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt index d57334faf6dd5..8f19b170eb37d 100644 --- a/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt +++ b/sycl/unittests/program_manager/DynamicLinking/CMakeLists.txt @@ -4,5 +4,6 @@ if (NOT LLVM_LIBCXX_USED) add_sycl_unittest(DynamicLinkingTests OBJECT DynamicLinking.cpp + DynamicLinkingCompileTarget.cpp ) endif() diff --git a/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp new file mode 100644 index 0000000000000..8e9a869050c2f --- /dev/null +++ b/sycl/unittests/program_manager/DynamicLinking/DynamicLinkingCompileTarget.cpp @@ -0,0 +1,211 @@ +#include + +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include +#include +#include +#include +#include + +// Check that the dependencies of an AOT image are resolved with images built +// for the same device when there are images for several devices. + +namespace DynamicLinkingCompileTargetTest { +class MultiArchKernel; +} // namespace DynamicLinkingCompileTargetTest + +namespace sycl { +inline namespace _V1 { +namespace detail { +template <> +struct KernelInfo + : public unittest::MockKernelInfoBase { + static constexpr const char *getName() { return "MultiArchKernel"; } +}; +} // namespace detail +} // namespace _V1 +} // namespace sycl + +namespace { +static constexpr unsigned TGLLP_PRG = 31; +static constexpr unsigned TGLLP_PRG_DEP = 37; +static constexpr unsigned PVC_PRG = 41; +static constexpr unsigned PVC_PRG_DEP = 43; + +// IP version of a PVC device. +static constexpr uint32_t PvcIp = 0x030f0000; + +sycl::unittest::MockDeviceImage +generateImage(std::initializer_list KernelNames, + const std::string &ExportedSymbol, + const std::string &ImportedSymbol, unsigned char Magic, + const std::string &CompileTarget) { + sycl::unittest::MockPropertySet PropSet; + + std::vector Data(8 + CompileTarget.size()); + std::copy(CompileTarget.begin(), CompileTarget.end(), Data.data() + 8); + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS, + sycl::unittest::MockProperty("compile_target", Data, + SYCL_PROPERTY_TYPE_BYTE_ARRAY)); + + auto makeSymbolProps = [](const std::string &Symbol) { + std::vector Props; + Props.emplace_back(Symbol, std::vector(sizeof(uint32_t)), + SYCL_PROPERTY_TYPE_UINT32); + return Props; + }; + if (!ExportedSymbol.empty()) + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS, + makeSymbolProps(ExportedSymbol)); + if (!ImportedSymbol.empty()) + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_IMPORTED_SYMBOLS, + makeSymbolProps(ImportedSymbol)); + + return sycl::unittest::MockDeviceImage{ + SYCL_DEVICE_BINARY_TYPE_NATIVE, + __SYCL_DEVICE_BINARY_TARGET_SPIRV64_GEN, + "", // Compile options + "", // Link options + std::vector{Magic}, + sycl::unittest::makeEmptyKernels(KernelNames), + std::move(PropSet)}; +} + +// The TGLLP images come first, so that picking the first image exporting the +// symbol would choose the wrong dependency for a PVC device. +static sycl::unittest::MockDeviceImage Imgs[] = { + generateImage({"MultiArchKernel"}, "", "MultiArchKernelDep", TGLLP_PRG, + "intel_gpu_tgllp"), + generateImage({"MultiArchKernelDep"}, "MultiArchKernelDep", "", + TGLLP_PRG_DEP, "intel_gpu_tgllp"), + generateImage({"MultiArchKernel"}, "", "MultiArchKernelDep", PVC_PRG, + "intel_gpu_pvc"), + generateImage({"MultiArchKernelDep"}, "MultiArchKernelDep", "", PVC_PRG_DEP, + "intel_gpu_pvc")}; +static sycl::unittest::MockDeviceImageArray<4> ImgArray{Imgs}; + +ur_result_t redefinedDeviceGetInfo(void *pParams) { + auto Params = *static_cast(pParams); + if (*Params.ppropName == UR_DEVICE_INFO_IP_VERSION && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = PvcIp; + if (*Params.ppropName == UR_DEVICE_INFO_TYPE && *Params.ppPropValue) + *static_cast(*Params.ppPropValue) = UR_DEVICE_TYPE_GPU; + return UR_RESULT_SUCCESS; +} +} // namespace + +TEST(DynamicLinking, AheadOfTimeDependencyForDevice) { + sycl::unittest::UrMock Mock; + setupRuntimeLinkingMock(); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::platform Plt = sycl::platform(); + sycl::queue Q(Plt.get_devices()[0]); + + CapturedLinkingData.clear(); + + Q.single_task([=]() {}); + ASSERT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PVC_PRG, PVC_PRG_DEP})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PVC_PRG * PVC_PRG_DEP); +} + +TEST(DynamicLinking, AheadOfTimeKernelBundleForDevice) { + sycl::unittest::UrMock Mock; + setupRuntimeLinkingMock(); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::platform Plt = sycl::platform(); + sycl::queue Q(Plt.get_devices()[0]); + + CapturedLinkingData.clear(); + + auto Bundle = sycl::get_kernel_bundle( + Q.get_context(), {Q.get_device()}, + {sycl::get_kernel_id< + DynamicLinkingCompileTargetTest::MultiArchKernel>()}); + Q.submit([&](sycl::handler &CGH) { + CGH.use_kernel_bundle(Bundle); + CGH.single_task([=]() {}); + }); + ASSERT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PVC_PRG, PVC_PRG_DEP})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PVC_PRG * PVC_PRG_DEP); +} + +// Images are also picked from a given set, e.g. for device_global accesses. +// The selection by the backend only sees the image targets, which are the same +// for all images here, so make it pick the TGLLP image whenever it is offered. +std::vector OfferedImgs; +ur_result_t redefinedDeviceSelectBinary(void *pParams) { + auto Params = *static_cast(pParams); + **Params.ppSelectedBinary = 0; + for (uint32_t I = 0; I < *Params.pNumBinaries; ++I) + if (OfferedImgs.size() == *Params.pNumBinaries && + *OfferedImgs[I]->getRawData().BinaryStart == TGLLP_PRG) + **Params.ppSelectedBinary = I; + return UR_RESULT_SUCCESS; +} + +TEST(DynamicLinking, AheadOfTimeImageFromSetForDevice) { + sycl::unittest::UrMock Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::queue Q(sycl::platform().get_devices()[0]); + auto &PM = sycl::detail::ProgramManager::getInstance(); + std::set KernelImgs = + PM.getRawDeviceImages({sycl::get_kernel_id< + DynamicLinkingCompileTargetTest::MultiArchKernel>()}); + ASSERT_EQ(KernelImgs.size(), 2u); + + std::unordered_set ImgSet( + KernelImgs.begin(), KernelImgs.end()); + // The images offered to the backend if none are filtered out. + OfferedImgs.assign(ImgSet.begin(), ImgSet.end()); + mock::getCallbacks().set_replace_callback("urDeviceSelectBinary", + &redefinedDeviceSelectBinary); + + const sycl::detail::RTDeviceBinaryImage &Img = + PM.getDeviceImage(ImgSet, *sycl::detail::getSyclObjImpl(Q.get_context()), + *sycl::detail::getSyclObjImpl(Q.get_device())); + EXPECT_EQ(*Img.getRawData().BinaryStart, PVC_PRG); +} + +TEST(DynamicLinking, AheadOfTimeImageFromSetForOtherDevice) { + sycl::unittest::UrMock Mock; + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::queue Q(sycl::platform().get_devices()[0]); + auto &PM = sycl::detail::ProgramManager::getInstance(); + std::unordered_set ImgSet; + for (const sycl::detail::RTDeviceBinaryImage *Img : + PM.getRawDeviceImages({sycl::get_kernel_id< + DynamicLinkingCompileTargetTest::MultiArchKernel>()})) + if (*Img->getRawData().BinaryStart != PVC_PRG) + ImgSet.insert(Img); + ASSERT_EQ(ImgSet.size(), 1u); + + // The only image is built for TGLLP, so it can't be used on PVC. + EXPECT_THROW(PM.getDeviceImage(ImgSet, + *sycl::detail::getSyclObjImpl(Q.get_context()), + *sycl::detail::getSyclObjImpl(Q.get_device())), + sycl::exception); +}