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

[Sergey Semenov via llvm-commits]

================
@@ -97,49 +107,58 @@ void ProgramAndKernelManager::unregisterFatBin(
       continue;
 
     for (auto EntriesIt = EntriesB; EntriesIt != EntriesE; ++EntriesIt) {
-      if (auto KernelIDIt = MKernelNameToID.find(EntriesIt->SymbolName);
-          KernelIDIt != MKernelNameToID.end()) {
-        MKernelIDToDevImageJIT.erase(KernelIDIt->second);
-        MKernelNameToID.erase(KernelIDIt);
+      if (auto KernelIt = MDeviceKernelInfoMap.find(EntriesIt->SymbolName);
+          KernelIt != MDeviceKernelInfoMap.end()) {
+        // Programs are attached to image and will be released with image
+        // destruction. Clear only kernel specific data by destroying its kernel
+        // info object.
+        MDeviceKernelInfoMap.erase(KernelIt);
       }
     }
 
-    MDeviceImageWrappers.erase(DevImageIt);
+    MDeviceImageManagers.erase(DevImageIt);
   }
 }
 
-static bool isImageTargetCompatible(const DeviceImageWrapper &Image,
-                                    const DeviceImpl &Device) {
+static bool isImageCompatible(const DeviceImageManager &Image,
+                              const DeviceImpl &Device) {
   sycl::backend BE = Device.getBackend();
   const char *Target = Image.getRawData().TripleString;
 
-  return (strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
-         (BE == sycl::backend::level_zero);
+  if (!(strcmp(Target, DeviceBinaryTripleSPIRV64) == 0) &&
+      (BE == sycl::backend::level_zero))
----------------
sergey-semenov wrote:

```suggestion
  if (!(strcmp(Target, DeviceBinaryTripleSPIRV64) == 0 &&
      BE == sycl::backend::level_zero))
```

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