Skip to content

[Clang][SYCL] Fix dangling pointers and ODR violation in free functions #20422

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

Merged
merged 4 commits into from
Oct 31, 2025
Merged

[Clang][SYCL] Fix dangling pointers and ODR violation in free functions #20422

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
7e043a9
[Clang][SYCL] Discard pointers to free function descriptors when libr…
uditagarwal97 Oct 21, 2025
2079eb2
Move GlobalMapUpdater into anynymous namespace
uditagarwal97 Oct 30, 2025
d31600e
Add comment
uditagarwal97 Oct 30, 2025
4df8683
Update clang/lib/Sema/SemaSYCL.cpp
uditagarwal97 Oct 31, 2025
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
20 changes: 14 additions & 6 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7302,22 +7302,30 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
}

if (FreeFunctionCount > 0) {
// GlobalMapUpdater has to be in an anonymous namespace.
// Otherwise, if multiple translation units include the same integration
// header, there will be multiple varying definitions of GlobalMapUpdater
// with the same name across translation units, violating the C++'s One
// Definition Rule. Putting it in an anonymous namespace gives each
// translation unit its own unique definition.

O << "\n#include <sycl/kernel_bundle.hpp>\n";
O << "#include <sycl/detail/kernel_global_info.hpp>\n";
O << "namespace sycl {\n";
O << "inline namespace _V1 {\n";
O << "namespace detail {\n";
O << "namespace {\n";
O << "struct GlobalMapUpdater {\n";
O << " GlobalMapUpdater() {\n";
O << " sycl::detail::free_function_info_map::add("
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
<< KernelDescs.size() << ");\n";
O << " }\n";
O << " ~GlobalMapUpdater() {\n";
O << " sycl::detail::free_function_info_map::remove("
<< "sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, "
<< KernelDescs.size() << ");\n";
O << " }\n";
O << "};\n";
O << "static GlobalMapUpdater updater;\n";
O << "} // namespace detail\n";
O << "} // namespace _V1\n";
O << "} // namespace sycl\n";
O << "} // namespace\n";
}
}

Expand Down
11 changes: 5 additions & 6 deletions clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1087,15 +1087,14 @@ namespace Testing::Tests {

// CHECK: #include <sycl/kernel_bundle.hpp>
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: inline namespace _V1 {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: namespace {
// CHECK-NEXT: struct GlobalMapUpdater {
// CHECK-NEXT: GlobalMapUpdater() {
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
// CHECK-NEXT: }
// CHECK-NEXT: ~GlobalMapUpdater() {
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 28);
// CHECK-NEXT: }
// CHECK-NEXT: };
// CHECK-NEXT: static GlobalMapUpdater updater;
// CHECK-NEXT: } // namespace detail
// CHECK-NEXT: } // namespace _V1
// CHECK-NEXT: } // namespace sycl
// CHECK-NEXT: }
11 changes: 5 additions & 6 deletions clang/test/CodeGenSYCL/free_function_int_header.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1567,15 +1567,14 @@ void ff_24(int arg) {

// CHECK: #include <sycl/kernel_bundle.hpp>
// CHECK-NEXT: #include <sycl/detail/kernel_global_info.hpp>
// CHECK-NEXT: namespace sycl {
// CHECK-NEXT: inline namespace _V1 {
// CHECK-NEXT: namespace detail {
// CHECK-NEXT: namespace {
// CHECK-NEXT: struct GlobalMapUpdater {
// CHECK-NEXT: GlobalMapUpdater() {
// CHECK-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
// CHECK-NEXT: }
// CHECK-NEXT: ~GlobalMapUpdater() {
// CHECK-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 33);
// CHECK-NEXT: }
// CHECK-NEXT: };
// CHECK-NEXT: static GlobalMapUpdater updater;
// CHECK-NEXT: } // namespace detail
// CHECK-NEXT: } // namespace _V1
// CHECK-NEXT: } // namespace sycl
// CHECK-NEXT: }
11 changes: 5 additions & 6 deletions clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -75,15 +75,14 @@ int main(){

// CHECK-NORTC: #include <sycl/kernel_bundle.hpp>
// CHECK-NORTC-NEXT: #include <sycl/detail/kernel_global_info.hpp>
// CHECK-NORTC-NEXT: namespace sycl {
// CHECK-NORTC-NEXT: inline namespace _V1 {
// CHECK-NORTC-NEXT: namespace detail {
// CHECK-NORTC-NEXT: namespace {
// CHECK-NORTC-NEXT: struct GlobalMapUpdater {
// CHECK-NORTC-NEXT: GlobalMapUpdater() {
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::add(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
// CHECK-NORTC-NEXT: }
// CHECK-NORTC-NEXT: ~GlobalMapUpdater() {
// CHECK-NORTC-NEXT: sycl::detail::free_function_info_map::remove(sycl::detail::kernel_names, sycl::detail::kernel_args_sizes, 3);
// CHECK-NORTC-NEXT: }
// CHECK-NORTC-NEXT: };
// CHECK-NORTC-NEXT: static GlobalMapUpdater updater;
// CHECK-NORTC-NEXT: } // namespace detail
// CHECK-NORTC-NEXT: } // namespace _V1
// CHECK-NORTC-NEXT: } // namespace sycl
// CHECK-NORTC-NEXT: }
9 changes: 3 additions & 6 deletions clang/test/SemaSYCL/Inputs/sycl/detail/kernel_global_info.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,14 +10,11 @@

#include <sycl/detail/export.hpp>

namespace sycl {
inline namespace _V1 {
namespace detail {
namespace {
namespace free_function_info_map {

__SYCL_EXPORT void add(const void *DeviceGlobalPtr, const char *UniqueId);
__SYCL_EXPORT void remove(const void *DeviceGlobalPtr, const char *UniqueId);
Copy link
Contributor

@premanandrao premanandrao Oct 31, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Were the changes in this file needed? (I understand it is the mock header, but I don't see this tested in SemaSYCL anywhere.)

Copy link
Contributor Author

@uditagarwal97 uditagarwal97 Oct 31, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@dklochkov-emb since this file was added in https://github.com/intel/llvm/pull/19517/files#diff-dfac134a49adec0ea0525b0941ce5b5e8b51c64e0692aae6c8f0105b9936106e, do you know how is it being used?

Copy link
Contributor

@dklochkov-emb dklochkov-emb Oct 31, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

as far as I know it is needed for SemaSYCL tests, i.e. it mocks real header. For example, test file
https://github.com/intel/llvm/blob/sycl/clang/test/SemaSYCL/free_function_kernel_params.cpp

contains -internal-isystem %S/Inputs. In my understanding, this file will be visible to include as the system file after that


} // namespace free_function_info_map
} // namespace detail
} // namespace _V1
} // namespace sycl
}
2 changes: 2 additions & 0 deletions sycl/include/sycl/detail/kernel_global_info.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,8 @@ namespace free_function_info_map {
__SYCL_EXPORT void add(const char *const *UniqueId,
const unsigned *DeviceGlobalPtr, unsigned Size);

__SYCL_EXPORT void remove(const char *const *UniqueId,
const unsigned *DeviceGlobalPtr, unsigned Size);
} // namespace free_function_info_map
} // namespace detail
} // namespace _V1
Expand Down
11 changes: 11 additions & 0 deletions sycl/source/detail/kernel_global_info.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,17 @@ __SYCL_EXPORT void add(const char *const *FreeFunctionNames,
std::move(GlobalInfoToCopy));
}

__SYCL_EXPORT void remove(const char *const *FreeFunctionNames,
const unsigned *FreeFunctionNumArgs, unsigned Size) {
std::unordered_map<std::string_view, unsigned> GlobalInfoToCopy;
for (size_t i = 0; i < Size; ++i) {
GlobalInfoToCopy[std::string_view{FreeFunctionNames[i]}] =
FreeFunctionNumArgs[i];
}
detail::ProgramManager::getInstance().unRegisterKernelGlobalInfo(
std::move(GlobalInfoToCopy));
}

} // namespace detail::free_function_info_map
} // namespace _V1
} // namespace sycl
14 changes: 12 additions & 2 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 @@ -2418,12 +2418,22 @@ void ProgramManager::registerKernelGlobalInfo(
if (m_FreeFunctionKernelGlobalInfo.empty())
m_FreeFunctionKernelGlobalInfo = std::move(GlobalInfoToCopy);
else {
for (auto &GlobalInfo : GlobalInfoToCopy) {
for (auto &GlobalInfo : GlobalInfoToCopy)
m_FreeFunctionKernelGlobalInfo.insert(GlobalInfo);
}
}
}

// Remove entries from m_FreeFunctionKernelGlobalInfo that matches
// the ones in GlobalInfoToCopy. This function is called when a shared
// library consisting of SYCL kernels is unloaded.
void ProgramManager::unRegisterKernelGlobalInfo(
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy) {
std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);

for (const auto &GlobalInfo : GlobalInfoToCopy)
m_FreeFunctionKernelGlobalInfo.erase(GlobalInfo.first);
}

std::optional<unsigned>
ProgramManager::getKernelGlobalInfoDesc(const char *UniqueId) {
std::lock_guard<std::mutex> Guard(MNativeProgramsMutex);
Expand Down
6 changes: 6 additions & 0 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 @@ -259,6 +259,12 @@ class ProgramManager {
void registerKernelGlobalInfo(
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);

// The function removes kernel global descriptors from the
// kernel global map when a shared library consisting of SYCL kernels
// is unloaded.
void unRegisterKernelGlobalInfo(
std::unordered_map<std::string_view, unsigned> &&GlobalInfoToCopy);

// The function returns a pointer to the kernel global desc identified by
// the unique ID from the kernel global map.
std::optional<unsigned> getKernelGlobalInfoDesc(const char *UniqueId);
Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3344,6 +3344,7 @@ _ZN4sycl3_V16detail21LocalAccessorBaseHostC1ENS0_5rangeILi3EEEiiRKNS0_13property
_ZN4sycl3_V16detail21LocalAccessorBaseHostC2ENS0_5rangeILi3EEEiiRKNS0_13property_listE
_ZN4sycl3_V16detail22addHostAccessorAndWaitEPNS1_16AccessorImplHostE
_ZN4sycl3_V16detail22free_function_info_map3addEPKPKcPKjj
_ZN4sycl3_V16detail22free_function_info_map6removeEPKPKcPKjj
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateERKSt8functionIFbRKSt10shared_ptrINS1_17device_image_implEEEE
_ZN4sycl3_V16detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKNS0_4spanIcLm18446744073709551615EEENS0_12bundle_stateE
Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_windows.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4393,6 +4393,7 @@
?release_external_semaphore@experimental@oneapi@ext@_V1@sycl@@YAXUexternal_semaphore@12345@AEBVqueue@45@@Z
?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVcontext@45@@Z
?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVqueue@45@@Z
?remove@free_function_info_map@detail@_V1@sycl@@YAXPEBQEBDPEBII@Z
?removeDuplicateDevices@detail@_V1@sycl@@YA?BV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV45@@Z
?remquo_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@0PEAH@Z
?remquo_impl@detail@_V1@sycl@@YAMMMPEAH@Z
Expand Down