Cperkins investigate unregister lib

clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp

@@ -67,6 +67,9 @@
 #include <string>
 #include <tuple>
 
+// CP
+#include <iostream>
+
 // For device image compression.
 #include <llvm/Support/Compression.h>
 
@@ -1254,6 +1257,10 @@ class BinaryWrapper {
                                                         : "__tgt_register_lib",
                               RegFuncTy);
 
+    // CP  -- this is output when COMPILING the application.
+    std::cout << "ClangOffloadWrapper createRegisterFunction.  SYCL?: "
+              << (Kind == OffloadKind::SYCL) << std::endl;
+
     // Construct function body
     IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func));
     Builder.CreateCall(RegFuncC, BinDesc);
@@ -1268,6 +1275,8 @@ class BinaryWrapper {
     appendToGlobalCtors(M, Func, /*Priority*/ 1);
   }
 
+
+
   void createUnregisterFunction(OffloadKind Kind, GlobalVariable *BinDesc) {
     auto *FuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
     auto *Func =
@@ -1283,6 +1292,11 @@ class BinaryWrapper {
                                   : "__tgt_unregister_lib",
         UnRegFuncTy);
 
+
+    // CP  -- this is output when COMPILING the application. 
+    std::cout << "ClangOffloadWrapper createUnregisterFunction.  SYCL?: " << (Kind == OffloadKind::SYCL) << std::endl;
+
+
     // Construct function body
     IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func));
     Builder.CreateCall(UnRegFuncC, BinDesc);
@@ -1293,6 +1307,57 @@ class BinaryWrapper {
     appendToGlobalDtors(M, Func, /*Priority*/ 1);
   }
 
+  void createSyclRegisterWithAtexitUnregister(GlobalVariable *BinDesc) {
+    auto *UnregFuncTy =
+        FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *UnregFunc =
+        Function::Create(UnregFuncTy, GlobalValue::InternalLinkage,
+                         "sycl.descriptor_unreg.atexit", &M);
+    UnregFunc->setSection(".text.startup");
+
+    // Declaration for __sycl_unregister_lib(void*).
+    auto *UnregTargetTy =
+        FunctionType::get(Type::getVoidTy(C), getPtrTy(), false);
+    FunctionCallee UnregTargetC =
+        M.getOrInsertFunction("__sycl_unregister_lib", UnregTargetTy);
+
+    // Body of the unregister wrapper.
+    IRBuilder<> UnregBuilder(BasicBlock::Create(C, "entry", UnregFunc));
+    UnregBuilder.CreateCall(UnregTargetC, BinDesc);
+    UnregBuilder.CreateRetVoid();
+
+    auto *RegFuncTy =
+        FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *RegFunc = Function::Create(RegFuncTy, GlobalValue::InternalLinkage,
+                                     "sycl.descriptor_reg", &M);
+    RegFunc->setSection(".text.startup");
+
+    auto *RegTargetTy =
+        FunctionType::get(Type::getVoidTy(C), getPtrTy(), false);
+    FunctionCallee RegTargetC =
+        M.getOrInsertFunction("__sycl_register_lib", RegTargetTy);
+
+    // `atexit` takes a `void(*)()` function pointer. In LLVM IR, this is
+    // typically represented as `i32 (ptr)`.
+    FunctionType *AtExitTy =
+        FunctionType::get(Type::getInt32Ty(C), getPtrTy(), false);
+    FunctionCallee AtExitC = M.getOrInsertFunction("atexit", AtExitTy);
+
+    // CP  -- this is output when COMPILING the application.
+    std::cout << "ClangOffloadWrapper createSyclRegisterWithAtexitUnregister. "
+              << std::endl;
+
+    // Body of the register function.
+    IRBuilder<> RegBuilder(BasicBlock::Create(C, "entry", RegFunc));
+    RegBuilder.CreateCall(RegTargetC, BinDesc);
+    RegBuilder.CreateCall(AtExitC, UnregFunc);
+    RegBuilder.CreateRetVoid();
+
+    // Add this function to global destructors.
+    // Match priority of __tgt_register_lib
+    appendToGlobalCtors(M, RegFunc, /*Priority*/ 1);
+  }
+
 public:
   BinaryWrapper(StringRef Target, StringRef ToolName,
                 StringRef SymPropBCFiles = "")
@@ -1370,8 +1435,13 @@ class BinaryWrapper {
 
       if (EmitRegFuncs) {
         GlobalVariable *Desc = *DescOrErr;
-        createRegisterFunction(Kind, Desc);
-        createUnregisterFunction(Kind, Desc);
+        if (Kind == OffloadKind::SYCL &&
+            Triple(M.getTargetTriple()).isOSWindows()) {
+          createSyclRegisterWithAtexitUnregister(Desc);
+        } else {
+          createRegisterFunction(Kind, Desc);
+          createUnregisterFunction(Kind, Desc);
+        }
       }
     }
     return &M;

llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp

@@ -34,10 +34,13 @@
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/LineIterator.h"
 #include "llvm/Support/PropertySetIO.h"
+#include "llvm/TargetParser/Triple.h"
 #include "llvm/Transforms/Utils/ModuleUtils.h"
 #include <memory>
 #include <string>
 #include <utility>
+// CP
+#include <iostream>
 
 using namespace llvm;
 using namespace llvm::offloading;
@@ -704,6 +707,9 @@ struct Wrapper {
     FunctionCallee RegFuncC =
         M.getOrInsertFunction("__sycl_register_lib", RegFuncTy);
 
+    // CP  -- 
+    std::cout << "SYCLOffloadWrapper createRegisterFatbinFunction. " << std::endl;
+
     // Construct function body
     IRBuilder Builder(BasicBlock::Create(C, "entry", Func));
     Builder.CreateCall(RegFuncC, FatbinDesc);
@@ -726,6 +732,9 @@ struct Wrapper {
     FunctionCallee UnRegFuncC =
         M.getOrInsertFunction("__sycl_unregister_lib", UnRegFuncTy);
 
+    // CP  -- 
+    std::cout << "SYCLOffloadWrapper createUnregisterFunction. " << std::endl;
+
     // Construct function body
     IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func));
     Builder.CreateCall(UnRegFuncC, FatbinDesc);
@@ -734,6 +743,56 @@ struct Wrapper {
     // Add this function to global destructors.
     appendToGlobalDtors(M, Func, /*Priority*/ 1);
   }
+
+  void createSyclRegisterWithAtexitUnregister(GlobalVariable *FatbinDesc) {
+    auto *UnregFuncTy =
+        FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *UnregFunc =
+        Function::Create(UnregFuncTy, GlobalValue::InternalLinkage,
+                         "sycl.descriptor_unreg.atexit", &M);
+    UnregFunc->setSection(".text.startup");
+
+    // Declaration for __sycl_unregister_lib(void*).
+    auto *UnregTargetTy =
+        FunctionType::get(Type::getVoidTy(C), PointerType::getUnqual(C), false);
+    FunctionCallee UnregTargetC =
+        M.getOrInsertFunction("__sycl_unregister_lib", UnregTargetTy);
+
+    // Body of the unregister wrapper.
+    IRBuilder<> UnregBuilder(BasicBlock::Create(C, "entry", UnregFunc));
+    UnregBuilder.CreateCall(UnregTargetC, FatbinDesc);
+    UnregBuilder.CreateRetVoid();
+
+    auto *RegFuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false);
+    auto *RegFunc = Function::Create(RegFuncTy, GlobalValue::InternalLinkage,
+                                     "sycl.descriptor_reg", &M);
+    RegFunc->setSection(".text.startup");
+
+    auto *RegTargetTy =
+        FunctionType::get(Type::getVoidTy(C), PointerType::getUnqual(C), false);
+    FunctionCallee RegTargetC =
+        M.getOrInsertFunction("__sycl_register_lib", RegTargetTy);
+
+    // `atexit` takes a `void(*)()` function pointer. In LLVM IR, this is
+    // typically represented as `i32 (ptr)`.
+    FunctionType *AtExitTy = FunctionType::get(
+        Type::getInt32Ty(C), PointerType::getUnqual(C), false);
+    FunctionCallee AtExitC = M.getOrInsertFunction("atexit", AtExitTy);
+
+    // CP  --
+    std::cout << "SYCLOffloadWrapper createSyclRegisterWithAtexitUnregister. "
+              << std::endl;
+
+    // Body of the register function.
+    IRBuilder<> RegBuilder(BasicBlock::Create(C, "entry", RegFunc));
+    RegBuilder.CreateCall(RegTargetC, FatbinDesc);
+    RegBuilder.CreateCall(AtExitC, UnregFunc);
+    RegBuilder.CreateRetVoid();
+
+    // Add to global constructors.
+    appendToGlobalCtors(M, RegFunc, /*Priority*/ 1);
+  }
+
 }; // end of Wrapper
 
 } // anonymous namespace
@@ -747,7 +806,11 @@ Error llvm::offloading::wrapSYCLBinaries(llvm::Module &M,
     return createStringError(inconvertibleErrorCode(),
                              "No binary descriptors created.");
 
-  W.createRegisterFatbinFunction(Desc);
-  W.createUnregisterFunction(Desc);
+  if (Triple(M.getTargetTriple()).isOSWindows()) {
+    W.createSyclRegisterWithAtexitUnregister(Desc);
+  } else {
+    W.createRegisterFatbinFunction(Desc);
+    W.createUnregisterFunction(Desc);
+  }
   return Error::success();
 }

sycl/source/detail/program_manager/program_manager.cpp

@@ -3844,15 +3844,18 @@ bool doesImageTargetMatchDevice(const RTDeviceBinaryImage &Img,
 } // namespace sycl
 
 extern "C" void __sycl_register_lib(sycl_device_binaries desc) {
+  std::cout << "__sycl_register_lib" << std::endl;
   sycl::detail::ProgramManager::getInstance().addImages(desc);
 }
 
 // Executed as a part of current module's (.exe, .dll) static initialization
 extern "C" void __sycl_unregister_lib(sycl_device_binaries desc) {
+   std::cout << "__sycl_unregister_lib" << std::endl;
   // Partial cleanup is not necessary at shutdown
 #ifndef _WIN32
   if (!sycl::detail::GlobalHandler::instance().isOkToDefer())
     return;
-  sycl::detail::ProgramManager::getInstance().removeImages(desc);
 #endif
+  sycl::detail::ProgramManager::getInstance().removeImages(desc);
+
 }