From 88e13ae6f59886a276147bca2a6e476dd928a3a0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 23 Jul 2025 12:48:36 -0700 Subject: [PATCH 1/6] bad boy does bad things Signed-off-by: Chris Perkins --- .../clang-offload-wrapper/ClangOffloadWrapper.cpp | 10 ++++++++++ sycl/source/detail/program_manager/program_manager.cpp | 2 ++ 2 files changed, 12 insertions(+) diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index ee913fc8d0ee..93ab28c529f0 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 @@ -1268,6 +1271,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 +1288,11 @@ class BinaryWrapper { : "__tgt_unregister_lib", UnRegFuncTy); + + // CP -- this is output when COMPILING the application. + std::cout << "createUnregisterFunction. SYCL?: " << (Kind == OffloadKind::SYCL) << std::endl; + + // Construct function body IRBuilder<> Builder(BasicBlock::Create(C, "entry", Func)); Builder.CreateCall(UnRegFuncC, BinDesc); diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 485b94e36f65..a16ac6d6248c 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -3844,11 +3844,13 @@ 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()) From 7631c5b49191ce69ec7a27d13d6e1b964da50dd5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 23 Jul 2025 15:01:51 -0700 Subject: [PATCH 2/6] sycloffloadwrapper Signed-off-by: Chris Perkins --- clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp | 2 +- llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp | 8 ++++++++ 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index 93ab28c529f0..847e1e418bfe 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -1290,7 +1290,7 @@ class BinaryWrapper { // CP -- this is output when COMPILING the application. - std::cout << "createUnregisterFunction. SYCL?: " << (Kind == OffloadKind::SYCL) << std::endl; + std::cout << "ClangOffloadWrapper createUnregisterFunction. SYCL?: " << (Kind == OffloadKind::SYCL) << std::endl; // Construct function body diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 3d227d0c2e05..4ac580fc53ca 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -38,6 +38,8 @@ #include #include #include +// CP +#include using namespace llvm; using namespace llvm::offloading; @@ -704,6 +706,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 +731,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); From f94879eac75da434c1b8eb614d38d9e1b5b098e3 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 25 Jul 2025 10:55:57 -0700 Subject: [PATCH 3/6] register atexit() on Win Signed-off-by: Chris Perkins --- .../ClangOffloadWrapper.cpp | 64 ++++++++++++++++++- 1 file changed, 62 insertions(+), 2 deletions(-) diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index 847e1e418bfe..2802e38de6dd 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -1257,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); @@ -1303,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 = "") @@ -1380,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; From a92b1fa312b27445a16f1868a368649abfaef9ff Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 25 Jul 2025 11:01:02 -0700 Subject: [PATCH 4/6] clang-format is having a stroke? Signed-off-by: Chris Perkins --- clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index 2802e38de6dd..f4ac50b183a1 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -1326,7 +1326,7 @@ class BinaryWrapper { UnregBuilder.CreateCall(UnregTargetC, BinDesc); UnregBuilder.CreateRetVoid(); - .auto *RegFuncTy = + auto *RegFuncTy = FunctionType::get(Type::getVoidTy(C), /*isVarArg*/ false); auto *RegFunc = Function::Create(RegFuncTy, GlobalValue::InternalLinkage, "sycl.descriptor_reg", &M); From 24a77ecaa94b04e0baf98d201948f7dcfe534e15 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 25 Jul 2025 11:43:49 -0700 Subject: [PATCH 5/6] unregister with atexit for Win in SYCLOffloadWrapper Signed-off-by: Chris Perkins --- .../Offloading/SYCLOffloadWrapper.cpp | 59 ++++++++++++++++++- 1 file changed, 57 insertions(+), 2 deletions(-) diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp index 4ac580fc53ca..993486aec12a 100644 --- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp +++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp @@ -34,6 +34,7 @@ #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 @@ -742,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 @@ -755,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(); } From 9c2fb3ae46b0ff10940b1593cf42e6926e010f20 Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Fri, 25 Jul 2025 12:28:10 -0700 Subject: [PATCH 6/6] program_manager.cpp sycl_unregister_lib() reenabled release of images --- sycl/source/detail/program_manager/program_manager.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index a16ac6d6248c..6a3ec7628182 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -3855,6 +3855,7 @@ extern "C" void __sycl_unregister_lib(sycl_device_binaries desc) { #ifndef _WIN32 if (!sycl::detail::GlobalHandler::instance().isOkToDefer()) return; - sycl::detail::ProgramManager::getInstance().removeImages(desc); #endif + sycl::detail::ProgramManager::getInstance().removeImages(desc); + }