[llvm] [libsycl] Add liboffload kernel creation (PR #188794)

Sergey Semenov via llvm-commits llvm-commits at lists.llvm.org
Tue Apr 14 06:43:40 PDT 2026


================
@@ -51,26 +65,22 @@ void ProgramAndKernelManager::registerFatBin(__sycl_tgt_bin_desc *FatbinDesc) {
     if (EntriesB == EntriesE)
       continue;
 
-    std::unique_ptr<DeviceImageWrapper> NewImageWrapper =
-        std::make_unique<DeviceImageWrapper>(RawDeviceImage);
+    std::unique_ptr<DeviceImageManager> NewImageWrapper =
+        std::make_unique<DeviceImageManager>(RawDeviceImage);
 
     for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
       auto Name = EntriesIt->SymbolName;
-      auto KernelIDIt = MKernelNameToID.find(Name);
-      if (KernelIDIt == MKernelNameToID.end()) {
-        sycl::kernel_id KernelID =
-            detail::createSyclObjFromImpl<sycl::kernel_id>(
-                std::make_shared<detail::KernelIdImpl>(Name));
-        KernelIDIt = MKernelNameToID.insert(
-            MKernelNameToID.end(),
-            std::make_pair(std::string_view(Name), KernelID));
-      }
 
-      MKernelIDToDevImageJIT.insert(
-          std::make_pair(KernelIDIt->second, NewImageWrapper.get()));
+      auto It = MDeviceKernelInfoMap.find(std::string_view(Name));
+      if (It == MDeviceKernelInfoMap.end()) {
+
+        std::ignore = MDeviceKernelInfoMap.emplace_hint(
+            It, std::piecewise_construct, std::forward_as_tuple(Name),
+            std::forward_as_tuple(Name, *NewImageWrapper));
+      }
----------------
sergey-semenov wrote:

```suggestion
     [[maybe_unused]] auto Result = MDeviceKernelInfoMap.emplace(std::piecewise_construct, std::forward_as_tuple(Name), std::forward_as_tuple(Name, *NewImageWrapper));
     assert(Result.second && "Kernel name found in multiple images");
```
Currently, the code assumes a single image per kernel, so I think we should make that assertion explicit.

https://github.com/llvm/llvm-project/pull/188794


More information about the llvm-commits mailing list