Skip to content

Commit df48d76

Browse files
author
Steffen Larsen
authored
[SYCL][NFCI] Move device_global map into separate structure (#19084)
This commit moves the implementation and protection of the device_global variable registration entries into a separate structure for cleaner separation. This should make it simpler to let kernel_bundle manage lifetime and isolate device_global variables for SYCLBIN in a future patch. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
1 parent 8589076 commit df48d76

5 files changed

Lines changed: 175 additions & 118 deletions

File tree

‎sycl/source/detail/context_impl.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -412,7 +412,7 @@ context_impl::initializeDeviceGlobals(ur_program_handle_t NativePrg,
412412

413413
// Device global map entry pointers will not die before the end of the
414414
// program and the pointers will stay the same, so we do not need
415-
// m_DeviceGlobalsMutex here.
415+
// to lock the device global map here.
416416
// The lifetimes of device global map entries representing globals in
417417
// runtime-compiled code will be tied to the kernel bundle, so the
418418
// assumption holds in that setting as well.

‎sycl/source/detail/device_global_map.hpp‎

Lines changed: 157 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,157 @@
1+
//==-------------------- device_global_map.hpp -----------------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <mutex>
12+
#include <unordered_map>
13+
14+
#include <detail/compiler.hpp>
15+
#include <detail/device_binary_image.hpp>
16+
#include <detail/device_global_map_entry.hpp>
17+
#include <sycl/detail/defines_elementary.hpp>
18+
#include <sycl/detail/kernel_name_str_t.hpp>
19+
20+
namespace sycl {
21+
inline namespace _V1 {
22+
namespace detail {
23+
24+
class DeviceGlobalMap {
25+
public:
26+
void initializeEntries(RTDeviceBinaryImage *Img) {
27+
const auto &DeviceGlobals = Img->getDeviceGlobals();
28+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
29+
for (const sycl_device_binary_property &DeviceGlobal : DeviceGlobals) {
30+
ByteArray DeviceGlobalInfo =
31+
DeviceBinaryProperty(DeviceGlobal).asByteArray();
32+
33+
// The supplied device_global info property is expected to contain:
34+
// * 8 bytes - Size of the property.
35+
// * 4 bytes - Size of the underlying type in the device_global.
36+
// * 4 bytes - 0 if device_global has device_image_scope and any value
37+
// otherwise.
38+
DeviceGlobalInfo.dropBytes(8);
39+
auto [TypeSize, DeviceImageScopeDecorated] =
40+
DeviceGlobalInfo.consume<std::uint32_t, std::uint32_t>();
41+
assert(DeviceGlobalInfo.empty() && "Extra data left!");
42+
43+
// Give the image pointer as an identifier for the image the
44+
// device-global is associated with.
45+
46+
auto ExistingDeviceGlobal = MDeviceGlobals.find(DeviceGlobal->Name);
47+
if (ExistingDeviceGlobal != MDeviceGlobals.end()) {
48+
// If it has already been registered we update the information.
49+
ExistingDeviceGlobal->second->initialize(Img, TypeSize,
50+
DeviceImageScopeDecorated);
51+
} else {
52+
// If it has not already been registered we create a new entry.
53+
// Note: Pointer to the device global is not available here, so it
54+
// cannot be set until registration happens.
55+
auto EntryUPtr = std::make_unique<DeviceGlobalMapEntry>(
56+
DeviceGlobal->Name, Img, TypeSize, DeviceImageScopeDecorated);
57+
MDeviceGlobals.emplace(DeviceGlobal->Name, std::move(EntryUPtr));
58+
}
59+
}
60+
}
61+
62+
void eraseEntries(const RTDeviceBinaryImage *Img) {
63+
const auto &DeviceGlobals = Img->getDeviceGlobals();
64+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
65+
for (const sycl_device_binary_property &DeviceGlobal : DeviceGlobals) {
66+
if (auto DevGlobalIt = MDeviceGlobals.find(DeviceGlobal->Name);
67+
DevGlobalIt != MDeviceGlobals.end()) {
68+
auto findDevGlobalByValue = std::find_if(
69+
MPtr2DeviceGlobal.begin(), MPtr2DeviceGlobal.end(),
70+
[&DevGlobalIt](
71+
const std::pair<const void *, DeviceGlobalMapEntry *> &Entry) {
72+
return Entry.second == DevGlobalIt->second.get();
73+
});
74+
if (findDevGlobalByValue != MPtr2DeviceGlobal.end())
75+
MPtr2DeviceGlobal.erase(findDevGlobalByValue);
76+
MDeviceGlobals.erase(DevGlobalIt);
77+
}
78+
}
79+
}
80+
81+
void addOrInitialize(const void *DeviceGlobalPtr, const char *UniqueId) {
82+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
83+
auto ExistingDeviceGlobal = MDeviceGlobals.find(UniqueId);
84+
if (ExistingDeviceGlobal != MDeviceGlobals.end()) {
85+
// Update the existing information and add the entry to the pointer map.
86+
ExistingDeviceGlobal->second->initialize(DeviceGlobalPtr);
87+
MPtr2DeviceGlobal.insert(
88+
{DeviceGlobalPtr, ExistingDeviceGlobal->second.get()});
89+
return;
90+
}
91+
92+
auto EntryUPtr =
93+
std::make_unique<DeviceGlobalMapEntry>(UniqueId, DeviceGlobalPtr);
94+
auto NewEntry = MDeviceGlobals.emplace(UniqueId, std::move(EntryUPtr));
95+
MPtr2DeviceGlobal.insert({DeviceGlobalPtr, NewEntry.first->second.get()});
96+
}
97+
98+
DeviceGlobalMapEntry *getEntry(const void *DeviceGlobalPtr) {
99+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
100+
auto Entry = MPtr2DeviceGlobal.find(DeviceGlobalPtr);
101+
assert(Entry != MPtr2DeviceGlobal.end() && "Device global entry not found");
102+
return Entry->second;
103+
}
104+
105+
DeviceGlobalMapEntry *tryGetEntry(const std::string &UniqueId,
106+
bool ExcludeDeviceImageScopeDecorated) {
107+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
108+
auto DeviceGlobalEntry = MDeviceGlobals.find(UniqueId);
109+
if (DeviceGlobalEntry != MDeviceGlobals.end() &&
110+
(!ExcludeDeviceImageScopeDecorated ||
111+
!DeviceGlobalEntry->second->MIsDeviceImageScopeDecorated))
112+
return DeviceGlobalEntry->second.get();
113+
return nullptr;
114+
}
115+
116+
std::vector<DeviceGlobalMapEntry *>
117+
getEntries(const std::vector<std::string> &UniqueIds,
118+
bool ExcludeDeviceImageScopeDecorated) {
119+
std::vector<DeviceGlobalMapEntry *> FoundEntries;
120+
FoundEntries.reserve(UniqueIds.size());
121+
122+
std::lock_guard<std::mutex> DeviceGlobalsGuard(MDeviceGlobalsMutex);
123+
for (const std::string &UniqueId : UniqueIds) {
124+
auto DeviceGlobalEntry = MDeviceGlobals.find(UniqueId);
125+
assert(DeviceGlobalEntry != MDeviceGlobals.end() &&
126+
"Device global not found in map.");
127+
if (!ExcludeDeviceImageScopeDecorated ||
128+
!DeviceGlobalEntry->second->MIsDeviceImageScopeDecorated)
129+
FoundEntries.push_back(DeviceGlobalEntry->second.get());
130+
}
131+
return FoundEntries;
132+
}
133+
134+
const std::unordered_map<const void *, DeviceGlobalMapEntry *>
135+
getPointerMap() const {
136+
return MPtr2DeviceGlobal;
137+
}
138+
139+
size_t size() const { return MDeviceGlobals.size(); }
140+
141+
size_t count(const KernelNameStrT &UniqueId) const {
142+
return MDeviceGlobals.count(UniqueId);
143+
}
144+
145+
private:
146+
// Maps between device_global identifiers and associated information.
147+
std::unordered_map<KernelNameStrT, std::unique_ptr<DeviceGlobalMapEntry>>
148+
MDeviceGlobals;
149+
std::unordered_map<const void *, DeviceGlobalMapEntry *> MPtr2DeviceGlobal;
150+
151+
/// Protects MDeviceGlobals and MPtr2DeviceGlobal.
152+
std::mutex MDeviceGlobalsMutex;
153+
};
154+
155+
} // namespace detail
156+
} // namespace _V1
157+
} // namespace sycl

‎sycl/source/detail/program_manager/program_manager.cpp‎

Lines changed: 8 additions & 95 deletions
Original file line numberDiff line numberDiff line change
@@ -2115,42 +2115,7 @@ void ProgramManager::addImage(sycl_device_binary RawImg,
21152115
KernelIDs->end());
21162116

21172117
// ... and initialize associated device_global information
2118-
{
2119-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2120-
2121-
auto DeviceGlobals = Img->getDeviceGlobals();
2122-
for (const sycl_device_binary_property &DeviceGlobal : DeviceGlobals) {
2123-
ByteArray DeviceGlobalInfo =
2124-
DeviceBinaryProperty(DeviceGlobal).asByteArray();
2125-
2126-
// The supplied device_global info property is expected to contain:
2127-
// * 8 bytes - Size of the property.
2128-
// * 4 bytes - Size of the underlying type in the device_global.
2129-
// * 4 bytes - 0 if device_global has device_image_scope and any value
2130-
// otherwise.
2131-
DeviceGlobalInfo.dropBytes(8);
2132-
auto [TypeSize, DeviceImageScopeDecorated] =
2133-
DeviceGlobalInfo.consume<std::uint32_t, std::uint32_t>();
2134-
assert(DeviceGlobalInfo.empty() && "Extra data left!");
2135-
2136-
// Give the image pointer as an identifier for the image the
2137-
// device-global is associated with.
2138-
2139-
auto ExistingDeviceGlobal = m_DeviceGlobals.find(DeviceGlobal->Name);
2140-
if (ExistingDeviceGlobal != m_DeviceGlobals.end()) {
2141-
// If it has already been registered we update the information.
2142-
ExistingDeviceGlobal->second->initialize(Img.get(), TypeSize,
2143-
DeviceImageScopeDecorated);
2144-
} else {
2145-
// If it has not already been registered we create a new entry.
2146-
// Note: Pointer to the device global is not available here, so it
2147-
// cannot be set until registration happens.
2148-
auto EntryUPtr = std::make_unique<DeviceGlobalMapEntry>(
2149-
DeviceGlobal->Name, Img.get(), TypeSize, DeviceImageScopeDecorated);
2150-
m_DeviceGlobals.emplace(DeviceGlobal->Name, std::move(EntryUPtr));
2151-
}
2152-
}
2153-
}
2118+
m_DeviceGlobals.initializeEntries(Img.get());
21542119
// ... and initialize associated host_pipe information
21552120
{
21562121
std::lock_guard<std::mutex> HostPipesGuard(m_HostPipesMutex);
@@ -2257,24 +2222,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
22572222
m_VFSet2BinImage.erase(SetName);
22582223
}
22592224

2260-
{
2261-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2262-
auto DeviceGlobals = Img->getDeviceGlobals();
2263-
for (const sycl_device_binary_property &DeviceGlobal : DeviceGlobals) {
2264-
if (auto DevGlobalIt = m_DeviceGlobals.find(DeviceGlobal->Name);
2265-
DevGlobalIt != m_DeviceGlobals.end()) {
2266-
auto findDevGlobalByValue = std::find_if(
2267-
m_Ptr2DeviceGlobal.begin(), m_Ptr2DeviceGlobal.end(),
2268-
[&DevGlobalIt](const std::pair<const void *,
2269-
DeviceGlobalMapEntry *> &Entry) {
2270-
return Entry.second == DevGlobalIt->second.get();
2271-
});
2272-
if (findDevGlobalByValue != m_Ptr2DeviceGlobal.end())
2273-
m_Ptr2DeviceGlobal.erase(findDevGlobalByValue);
2274-
m_DeviceGlobals.erase(DevGlobalIt);
2275-
}
2276-
}
2277-
}
2225+
m_DeviceGlobals.eraseEntries(Img);
22782226

22792227
{
22802228
std::lock_guard<std::mutex> HostPipesGuard(m_HostPipesMutex);
@@ -2468,21 +2416,7 @@ kernel_id ProgramManager::getBuiltInKernelID(KernelNameStrRefT KernelName) {
24682416

24692417
void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
24702418
const char *UniqueId) {
2471-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2472-
2473-
auto ExistingDeviceGlobal = m_DeviceGlobals.find(UniqueId);
2474-
if (ExistingDeviceGlobal != m_DeviceGlobals.end()) {
2475-
// Update the existing information and add the entry to the pointer map.
2476-
ExistingDeviceGlobal->second->initialize(DeviceGlobalPtr);
2477-
m_Ptr2DeviceGlobal.insert(
2478-
{DeviceGlobalPtr, ExistingDeviceGlobal->second.get()});
2479-
return;
2480-
}
2481-
2482-
auto EntryUPtr =
2483-
std::make_unique<DeviceGlobalMapEntry>(UniqueId, DeviceGlobalPtr);
2484-
auto NewEntry = m_DeviceGlobals.emplace(UniqueId, std::move(EntryUPtr));
2485-
m_Ptr2DeviceGlobal.insert({DeviceGlobalPtr, NewEntry.first->second.get()});
2419+
m_DeviceGlobals.addOrInitialize(DeviceGlobalPtr, UniqueId);
24862420
}
24872421

24882422
std::set<RTDeviceBinaryImage *>
@@ -2499,42 +2433,21 @@ ProgramManager::getRawDeviceImages(const std::vector<kernel_id> &KernelIDs) {
24992433

25002434
DeviceGlobalMapEntry *
25012435
ProgramManager::getDeviceGlobalEntry(const void *DeviceGlobalPtr) {
2502-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2503-
auto Entry = m_Ptr2DeviceGlobal.find(DeviceGlobalPtr);
2504-
assert(Entry != m_Ptr2DeviceGlobal.end() && "Device global entry not found");
2505-
return Entry->second;
2436+
return m_DeviceGlobals.getEntry(DeviceGlobalPtr);
25062437
}
25072438

25082439
DeviceGlobalMapEntry *
25092440
ProgramManager::tryGetDeviceGlobalEntry(const std::string &UniqueId,
25102441
bool ExcludeDeviceImageScopeDecorated) {
2511-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2512-
auto DeviceGlobalEntry = m_DeviceGlobals.find(UniqueId);
2513-
assert(DeviceGlobalEntry != m_DeviceGlobals.end() &&
2514-
"Device global not found in map.");
2515-
if (DeviceGlobalEntry != m_DeviceGlobals.end() &&
2516-
(!ExcludeDeviceImageScopeDecorated ||
2517-
!DeviceGlobalEntry->second->MIsDeviceImageScopeDecorated))
2518-
return DeviceGlobalEntry->second.get();
2519-
return nullptr;
2442+
return m_DeviceGlobals.tryGetEntry(UniqueId,
2443+
ExcludeDeviceImageScopeDecorated);
25202444
}
25212445

25222446
std::vector<DeviceGlobalMapEntry *> ProgramManager::getDeviceGlobalEntries(
25232447
const std::vector<std::string> &UniqueIds,
25242448
bool ExcludeDeviceImageScopeDecorated) {
2525-
std::vector<DeviceGlobalMapEntry *> FoundEntries;
2526-
FoundEntries.reserve(UniqueIds.size());
2527-
2528-
std::lock_guard<std::mutex> DeviceGlobalsGuard(m_DeviceGlobalsMutex);
2529-
for (const std::string &UniqueId : UniqueIds) {
2530-
auto DeviceGlobalEntry = m_DeviceGlobals.find(UniqueId);
2531-
assert(DeviceGlobalEntry != m_DeviceGlobals.end() &&
2532-
"Device global not found in map.");
2533-
if (!ExcludeDeviceImageScopeDecorated ||
2534-
!DeviceGlobalEntry->second->MIsDeviceImageScopeDecorated)
2535-
FoundEntries.push_back(DeviceGlobalEntry->second.get());
2536-
}
2537-
return FoundEntries;
2449+
return m_DeviceGlobals.getEntries(UniqueIds,
2450+
ExcludeDeviceImageScopeDecorated);
25382451
}
25392452

25402453
void ProgramManager::addOrInitHostPipeEntry(const void *HostPipePtr,

‎sycl/source/detail/program_manager/program_manager.hpp‎

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
#pragma once
1010
#include <detail/cg.hpp>
1111
#include <detail/device_binary_image.hpp>
12+
#include <detail/device_global_map.hpp>
1213
#include <detail/device_global_map_entry.hpp>
1314
#include <detail/host_pipe_map_entry.hpp>
1415
#include <detail/kernel_arg_mask.hpp>
@@ -530,12 +531,7 @@ class ProgramManager {
530531
SanitizerType m_SanitizerFoundInImage;
531532

532533
// Maps between device_global identifiers and associated information.
533-
std::unordered_map<KernelNameStrT, std::unique_ptr<DeviceGlobalMapEntry>>
534-
m_DeviceGlobals;
535-
std::unordered_map<const void *, DeviceGlobalMapEntry *> m_Ptr2DeviceGlobal;
536-
537-
/// Protects m_DeviceGlobals and m_Ptr2DeviceGlobal.
538-
std::mutex m_DeviceGlobalsMutex;
534+
DeviceGlobalMap m_DeviceGlobals;
539535

540536
// Maps between host_pipe identifiers and associated information.
541537
std::unordered_map<std::string, std::unique_ptr<HostPipeMapEntry>>

‎sycl/unittests/program_manager/Cleanup.cpp‎

Lines changed: 7 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,7 @@
11
#include <sycl/sycl.hpp>
22

33
#include <detail/device_binary_image.hpp>
4+
#include <detail/device_global_map.hpp>
45
#include <detail/device_image_impl.hpp>
56
#include <detail/program_manager/program_manager.hpp>
67
#include <helpers/MockDeviceImage.hpp>
@@ -85,16 +86,7 @@ class ProgramManagerExposed : public sycl::detail::ProgramManager {
8586
return m_Ptr2HostPipe;
8687
}
8788

88-
std::unordered_map<sycl::detail::KernelNameStrT,
89-
std::unique_ptr<sycl::detail::DeviceGlobalMapEntry>> &
90-
getDeviceGlobals() {
91-
return m_DeviceGlobals;
92-
}
93-
94-
std::unordered_map<const void *, sycl::detail::DeviceGlobalMapEntry *> &
95-
getPtrToDeviceGlobal() {
96-
return m_Ptr2DeviceGlobal;
97-
}
89+
sycl::detail::DeviceGlobalMap &getDeviceGlobals() { return m_DeviceGlobals; }
9890
};
9991

10092
namespace {
@@ -312,15 +304,14 @@ void checkAllInvolvedContainers(ProgramManagerExposed &PM, size_t ExpectedCount,
312304
EXPECT_EQ(PM.getKernelImplicitLocalArgPos().size(), ExpectedCount) << Comment;
313305

314306
{
315-
EXPECT_EQ(PM.getDeviceGlobals().size(), ExpectedCount) << Comment;
316-
EXPECT_TRUE(
317-
PM.getDeviceGlobals().count(generateRefName("A", "DeviceGlobal")) > 0)
307+
sycl::detail::DeviceGlobalMap &DeviceGlobalMap = PM.getDeviceGlobals();
308+
EXPECT_EQ(DeviceGlobalMap.size(), ExpectedCount) << Comment;
309+
EXPECT_TRUE(DeviceGlobalMap.count(generateRefName("A", "DeviceGlobal")) > 0)
318310
<< Comment;
319-
EXPECT_TRUE(
320-
PM.getDeviceGlobals().count(generateRefName("B", "DeviceGlobal")) > 0)
311+
EXPECT_TRUE(DeviceGlobalMap.count(generateRefName("B", "DeviceGlobal")) > 0)
321312
<< Comment;
313+
EXPECT_EQ(DeviceGlobalMap.getPointerMap().size(), ExpectedCount) << Comment;
322314
}
323-
EXPECT_EQ(PM.getPtrToDeviceGlobal().size(), ExpectedCount) << Comment;
324315

325316
{
326317
EXPECT_EQ(PM.getHostPipes().size(), ExpectedCount) << Comment;

0 commit comments

Comments
 (0)