[sycl-post-link] Do not drop SYCL_EXTERNAL functions #3793

DenisBakhvalov · 2021-05-20T23:50:27Z

There are a few use cases where unreferenced device functions should go through sycl-post-link. Such functions may be referenced later, for example on a SPIRV level.
Dynamic linking of device binary images is one such case. I.e. when the kernel will have an import list with functionality implemented in a dynamic library. DPCPP RT will resolve imported functions, which could be implemented as SYCL_EXTERNAL functions. Thus we need to allow SYCL_EXTERNAL to go through sycl-post-link to make it work.
ESIMD/ISPC interoperability requested by @aneshlya is another motivation for this functionality.

DenisBakhvalov · 2021-05-20T23:56:49Z

I'm not 100% sure it always safe to allow unreferenced functions to go through sycl-post-link. One case I have in mind is when we split per-kernel. In such a scenario, sycl-post-link will place unreferenced functions in a separate module. So far, I believe the DPCPP runtime expects to see kernels in each SPIRV module, but there might be none. I know that RT claims binary images based on symbols they contain, but still, I'm not 100% sure about it. @romanovvlad, do you have any comments here?

DenisBakhvalov · 2021-05-21T03:59:38Z

I will address CI failures shortly.
Does anyone have a strong opinion on whether I should put this functionality under a dedicated option?

romanovvlad · 2021-05-21T19:42:47Z

I'm not 100% sure it always safe to allow unreferenced functions to go through sycl-post-link. One case I have in mind is when we split per-kernel. In such a scenario, sycl-post-link will place unreferenced functions in a separate module. So far, I believe the DPCPP runtime expects to see kernels in each SPIRV module, but there might be none. I know that RT claims binary images based on symbols they contain, but still, I'm not 100% sure about it. @romanovvlad, do you have any comments here?

Currently SYCL RT will assume a device images has all kernels if the list of kernel in the device image is zero.
But it uses such a device image if only a kernel is not found in images where list of kernels is not zero.

llvm/test/tools/sycl-post-link/unreferenced_funcs/split-global.ll

llvm/test/tools/sycl-post-link/unreferenced_funcs/split-per-source.ll

llvm/tools/sycl-post-link/sycl-post-link.cpp

AlexeySachkov · 2021-05-24T10:52:25Z

llvm/tools/sycl-post-link/sycl-post-link.cpp

+
+  std::vector<llvm::Function *> UnreferencedFuncs;
+  for (auto &F : M.functions()) {
+    if (!F.isDeclaration() && !ReferencedFuncs.count(&F))


What if we have an internal (I mean linkage type) unreferenced function here? Could it cause any conflicts later?

I think we can ignore unreferenced internal functions, just like we did before. AFAIK if two LLVM modules defined internal functions with same names, they won't cause conflicts when these two LLVM modules linked. I'm not sure about SPIR-V though, I can't find any analog of internal linkage type in SPIR-V spec, but the translator seems to support some Internal linkage type. (@AlexeySotkin , any idea why?).

There might be cases when we want to preserve unreferenced functions in SPIR-V, e.g. those which are mentioned in @llvm.used variable.
Whether they will conflict or not depends on linkage type.
internal linkage type is the default in SPIR-V, i.e. when global variable or function doesn't have LinkageType decoration it has internal linkage.

Fznamznon · 2021-05-24T13:38:37Z

The idea behind dynamic linking support was to consider SYCL_EXTERNAL functions as entry points. I.e. SYCL_EXTERNAL functions would be processed as kernels. Kernels don't have references as well, so the algorithm will require smaller change than you have now. The rest non-SYCL_EXTERNAL functions are still considered as internal functions in device code, so they can be removed from the module if not referenced.
Support for dynamic linking would require adding sycl-module-id attribute to SYCL_EXTERNAL functions. Applying of this idea can also help to fix some lit fails because the reason is missing sycl-module-id attribute.

DenisBakhvalov · 2021-05-25T00:43:42Z

The idea behind dynamic linking support was to consider SYCL_EXTERNAL functions as entry points. I.e. SYCL_EXTERNAL functions would be processed as kernels.

That's exactly what I'm trying to do in this patch.

Kernels don't have references as well, so the algorithm will require smaller change than you have now. The rest non-SYCL_EXTERNAL functions are still considered as internal functions in device code, so they can be removed from the module if not referenced.

I thought that non-SYCL_EXTERNAL functions are filtered by FE, no?

Support for dynamic linking would require adding sycl-module-id attribute to SYCL_EXTERNAL functions. Applying of this idea can also help to fix some lit fails because the reason is missing sycl-module-id attribute.

Do I understand correctly that your idea is to distinguish roots (entry points) by the presence of sycl-module-id attributes, correct? So, we essentially can replace:
if (F.getCallingConv() == CallingConv::SPIR_KERNEL)
with
if (F->hasFnAttribute(ATTR_SYCL_MODULE_ID))
right?

Also, there is a comment in clang/lib/CodeGen/CodeGenFunction.cpp that says:

// TODO Module identifier is not reliable for this purpose since two modules
// can have the same ID, needs improvement
if (getLangOpts().SYCLIsDevice)
  Fn->addFnAttr("sycl-module-id", Fn->getParent()->getModuleIdentifier());

Fznamznon · 2021-05-25T07:54:23Z

I thought that non-SYCL_EXTERNAL functions are filtered by FE, no?

Yes, FE ignores functions that is not referenced from kernels/SYCL_EXTERNALs (I'll call it device context), but with early optimizations there is the case when function was originally referenced from device context, but use of it was optimized away. In this case we get unreferenced non-SYCL_EXTERNAL function in the module on sycl-post-link stage. I guess this is happening in failed tests.

Do I understand correctly that your idea is to distinguish roots (entry points) by the presence of sycl-module-id attributes, correct?

Yes, that is right.

Also, there is a comment in clang/lib/CodeGen/CodeGenFunction.cpp that says:

Yes, there is still a room for improvement, I don't remember if anyone ever returned to it. But it shouldn't block us from re-using the attribute for SYCL_EXTERNAL functions. The only thing that might be affected by collision in sycl-module-id attribute values is per-source device code split distribution.

DenisBakhvalov · 2021-05-25T22:56:02Z

Thanks, @Fznamznon,
I found that FE already generates the sycl-module-id attribute for SYCL_EXTERNAL functions. I updated the PR according to the previous comments.

Fznamznon · 2021-05-27T12:57:12Z

llvm/tools/sycl-post-link/sycl-post-link.cpp

+      "  any kernel are dropped from the resulting module(s), except from\n"
+      "  SYCL_EXTERNAL functions.\n"


How about "Functions unreachable from any kernel or SYCL_EXTERNAL function are dropped from the resulting module(s)."?

Fznamznon · 2021-05-27T13:15:48Z

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source.ll

+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir64-unknown-linux-sycldevice"
+
+define dso_local spir_func void @externalDeviceFunc() #0 {


Can we add a test where some SYCL_EXTERNAL function foo is defined in source "a.cpp" but it is also referenced by a kernel defined in source "b.cpp"? In this case design of dynamic code linking says that definition of function foo should be present in both resulting modules when per-source split is requested.

DenisBakhvalov · 2021-05-27T23:30:44Z

This PR is supposed to fix #3239.

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-and-lower-esimd.ll

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-global.ll

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-kernel.ll

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source1.ll

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-per-source2.ll

llvm/tools/sycl-post-link/sycl-post-link.cpp

This is an NFC patch, a preparation for further changes.

…ribute

llvm/test/tools/sycl-post-link/sycl-external-funcs/drop-unref-funcs.ll

DenisBakhvalov · 2021-06-03T00:29:28Z

Looks like CUDA AOT mode generates call without SPIR_KERNEL or SPIR_FUNC and still uses sycl-post-link. So, I removed the check for unsupported calling conventions.

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-global.ll

DenisBakhvalov · 2021-06-04T22:30:10Z

I have a new issue that popped up in the CI testing (AOT/multiple-devices.cpp). The problem there is that DPCPP RT does not allow device binary images to have a disjoint set of kernels&functions in them. With the changes in the current patch, it is the case: we compile the same code for 3 different targets in AOT mode: cpu, gpu, and fpga. For cpu and gpu we compile with -mllvm -sycl-opt, while for fpga we compile with -fno-sycl-early-optimizations. This causes SYCL_EXTERNAL size_t __spirv_GlobalInvocationId_x() to be not inlined for FPGA and later recognized as external symbol by sycl-post-link. So, in the end we have something like:

# cpu symbols:
_ZTS10SimpleVaddIiE

# gpu symbols:
_ZTS10SimpleVaddIiE

# fpga symbols:
_ZTS10SimpleVaddIiE
_Z28__spirv_GlobalInvocationId_xv

And this is why DPCPP RT asserts that the set of kernel symbols are disjoint across all the images in the device binary.

I tend to think we should filter SPIRV ("_spirv") and SYCL ("_sycl") built-in functions in sycl-post-link. Such functions should not be exposed to the outside even though they are declared as SYCL_EXTERNAL. So, I think we don't need to add them to the ResKernelModuleMap in sycl-post-link. If we do that, _Z28__spirv_GlobalInvocationId_xv will be gone from the fpga symbols and the problem will be solved.

@Fznamznon, @AlexeySachkov, @kbobrovs any thoughts?

Fznamznon · 2021-06-07T12:53:35Z

The problem there is that DPCPP RT does not allow device binary images to have a disjoint set of kernels&functions in them.

Is there any check for functions? I remember RT has a thing called kernel set id, it is based on kernels presented in the device image. Are you saying with your changes, we add names of SYCL_EXTERNAL functions to entries table (

llvm/sycl/include/CL/sycl/detail/pi.h

Line 728 in a3fc51a

_pi_offload_entry EntriesBegin;

) as well?
The idea behind dynamic linking was to have a separate property set for SYCL_EXTERNAL function names and not add them to entries table, so this thing wouldn't break. However I think we actually could use either only entries table or only property table to save kernel names and SYCL_EXTERNAL names both, but we will need to re-do RT part with kernel set id. If I understand it correct, this kernel set id is actually not a necessity, but some choice made by RT developers some time ago. @s-kanaev , WDYT?

I tend to think we should filter SPIRV ("_spirv") and SYCL ("_sycl") built-in functions in sycl-post-link. Such functions should not be exposed to the outside even though they are declared as SYCL_EXTERNAL.

Not sure about __sycl built-ins, but we should definitely not treat __spirv functions as external symbols, since they disappear after SPIR-V translation. AFAIK some of __sycl built-ins are used for spec constants, @AlexeySachkov , could you please comment?

AlexeySachkov · 2021-06-08T11:50:08Z

Not sure about __sycl built-ins, but we should definitely not treat __spirv functions as external symbols, since they disappear after SPIR-V translation. AFAIK some of __sycl built-ins are used for spec constants, @AlexeySachkov , could you please comment?

We do have a few __sycl_ pseudo-intrinsics, but all of them are replaced with something else by sycl-post-link.

kbobrovs

@DenisBakhvalov, looks like there is a real regression: [SYCL :: AOT/multiple-devices.cpp]

DenisBakhvalov · 2021-06-11T22:53:22Z

@DenisBakhvalov, looks like there is a real regression: [SYCL :: AOT/multiple-devices.cpp]

Right. I didn't have the time to fix it yet. But it's still under my radar. :)

…oints

DenisBakhvalov · 2021-06-15T00:32:10Z

Sorry for the late response.

@Fznamznon ,

Are you saying with your changes, we add names of SYCL_EXTERNAL functions to entries table as well?

Yes, here is the assert that gets triggered:

sycl/source/detail/program_manager/program_manager.cpp:965: void cl::sycl::detail::ProgramManager::addImages(pi_device_binaries): Assertion `KSIdMap[EntriesIt->name] == KSIdIt->second && "Kernel sets are not disjoint"' failed.

we should definitely not treat __spirv functions as external symbols, since they disappear after SPIR-V translation.

So, I went ahead and added this exception.

DenisBakhvalov · 2021-06-16T20:59:53Z

@Fznamznon, @AlexeySachkov, please review.

Fznamznon · 2021-06-18T08:08:50Z

llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll

@@ -21,25 +21,27 @@

 declare dso_local spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)*)

-define dso_local spir_kernel void @KERNEL_AAA() {
+define dso_local spir_kernel void @KERNEL_AAA() #0 {


I'm curious, why do we need to add the attribute to this test. Why It wasn't here before? Is it not possible to process the code without the attribute?

If I understand correctly, this is how we detect whether the function should be considered as an entry point or not

That's correct.

Fznamznon · 2021-06-18T08:10:49Z

llvm/test/tools/sycl-post-link/sym_but_no_split.ll

@@ -6,20 +6,22 @@
 ; RUN: FileCheck %s -input-file=%t.files.table --check-prefixes CHECK-TABLE
 ; RUN: FileCheck %s -input-file=%t.files_0.sym --match-full-lines --check-prefixes CHECK-SYM

-define dso_local spir_kernel void @KERNEL_AAA() {
+define dso_local spir_kernel void @KERNEL_AAA() #0 {


AlexeySachkov · 2021-06-18T12:55:59Z

llvm/test/tools/sycl-post-link/spec-constants/spec_const_and_split.ll

@@ -21,25 +21,27 @@

 declare dso_local spir_func zeroext i1 @_Z33__sycl_getScalarSpecConstantValueIbET_PKc(i8 addrspace(4)*)

-define dso_local spir_kernel void @KERNEL_AAA() {
+define dso_local spir_kernel void @KERNEL_AAA() #0 {


If I understand correctly, this is how we detect whether the function should be considered as an entry point or not

DenisBakhvalov · 2021-06-18T15:33:56Z

@kbobrovs , @mlychkov please review/approve.

DenisBakhvalov · 2021-06-18T17:34:44Z

@pvchupin , @kbobrovs , I have all the approvals, please merge.

DenisBakhvalov requested review from aneshlya, AlexeySachkov and kbobrovs May 20, 2021 23:50

DenisBakhvalov requested a review from mlychkov as a code owner May 20, 2021 23:50

DenisBakhvalov requested a review from romanovvlad May 20, 2021 23:54

bader requested a review from Fznamznon May 21, 2021 08:32

AlexeySachkov reviewed May 24, 2021

View reviewed changes

DenisBakhvalov force-pushed the external_funcs_post_link branch from 210d80b to d76610c Compare May 25, 2021 22:39

DenisBakhvalov changed the title ~~[sycl-post-link] Do not drop unreferenced functions~~ [sycl-post-link] Do not drop SYCL_EXTERNAL functions May 25, 2021

Fznamznon reviewed May 27, 2021

View reviewed changes

kbobrovs reviewed May 27, 2021

View reviewed changes

DenisBakhvalov and others added 11 commits May 27, 2021 17:27

[sycl-post-link] Do not drop unreferenced functions

9e41ec7

This is an NFC patch, a preparation for further changes.

[sycl-post-link] Do not drop unreferenced functions

7294ed5

[sycl-post-link] Support ESIMD lowering of unreferenced functions

966e154

Distinguish module entry points by the presence of sycl-module-id att…

253fc35

…ribute

Update erase_used.ll

15ebc4b

Update split-and-lower-esimd.ll

6cc11f0

Update split-per-kernel.ll

e1a24e7

Update split-per-source.ll

e9d0e27

Fixed help text

0313636

fixed review comments

c59db7d

Fixed code review comments

b3f5600

Fznamznon reviewed May 31, 2021

View reviewed changes

llvm/test/tools/sycl-post-link/sycl-external-funcs/drop-unref-funcs.ll Show resolved Hide resolved

Update drop-unref-funcs.ll

cdffa9f

Fznamznon previously approved these changes Jun 2, 2021

View reviewed changes

Fznamznon requested review from AlexeySachkov and kbobrovs June 2, 2021 18:37

Backed of the check for unsupported calling conventions

1eec0c4

DenisBakhvalov dismissed Fznamznon’s stale review via 1eec0c4 June 3, 2021 00:27

AlexeySachkov reviewed Jun 3, 2021

View reviewed changes

llvm/test/tools/sycl-post-link/sycl-external-funcs/split-global.ll Show resolved Hide resolved

kbobrovs previously approved these changes Jun 11, 2021

View reviewed changes

SPIRV and SYCL builtin functions are not considered as module entry p…

a54f245

…oints

DenisBakhvalov dismissed kbobrovs’s stale review via a54f245 June 15, 2021 00:30

DenisBakhvalov requested review from kbobrovs and Fznamznon June 15, 2021 03:03

Fznamznon reviewed Jun 18, 2021

View reviewed changes

AlexeySachkov approved these changes Jun 18, 2021

View reviewed changes

DenisBakhvalov requested a review from Fznamznon June 18, 2021 15:32

mlychkov approved these changes Jun 18, 2021

View reviewed changes

Fznamznon approved these changes Jun 18, 2021

View reviewed changes

kbobrovs approved these changes Jun 18, 2021

View reviewed changes

pvchupin merged commit 34e1f7c into intel:sycl Jun 18, 2021

Fznamznon mentioned this pull request May 21, 2024

[sycl-post-link] Handle SYCL_EXTERNAL functions in sycl-post-link #3239

Closed

		" any kernel are dropped from the resulting module(s), except from\n"
		" SYCL_EXTERNAL functions.\n"

[sycl-post-link] Do not drop SYCL_EXTERNAL functions #3793

[sycl-post-link] Do not drop SYCL_EXTERNAL functions #3793

Uh oh!

Conversation

DenisBakhvalov commented May 20, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

DenisBakhvalov commented May 20, 2021

Uh oh!

DenisBakhvalov commented May 21, 2021

Uh oh!

romanovvlad commented May 21, 2021

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon commented May 24, 2021

Uh oh!

DenisBakhvalov commented May 25, 2021

Uh oh!

Fznamznon commented May 25, 2021

Uh oh!

DenisBakhvalov commented May 25, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DenisBakhvalov commented May 27, 2021 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

DenisBakhvalov commented Jun 3, 2021

Uh oh!

Uh oh!

DenisBakhvalov commented Jun 4, 2021

Uh oh!

Fznamznon commented Jun 7, 2021

Uh oh!

AlexeySachkov commented Jun 8, 2021

Uh oh!

kbobrovs left a comment • edited by DenisBakhvalov Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

DenisBakhvalov commented Jun 11, 2021

Uh oh!

DenisBakhvalov commented Jun 15, 2021

Uh oh!

DenisBakhvalov commented Jun 16, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

DenisBakhvalov commented May 20, 2021 •

edited

Loading

DenisBakhvalov commented May 27, 2021 •

edited

Loading

kbobrovs left a comment •

edited by DenisBakhvalov

Loading