Skip to content

Commit 6b4e685

Browse files
uditagarwal97premanandrao
uditagarwal97
and
premanandrao
authored
[Clang][SYCL] Fix dangling pointers and ODR violation in free functions (#20422)
**Problems** Problem 1: When a library consisting of free function kernels is registered with SYCL RT, we store pointers (as `string_view`) to free function names in `m_FreeFunctionKernelGlobalInfo` but we do not remove them from `m_FreeFunctionKernelGlobalInfo` when the library is unloaded. Thus, we end up holding dangling pointers and any further operation on `m_FreeFunctionKernelGlobalInfo` might segfault. Problem 2: Consider the case when you have multiple TUs with free functions and they are compiled separately but linked together into a single shared lib. In that case, we will have multiple definition of `static GlobalMapUpdater updater` in the shared lib => violating ODR **Solution** Discard pointers to free function names when library is unloaded and have `GlobalMapUpdater` defined in anonymous namespace, instead of `sycl::v1::detail` --------- Co-authored-by: premanandrao <premanand.m.rao@intel.com>
1 parent faadbce commit 6b4e685

File tree

11 files changed

+65
-32
lines changed

11 files changed

+65
-32
lines changed

‎clang/lib/Sema/SemaSYCL.cpp‎

Lines changed: 14 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7302,22 +7302,30 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
73027302
}
73037303

73047304
if (FreeFunctionCount > 0) {
7305+
// GlobalMapUpdater has to be in an anonymous namespace.
7306+
// Otherwise, if multiple translation units include the same integration
7307+
// header, there will be multiple varying definitions of GlobalMapUpdater
7308+
// with the same name across translation units, violating the C++'s One
7309+
// Definition Rule. Putting it in an anonymous namespace gives each
7310+
// translation unit its own unique definition.
7311+
73057312
O << "\n#include <sycl/kernel_bundle.hpp>\n";
73067313
O << "#include <sycl/detail/kernel_global_info.hpp>\n";
7307-
O << "namespace sycl {\n";
7308-
O << "inline namespace _V1 {\n";
7309-
O << "namespace detail {\n";
7314+
O << "namespace {\n";
73107315
O << "struct GlobalMapUpdater {\n";
73117316
O << " GlobalMapUpdater() {\n";
73127317
O << " sycl::detail::free_function_info_map::add("
73137318
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
73147319
<< KernelDescs.size() << ");\n";
73157320
O << " }\n";
7321+
O << " ~GlobalMapUpdater() {\n";
7322+
O << " sycl::detail::free_function_info_map::remove("
7323+
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
7324+
<< KernelDescs.size() << ");\n";
7325+
O << " }\n";
73167326
O << "};\n";
73177327
O << "static GlobalMapUpdater updater;\n";
7318-
O << "} // namespace detail\n";
7319-
O << "} // namespace _V1\n";
7320-
O << "} // namespace sycl\n";
7328+
O << "} // namespace\n";
73217329
}
73227330
}
73237331

‎clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1087,15 +1087,14 @@ namespace Testing::Tests {
10871087

10881088
// CHECK: #include <sycl/kernel_bundle.hpp>
10891089
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
1090-
// CHECK-NEXT: namespace sycl {
1091-
// CHECK-NEXT: inline namespace _V1 {
1092-
// CHECK-NEXT: namespace detail {
1090+
// CHECK-NEXT: namespace {
10931091
// CHECK-NEXT: struct GlobalMapUpdater {
10941092
// CHECK-NEXT: GlobalMapUpdater() {
10951093
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
10961094
// CHECK-NEXT: }
1095+
// CHECK-NEXT: ~GlobalMapUpdater() {
1096+
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
1097+
// CHECK-NEXT: }
10971098
// CHECK-NEXT: };
10981099
// CHECK-NEXT: static GlobalMapUpdater updater;
1099-
// CHECK-NEXT: } // namespace detail
1100-
// CHECK-NEXT: } // namespace _V1
1101-
// CHECK-NEXT: } // namespace sycl
1100+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_int_header.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1567,15 +1567,14 @@ void ff_24(int arg) {
15671567

15681568
// CHECK: #include <sycl/kernel_bundle.hpp>
15691569
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
1570-
// CHECK-NEXT: namespace sycl {
1571-
// CHECK-NEXT: inline namespace _V1 {
1572-
// CHECK-NEXT: namespace detail {
1570+
// CHECK-NEXT: namespace {
15731571
// CHECK-NEXT: struct GlobalMapUpdater {
15741572
// CHECK-NEXT: GlobalMapUpdater() {
15751573
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
15761574
// CHECK-NEXT: }
1575+
// CHECK-NEXT: ~GlobalMapUpdater() {
1576+
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
1577+
// CHECK-NEXT: }
15771578
// CHECK-NEXT: };
15781579
// CHECK-NEXT: static GlobalMapUpdater updater;
1579-
// CHECK-NEXT: } // namespace detail
1580-
// CHECK-NEXT: } // namespace _V1
1581-
// CHECK-NEXT: } // namespace sycl
1580+
// CHECK-NEXT: }

‎clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -75,15 +75,14 @@ int main(){
7575

7676
// CHECK-NORTC: #include <sycl/kernel_bundle.hpp>
7777
// CHECK-NORTC-NEXT: #include <sycl/detail/kernel_global_info.hpp>
78-
// CHECK-NORTC-NEXT: namespace sycl {
79-
// CHECK-NORTC-NEXT: inline namespace _V1 {
80-
// CHECK-NORTC-NEXT: namespace detail {
78+
// CHECK-NORTC-NEXT: namespace {
8179
// CHECK-NORTC-NEXT: struct GlobalMapUpdater {
8280
// CHECK-NORTC-NEXT: GlobalMapUpdater() {
8381
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
8482
// CHECK-NORTC-NEXT: }
83+
// CHECK-NORTC-NEXT: ~GlobalMapUpdater() {
84+
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
85+
// CHECK-NORTC-NEXT: }
8586
// CHECK-NORTC-NEXT: };
8687
// CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
87-
// CHECK-NORTC-NEXT: } // namespace detail
88-
// CHECK-NORTC-NEXT: } // namespace _V1
89-
// CHECK-NORTC-NEXT: } // namespace sycl
88+
// CHECK-NORTC-NEXT: }

‎clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp‎

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -10,14 +10,11 @@
1010

1111
#include <sycl/detail/export.hpp>
1212

13-
namespace sycl {
14-
inline namespace _V1 {
15-
namespace detail {
13+
namespace {
1614
namespace free_function_info_map {
1715

1816
__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
17+
__SYCL_EXPORT void remove(const void *DeviceGlobalPtr, const char *UniqueId);
1918

2019
} // namespace free_function_info_map
21-
} // namespace detail
22-
} // namespace _V1
23-
} // namespace sycl
20+
}

‎sycl/include/sycl/detail/kernel_global_info.hpp‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,8 @@ namespace free_function_info_map {
1818
__SYCL_EXPORT void add(const char *const *UniqueId,
1919
const unsigned *DeviceGlobalPtr, unsigned Size);
2020

21+
__SYCL_EXPORT void remove(const char *const *UniqueId,
22+
const unsigned *DeviceGlobalPtr, unsigned Size);
2123
} // namespace free_function_info_map
2224
} // namespace detail
2325
} // namespace _V1

‎sycl/source/detail/kernel_global_info.cpp‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,17 @@ __SYCL_EXPORT void add(const char *const *FreeFunctionNames,
2323
std::move(GlobalInfoToCopy));
2424
}
2525

26+
__SYCL_EXPORT void remove(const char *const *FreeFunctionNames,
27+
const unsigned *FreeFunctionNumArgs, unsigned Size) {
28+
std::unordered_map<std::string_view, unsigned> GlobalInfoToCopy;
29+
for (size_t i = 0; i < Size; ++i) {
30+
GlobalInfoToCopy[std::string_view{FreeFunctionNames[i]}] =
31+
FreeFunctionNumArgs[i];
32+
}
33+
detail::ProgramManager::getInstance().unRegisterKernelGlobalInfo(
34+
std::move(GlobalInfoToCopy));
35+
}
36+
2637
} // namespace detail::free_function_info_map
2738
} // namespace _V1
2839
} // namespace sycl

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

Lines changed: 12 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2418,12 +2418,22 @@ void ProgramManager::registerKernelGlobalInfo(
24182418
if (m_FreeFunctionKernelGlobalInfo.empty())
24192419
m_FreeFunctionKernelGlobalInfo = std::move(GlobalInfoToCopy);
24202420
else {
2421-
for (auto &GlobalInfo : GlobalInfoToCopy) {
2421+
for (auto &GlobalInfo : GlobalInfoToCopy)
24222422
m_FreeFunctionKernelGlobalInfo.insert(GlobalInfo);
2423-
}
24242423
}
24252424
}
24262425

2426+
// Remove entries from m_FreeFunctionKernelGlobalInfo that matches
2427+
// the ones in GlobalInfoToCopy. This function is called when a shared
2428+
// library consisting of SYCL kernels is unloaded.
2429+
void ProgramManager::unRegisterKernelGlobalInfo(
2430+
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy) {
2431+
std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
2432+
2433+
for (const auto &GlobalInfo : GlobalInfoToCopy)
2434+
m_FreeFunctionKernelGlobalInfo.erase(GlobalInfo.first);
2435+
}
2436+
24272437
std::optional<unsigned>
24282438
ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
24292439
std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);

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

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -259,6 +259,12 @@ class ProgramManager {
259259
void registerKernelGlobalInfo(
260260
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);
261261

262+
// The function removes kernel global descriptors from the
263+
// kernel global map when a shared library consisting of SYCL kernels
264+
// is unloaded.
265+
void unRegisterKernelGlobalInfo(
266+
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);
267+
262268
// The function returns a pointer to the kernel global desc identified by
263269
// the unique ID from the kernel global map.
264270
std::optional<unsigned> getKernelGlobalInfoDesc(const char *UniqueId);

‎sycl/test/abi/sycl_symbols_linux.dump‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3344,6 +3344,7 @@ _ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEiiRKNS0_13property
33443344
_ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEiiRKNS0_13property_listE
33453345
_ZN4sycl3_V16detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE
33463346
_ZN4sycl3_V16detail22free_function_info_map3addEPKPKcPKjj
3347+
_ZN4sycl3_V16detail22free_function_info_map6removeEPKPKcPKjj
33473348
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE
33483349
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateERKSt8functionIFbRKSt10shared_ptrINS1_17device_image_implEEEE
33493350
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKNS0_4spanIcLm18446744073709551615EEENS0_12bundle_stateE

0 commit comments

Comments
 (0)