diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index ee913fc8d0eeb..f4ac50b183a1b 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -67,6 +67,9 @@ #include #include +// CP +#include + // For device image compression. #include @@ -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; diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 3d227d0c2e050..993486aec12ad 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/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 #include #include +// CP +#include 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(); } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 485b94e36f658..6a3ec7628182f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/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); + }