-
Notifications
You must be signed in to change notification settings - Fork 790
[SYCL] fix for __sycl_unregister_lib() on Windows and tests #19633
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from 16 commits
c480714
8c1905e
8688a67
6742db0
154eaf0
44bef0d
d8bc95a
496c746
d1c48c8
c415c15
455bdf8
e2458f3
c6afa75
843a961
a6ef7e0
8c7d1d5
2a193e0
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
|
@@ -1293,6 +1293,49 @@ 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); | ||
|
||
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); | ||
|
||
IRBuilder<> RegBuilder(BasicBlock::Create(C, "entry", RegFunc)); | ||
RegBuilder.CreateCall(RegTargetC, BinDesc); | ||
RegBuilder.CreateCall(AtExitC, UnregFunc); | ||
RegBuilder.CreateRetVoid(); | ||
|
||
// Add this function to global destructors. | ||
appendToGlobalCtors(M, RegFunc, /*Priority*/ 1); | ||
YuriPlyakhin marked this conversation as resolved.
Show resolved
Hide resolved
|
||
} | ||
|
||
public: | ||
BinaryWrapper(StringRef Target, StringRef ToolName, | ||
StringRef SymPropBCFiles = "") | ||
|
@@ -1370,8 +1413,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); | ||
} | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. It doesn't necessarily have to be part of this patch, but could you please document this approach in a design document? |
||
} | ||
} | ||
return &M; | ||
|
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -1,3 +1,6 @@ | ||
// UNSUPPORTED: hip | ||
// UNSUPPORTED-TRACKER: CMPLRLLVM-69478 | ||
|
||
// RUN: %{build} -o %t.out | ||
// RUN: %{run} %t.out | ||
|
||
|
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,38 @@ | ||
#include <sycl/detail/core.hpp> | ||
|
||
#if defined(_WIN32) | ||
#define API_EXPORT __declspec(dllexport) | ||
#else | ||
#define API_EXPORT | ||
#endif | ||
|
||
#ifndef INC | ||
#define INC 1 | ||
#endif | ||
|
||
#ifndef CLASSNAME | ||
#define CLASSNAME same | ||
#endif | ||
|
||
#ifdef WITH_DEVICE_GLOBALS | ||
// Using device globals within the shared libraries only | ||
// works if the names do not collide. Note that we cannot | ||
// load a library multiple times if it has a device global. | ||
#define CONCAT_HELPER(a, b) a##b | ||
#define CONCAT(a, b) CONCAT_HELPER(a, b) | ||
|
||
using SomeProperties = decltype(sycl::ext::oneapi::experimental::properties{}); | ||
sycl::ext::oneapi::experimental::device_global<int, SomeProperties> | ||
CONCAT(DGVar, CLASSNAME) __attribute__((visibility("default"))); | ||
|
||
#endif // WITH_DEVICE_GLOBALS | ||
|
||
extern "C" API_EXPORT void performIncrementation(sycl::queue &q, | ||
sycl::buffer<int, 1> &buf) { | ||
sycl::range<1> r = buf.get_range(); | ||
q.submit([&](sycl::handler &cgh) { | ||
auto acc = buf.get_access<sycl::access::mode::write>(cgh); | ||
cgh.parallel_for<class CLASSNAME>( | ||
r, [=](sycl::id<1> idx) { acc[idx] += INC; }); | ||
}); | ||
} |
Uh oh!
There was an error while loading. Please reload this page.