Skip to content

Use device_images_range #3

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed

Use device_images_range #3

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 16 additions & 0 deletions sycl/source/detail/device_image_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1354,13 +1354,29 @@ class device_image_impl
using device_images_iterator =
variadic_iterator<device_image_plain,
std::vector<device_image_plain>::const_iterator,
const device_image_plain *,
std::set<device_image_impl *>::const_iterator>;
class device_images_range : public iterator_range<device_images_iterator> {
private:
using Base = iterator_range<device_images_iterator>;

public:
using Base::Base;
template <typename Container>
decltype(std::declval<Base>().to<Container>()) to() const {
return this->Base::to<Container>();
}

template <typename Container>
std::enable_if_t<std::is_same_v<Container, std::vector<ur_program_handle_t>>,
Container>
to() const {
std::vector<ur_program_handle_t> ProgramHandles;
ProgramHandles.reserve(size());
std::transform(begin(), end(), std::back_inserter(ProgramHandles),
[](device_image_impl &Img) { return Img.get_ur_program(); });
return ProgramHandles;
}
};

} // namespace detail
Expand Down
18 changes: 6 additions & 12 deletions sycl/source/detail/kernel_bundle_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,12 +114,11 @@ CreateLinkGraph(const std::vector<device_image_plain> &DevImages) {
}

inline void
ThrowIfConflictingKernels(const std::vector<device_image_plain> &DevImages) {
ThrowIfConflictingKernels(device_images_range DevImages) {
std::set<std::string_view, std::less<>> SeenKernelNames;
std::set<std::string_view, std::less<>> Conflicts;
for (const device_image_plain &DevImage : DevImages) {
const KernelNameSetT &KernelNames =
getSyclObjImpl(DevImage)->getKernelNames();
for (const device_image_impl &DevImage : DevImages) {
const KernelNameSetT &KernelNames = DevImage.getKernelNames();
std::vector<std::string_view> Intersect;
std::set_intersection(SeenKernelNames.begin(), SeenKernelNames.end(),
KernelNames.begin(), KernelNames.end(),
Expand Down Expand Up @@ -513,9 +512,7 @@ class kernel_bundle_impl
// In dynamic linking, AOT binaries count as results as well.
LinkedResults.insert(LinkedResults.end(), AOTImgs.begin(),
AOTImgs.end());
sycl::span<device_image_plain, dynamic_extent> LinkedResultsSpan(
LinkedResults.data(), LinkedResults.size());
detail::ProgramManager::getInstance().dynamicLink(LinkedResultsSpan);
detail::ProgramManager::getInstance().dynamicLink(LinkedResults);
}

MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
Expand All @@ -536,12 +533,9 @@ class kernel_bundle_impl
}))
continue;

const std::vector<device_image_plain> &AllDevImgs =
DeviceImageWithDeps->getAll();
sycl::span<const device_image_plain> AllDevImgsSpan(AllDevImgs);
std::vector<device_image_plain> LinkedResults =
detail::ProgramManager::getInstance().link(AllDevImgsSpan, MDevices,
PropList);
detail::ProgramManager::getInstance().link(
DeviceImageWithDeps->getAll(), MDevices, PropList);
MDeviceImages.insert(MDeviceImages.end(), LinkedResults.begin(),
LinkedResults.end());
MUniqueDeviceImages.insert(MUniqueDeviceImages.end(),
Expand Down
65 changes: 27 additions & 38 deletions sycl/source/detail/program_manager/program_manager.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2919,24 +2919,22 @@ ProgramManager::compile(const DevImgPlainWithDeps &ImgWithDeps,
// Returns a merged device binary image, new set of kernel IDs and new
// specialization constant data.
static const RTDeviceBinaryImage *
mergeImageData(sycl::span<const device_image_plain> Imgs,
std::vector<kernel_id> &KernelIDs,
mergeImageData(device_images_range Imgs, std::vector<kernel_id> &KernelIDs,
std::vector<unsigned char> &NewSpecConstBlob,
device_image_impl::SpecConstMapT &NewSpecConstMap,
std::unique_ptr<DynRTDeviceBinaryImage> &MergedImageStorage) {
for (const device_image_plain &Img : Imgs) {
device_image_impl &DeviceImageImpl = *getSyclObjImpl(Img);
for (device_image_impl &Img : Imgs) {
// Duplicates are not expected here, otherwise urProgramLink should fail
KernelIDs.insert(KernelIDs.end(), DeviceImageImpl.get_kernel_ids().begin(),
DeviceImageImpl.get_kernel_ids().end());
KernelIDs.insert(KernelIDs.end(), Img.get_kernel_ids().begin(),
Img.get_kernel_ids().end());
// To be able to answer queries about specialziation constants, the new
// device image should have the specialization constants from all the linked
// images.
const std::lock_guard<std::mutex> SpecConstLock(
DeviceImageImpl.get_spec_const_data_lock());
Img.get_spec_const_data_lock());
// Copy all map entries to the new map. Since the blob will be copied to
// the end of the new blob we need to move the blob offset of each entry.
for (const auto &SpecConstIt : DeviceImageImpl.get_spec_const_data_ref()) {
for (const auto &SpecConstIt : Img.get_spec_const_data_ref()) {
std::vector<device_image_impl::SpecConstDescT> &NewDescEntries =
NewSpecConstMap[SpecConstIt.first];

Expand All @@ -2954,21 +2952,21 @@ mergeImageData(sycl::span<const device_image_plain> Imgs,
// Copy the blob from the device image into the new blob. This moves the
// offsets of the following blobs.
NewSpecConstBlob.insert(NewSpecConstBlob.end(),
DeviceImageImpl.get_spec_const_blob_ref().begin(),
DeviceImageImpl.get_spec_const_blob_ref().end());
Img.get_spec_const_blob_ref().begin(),
Img.get_spec_const_blob_ref().end());
}
// device_image_impl expects kernel ids to be sorted for fast search
std::sort(KernelIDs.begin(), KernelIDs.end(), LessByHash<kernel_id>{});

// If there is only a single image, use it as the result.
if (Imgs.size() == 1)
return getSyclObjImpl(Imgs[0])->get_bin_image_ref();
return Imgs.front().get_bin_image_ref();

// Otherwise we create a dynamic image with the merged information.
std::vector<const RTDeviceBinaryImage *> BinImgs;
BinImgs.reserve(Imgs.size());
for (const device_image_plain &Img : Imgs) {
auto ImgBinRef = getSyclObjImpl(Img)->get_bin_image_ref();
for (device_image_impl &Img : Imgs) {
auto ImgBinRef = Img.get_bin_image_ref();
// For some cases, like SYCL kernel compiler binaries, we don't have
// binaries. For these we assume no properties associated, so they can be
// safely ignored.
Expand All @@ -2980,26 +2978,22 @@ mergeImageData(sycl::span<const device_image_plain> Imgs,
}

std::vector<device_image_plain>
ProgramManager::link(sycl::span<const device_image_plain> Imgs,
devices_range Devs, const property_list &PropList,
ProgramManager::link(device_images_range Imgs, devices_range Devs,
const property_list &PropList,
bool AllowUnresolvedSymbols) {
{
auto NoAllowedPropertiesCheck = [](int) { return false; };
detail::PropertyValidator::checkPropsAndThrow(
PropList, NoAllowedPropertiesCheck, NoAllowedPropertiesCheck);
}

std::vector<ur_program_handle_t> URPrograms;
URPrograms.reserve(Imgs.size());
for (const device_image_plain &Img : Imgs)
URPrograms.push_back(getSyclObjImpl(Img)->get_ur_program());

auto URPrograms = Imgs.to<std::vector<ur_program_handle_t>>();
auto URDevices = Devs.to<std::vector<ur_device_handle_t>>();

// FIXME: Linker options are picked from the first object, but is that safe?
std::string LinkOptionsStr;
applyLinkOptionsFromEnvironment(LinkOptionsStr);
device_image_impl &FirstImgImpl = *getSyclObjImpl(Imgs[0]);
device_image_impl &FirstImgImpl = Imgs.front();
if (LinkOptionsStr.empty() && FirstImgImpl.get_bin_image_ref())
appendLinkOptionsFromImage(LinkOptionsStr,
*(FirstImgImpl.get_bin_image_ref()));
Expand Down Expand Up @@ -3061,8 +3055,8 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
// underlying program disposed of). Protecting from incorrect values by
// removal of map entries with same handle (obviously invalid entries).
std::ignore = NativePrograms.erase(LinkedProg);
for (const device_image_plain &Img : Imgs) {
if (auto BinImageRef = getSyclObjImpl(Img)->get_bin_image_ref())
for (device_image_impl &Img : Imgs) {
if (auto BinImageRef = Img.get_bin_image_ref())
NativePrograms.insert(
{LinkedProg, {ContextImpl.shared_from_this(), BinImageRef}});
}
Expand All @@ -3078,15 +3072,14 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
KernelNameSetT MergedKernelNames;
std::map<std::string, KernelArgMask, std::less<>>
MergedEliminatedKernelArgMasks;
for (const device_image_plain &DevImg : Imgs) {
device_image_impl &DevImgImpl = *getSyclObjImpl(DevImg);
CombinedOrigins |= DevImgImpl.getOriginMask();
RTCInfoPtrs.emplace_back(&(DevImgImpl.getRTCInfo()));
MergedKernelNames.insert(DevImgImpl.getKernelNames().begin(),
DevImgImpl.getKernelNames().end());
for (device_image_impl &DevImg : Imgs) {
CombinedOrigins |= DevImg.getOriginMask();
RTCInfoPtrs.emplace_back(&(DevImg.getRTCInfo()));
MergedKernelNames.insert(DevImg.getKernelNames().begin(),
DevImg.getKernelNames().end());
MergedEliminatedKernelArgMasks.insert(
DevImgImpl.getEliminatedKernelArgMasks().begin(),
DevImgImpl.getEliminatedKernelArgMasks().end());
DevImg.getEliminatedKernelArgMasks().begin(),
DevImg.getEliminatedKernelArgMasks().end());
}
auto MergedRTCInfo = detail::KernelCompilerBinaryInfo::Merge(RTCInfoPtrs);

Expand All @@ -3100,16 +3093,12 @@ ProgramManager::link(sycl::span<const device_image_plain> Imgs,
std::move(MergedImageStorage)))};
}

void ProgramManager::dynamicLink(sycl::span<const device_image_plain> Imgs) {
void ProgramManager::dynamicLink(device_images_range Imgs) {
if (Imgs.empty())
return;

std::vector<ur_program_handle_t> URPrograms;
URPrograms.reserve(Imgs.size());
for (const device_image_plain &Img : Imgs)
URPrograms.push_back(getSyclObjImpl(Img)->get_ur_program());

device_image_impl &FirstImgImpl = *getSyclObjImpl(Imgs[0]);
auto URPrograms = Imgs.to<std::vector<ur_program_handle_t>>();
device_image_impl &FirstImgImpl = Imgs.front();
auto [URCtx, Adapter] =
get_ur_handles(*getSyclObjImpl(FirstImgImpl.get_context()));
Adapter->call<UrApiKind::urProgramDynamicLinkExp>(URCtx, URPrograms.size(),
Expand Down
11 changes: 6 additions & 5 deletions sycl/source/detail/program_manager/program_manager.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,6 @@
#include <sycl/detail/util.hpp>
#include <sycl/device.hpp>
#include <sycl/kernel_bundle.hpp>
#include <sycl/sycl_span.hpp>

#include <array>
#include <cstdint>
Expand Down Expand Up @@ -77,6 +76,7 @@ class device_impl;
class devices_range;
class queue_impl;
class event_impl;
class device_images_range;
// DeviceLibExt is shared between sycl runtime and sycl-post-link tool.
// If any update is made here, need to sync with DeviceLibExt definition
// in llvm/tools/sycl-post-link/sycl-post-link.cpp
Expand Down Expand Up @@ -356,12 +356,13 @@ class ProgramManager {

// Produces set of device images by convering input device images to object
// the executable state
std::vector<device_image_plain>
link(sycl::span<const device_image_plain> Imgs, devices_range Devs,
const property_list &PropList, bool AllowUnresolvedSymbols = false);
std::vector<device_image_plain> link(device_images_range Imgs,
devices_range Devs,
const property_list &PropList,
bool AllowUnresolvedSymbols = false);

// Dynamically links images in executable state.
void dynamicLink(sycl::span<const device_image_plain> Imgs);
void dynamicLink(device_images_range Imgs);

// Produces new device image by converting input device image to the
// executable state
Expand Down