diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index 633a4269e1e78..0352e66c5aa0c 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -52,7 +52,7 @@ std::ostream &operator<<(std::ostream &Out, const DeviceBinaryProperty &P) { break; } case SYCL_PROPERTY_TYPE_STRING: - Out << P.asCString(); + Out << P.asStringView(); break; default: assert(false && "Unsupported property"); @@ -77,14 +77,14 @@ ByteArray DeviceBinaryProperty::asByteArray() const { return {Data, Prop->ValSize}; } -const char *DeviceBinaryProperty::asCString() const { +std::string_view DeviceBinaryProperty::asStringView() const { assert((Prop->Type == SYCL_PROPERTY_TYPE_STRING || Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY) && "property type mismatch"); assert(Prop->ValSize > 0 && "property size mismatch"); // Byte array stores its size in first 8 bytes size_t Shift = Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY ? 8 : 0; - return ur::cast(Prop->ValAddr) + Shift; + return {ur::cast(Prop->ValAddr) + Shift, Prop->ValSize - Shift}; } void RTDeviceBinaryImage::PropertyRange::init(sycl_device_binary Bin, diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp index 9aab698ef2749..546f88d2216d4 100644 --- a/sycl/source/detail/device_binary_image.hpp +++ b/sycl/source/detail/device_binary_image.hpp @@ -73,7 +73,7 @@ class DeviceBinaryProperty { uint32_t asUint32() const; ByteArray asByteArray() const; - const char *asCString() const; + std::string_view asStringView() const; protected: friend std::ostream &operator<<(std::ostream &Out, diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 035ca965ce2e5..3abcae421fcac 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -692,10 +692,11 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions( std::set HandledSets; std::queue WorkList; for (const sycl_device_binary_property &VFProp : Img.getVirtualFunctions()) { - std::string StrValue = DeviceBinaryProperty(VFProp).asCString(); + std::string_view StrValue = DeviceBinaryProperty(VFProp).asStringView(); // Device image passed to this function is expected to contain SYCL kernels // and therefore it may only use virtual function sets, but cannot provide - // them. We expect to see just a single property here + // them. Additionally, it cannot be a dummy image. + // We expect to see just a single property here assert(std::string(VFProp->Name) == "uses-virtual-functions-set" && "Unexpected virtual function property"); for (const auto &SetName : detail::split_string(StrValue, ',')) { @@ -716,22 +717,38 @@ ProgramManager::collectDependentDeviceImagesForVirtualFunctions( // virtual-functions-set properties, but their handling is the same: we // just grab all sets they reference and add them for consideration if // we haven't done so already. + bool isDummyImage = false; for (const sycl_device_binary_property &VFProp : BinImage->getVirtualFunctions()) { - std::string StrValue = DeviceBinaryProperty(VFProp).asCString(); + if (VFProp->Name == std::string_view("dummy-image")) { + isDummyImage = true; + continue; + } + std::string_view StrValue = DeviceBinaryProperty(VFProp).asStringView(); for (const auto &SetName : detail::split_string(StrValue, ',')) { if (HandledSets.insert(SetName).second) WorkList.push(SetName); } } - // TODO: Complete this part about handling of incompatible device images. // If device image uses the same virtual function set, then we only // link it if it is compatible. // However, if device image provides 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)) + // Note: we only link when exactly one of + // doesDevSupportDeviceRequirements(Dev, *BinImage) and + // isDummyImage is true. We don't want to link every dummy image, + // otherwise we could run into linking errors defining the same symbol + // multiple times. For every image providing virtual functions that has + // a dummy image, the dummy image will have the same device requirements + // as the original image. So when the dummy image does support the + // device requirements, we know that the corresponding image providing + // actual definitions will be linked and not the dummy. And vice versa: + // when the dummy image does not support the device requirements, we + // know the corresponding image providing virtual functions was not + // linked and we must link the dummy image. + if (doesDevSupportDeviceRequirements(Dev, *BinImage) ^ isDummyImage) DeviceImagesToLink.insert(BinImage); } } @@ -1850,7 +1867,9 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) { // Record mapping between virtual function sets and device images for (const sycl_device_binary_property &VFProp : Img->getVirtualFunctions()) { - std::string StrValue = DeviceBinaryProperty(VFProp).asCString(); + if (VFProp->Name == std::string_view("dummy-image")) + continue; + std::string_view StrValue = DeviceBinaryProperty(VFProp).asStringView(); for (const auto &SetName : detail::split_string(StrValue, ',')) m_VFSet2BinImage[SetName].insert(Img.get()); } diff --git a/sycl/unittests/Extensions/VirtualFunctions/RuntimeLinking.cpp b/sycl/unittests/Extensions/VirtualFunctions/RuntimeLinking.cpp index 61e0c5f5f32a5..3ebc7d6ce7fd8 100644 --- a/sycl/unittests/Extensions/VirtualFunctions/RuntimeLinking.cpp +++ b/sycl/unittests/Extensions/VirtualFunctions/RuntimeLinking.cpp @@ -18,6 +18,7 @@ class KernelD; class KernelE; class KernelF; class KernelG; +class KernelH; } // namespace VirtualFunctionsTest @@ -39,6 +40,7 @@ KERNEL_INFO(KernelD) KERNEL_INFO(KernelE) KERNEL_INFO(KernelF) KERNEL_INFO(KernelG) +KERNEL_INFO(KernelH) #undef KERNEL_INFO @@ -48,9 +50,13 @@ KERNEL_INFO(KernelG) static sycl::unittest::MockDeviceImage generateImage(std::initializer_list KernelNames, - const std::string &VFSets, bool UsesVFSets, unsigned char Magic) { + const std::string &VFSets, bool UsesVFSets, unsigned char Magic, + bool IsDummyImage = false, + std::vector Aspects = {}) { sycl::unittest::MockPropertySet PropSet; - std::vector Props; + + // Construct virtual function properties + std::vector VFProps; uint64_t PropSize = VFSets.size(); std::vector Storage(/* bytes for size */ 8 + PropSize + /* null terminator */ 1); @@ -64,9 +70,22 @@ generateImage(std::initializer_list KernelNames, sycl::unittest::MockProperty Prop(PropName, Storage, SYCL_PROPERTY_TYPE_BYTE_ARRAY); - Props.push_back(Prop); - PropSet.insert(__SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS, std::move(Props)); + VFProps.push_back(Prop); + if (IsDummyImage) + VFProps.emplace_back("dummy-image", std::vector(4), + SYCL_PROPERTY_TYPE_UINT32); + + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_VIRTUAL_FUNCTIONS, + std::move(VFProps)); + + // Construct device requirement properties + std::vector DeviceRequirmentsProps; + DeviceRequirmentsProps.emplace_back(sycl::unittest::makeAspectsProp(Aspects)); + PropSet.insert(__SYCL_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS, + std::move(DeviceRequirmentsProps)); + + // Assemble final device image std::vector Bin{Magic}; std::vector Entries = @@ -99,6 +118,9 @@ static constexpr unsigned PROGRAM_E0 = 37; static constexpr unsigned PROGRAM_F = 41; static constexpr unsigned PROGRAM_F0 = 47; static constexpr unsigned PROGRAM_F1 = 53; +static constexpr unsigned PROGRAM_H = 59; +static constexpr unsigned PROGRAM_H0 = 61; +static constexpr unsigned PROGRAM_H0d = 67; // Device images with no entires are ignored by SYCL RT during registration. // Therefore, we have to provide some kernel names to make the test work, even @@ -128,10 +150,16 @@ static sycl::unittest::MockDeviceImage Imgs[] = { generateImage({"KernelF"}, "set-f", /* uses vf set */ true, PROGRAM_F), generateImage({"DummyKernel7"}, "set-f", /* provides vf set */ false, PROGRAM_F0), - generateImage({"KernelG"}, "set-f", /* uses vf set */ true, PROGRAM_F1)}; + generateImage({"KernelG"}, "set-f", /* uses vf set */ true, PROGRAM_F1), + generateImage({"KernelH"}, "set-h", /* uses vf set */ true, PROGRAM_H, + false, {}), + generateImage({"DummyKernel7"}, "set-h", /* provides vf set */ false, + PROGRAM_H0, /* isDummy */ false, {sycl::aspect::fp64}), + generateImage({"DummyKernel7d"}, "set-h", /* provides vf set */ false, + PROGRAM_H0d, /* isDummy */ true, {sycl::aspect::fp64})}; // Registers mock devices images in the SYCL RT -static sycl::unittest::MockDeviceImageArray<15> ImgArray{Imgs}; +static sycl::unittest::MockDeviceImageArray ImgArray{Imgs}; TEST(VirtualFunctions, SingleKernelUsesSingleVFSet) { sycl::unittest::UrMock<> Mock; @@ -262,4 +290,101 @@ TEST(VirtualFunctions, TwoKernelsShareTheSameSet) { PROGRAM_F * PROGRAM_F0 * PROGRAM_F1); } +struct MockDeviceData { + std::string Extensions; + ur_device_handle_t getHandle() { + return reinterpret_cast(this); + } + static MockDeviceData *fromHandle(ur_device_handle_t handle) { + return reinterpret_cast(handle); + } +}; + +MockDeviceData MockDevices[] = { + {"cl_khr_fp64"}, + {""}, +}; + +static ur_result_t redefinedDeviceGet(void *pParams) { + auto params = *static_cast(pParams); + if (*params.ppNumDevices) { + **params.ppNumDevices = static_cast(std::size(MockDevices)); + return UR_RESULT_SUCCESS; + } + + if (*params.pphDevices) { + assert(*params.pNumEntries <= std::size(MockDevices)); + for (uint32_t i = 0; i < *params.pNumEntries; ++i) { + (*params.pphDevices)[i] = MockDevices[i].getHandle(); + } + } + + return UR_RESULT_SUCCESS; +} + +static ur_result_t redefinedDeviceGetInfo(void *pParams) { + auto *params = reinterpret_cast(pParams); + if (*params->ppropName == UR_DEVICE_INFO_EXTENSIONS) { + const std::string &Extensions = + MockDeviceData::fromHandle(*params->phDevice)->Extensions; + if (*params->ppPropValue) { + assert(*params->ppropSize >= Extensions.size() + 1); + std::memcpy(*params->ppPropValue, Extensions.data(), + Extensions.size() + 1); + } + if (*params->ppPropSizeRet && + **params->ppPropSizeRet < Extensions.size() + 1) + **params->ppPropSizeRet = Extensions.size() + 1; + return UR_RESULT_SUCCESS; + } + return UR_RESULT_SUCCESS; +} + +TEST(VirtualFunctions, DummyImages) { + sycl::unittest::UrMock<> Mock; + setupRuntimeLinkingMock(); + mock::getCallbacks().set_after_callback("urDeviceGet", &redefinedDeviceGet); + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &redefinedDeviceGetInfo); + + sycl::platform Plt = sycl::platform(); + sycl::queue Q(sycl::aspect_selector({sycl::aspect::fp64})); + EXPECT_TRUE(Q.get_device().has(sycl::aspect::fp64)); + + CapturedLinkingData.clear(); + + // KernelF uses set "set-h" that is also used by KernelG + Q.single_task([=]() {}); + // When we submit this kernel, we expect that two programs were created (one + // for KernelH, another providing "set-h" + EXPECT_EQ(CapturedLinkingData.NumOfUrProgramCreateCalls, 2u); + // Both programs should be linked together. + EXPECT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + // The module providing set-h is set up to use fp64, + // and since the device support fp64, we link the + // non-dummy version that provides set-h. + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PROGRAM_H, PROGRAM_H0})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PROGRAM_H * PROGRAM_H0); + + CapturedLinkingData.clear(); + + EXPECT_EQ(Plt.get_devices().size(), 2u); + sycl::queue Q2(sycl::aspect_selector({}, {sycl::aspect::fp64})); + + // We now repeat what we did launching KernelH but on another + // device that does not support fp64. + Q2.single_task([=]() {}); + EXPECT_EQ(CapturedLinkingData.NumOfUrProgramCreateCalls, 2u); + EXPECT_EQ(CapturedLinkingData.NumOfUrProgramLinkCalls, 1u); + + // However, this time, we expect the dummy image to be linked + // as the device does not support fp64. + EXPECT_TRUE( + CapturedLinkingData.LinkedProgramsContains({PROGRAM_H, PROGRAM_H0d})); + EXPECT_EQ(CapturedLinkingData.ProgramUsedToCreateKernel, + PROGRAM_H * PROGRAM_H0d); +} + // TODO: Add test cases for kernel_bundle usage diff --git a/sycl/unittests/helpers/MockDeviceImage.hpp b/sycl/unittests/helpers/MockDeviceImage.hpp index fea80d6b08c3e..0c2d375c3e810 100644 --- a/sycl/unittests/helpers/MockDeviceImage.hpp +++ b/sycl/unittests/helpers/MockDeviceImage.hpp @@ -558,7 +558,7 @@ inline MockProperty makeAspectsProp(const std::vector &Aspects) { uint64_t ValDataSize = ValData.size(); std::uninitialized_copy(&ValDataSize, &ValDataSize + sizeof(uint64_t), ValData.data()); - auto *AspectsPtr = reinterpret_cast(&Aspects[0]); + auto *AspectsPtr = reinterpret_cast(Aspects.data()); std::uninitialized_copy(AspectsPtr, AspectsPtr + Aspects.size(), ValData.data() + BYTES_FOR_SIZE); return {"aspects", ValData, SYCL_PROPERTY_TYPE_BYTE_ARRAY};